共享内存
共享内存和L1缓存同属为一块物理内存，所以他的访问和写入速度是和L1缓存相当的。并且可以通过CUDA提供的API，设置L1缓存和共享内存的大小，调整他们所占的比例，如果你几乎用不到共享内存，就可以将共享内存的大小调到最小，这样可以获得更大的L1缓存。
共享内存对线程束占用率的影响
和L1缓存一样，共享内存也是每个SM单元都独有一块共享内存块，并且在核函数被调用时，共享内存是在线程块中被共享的。
也就是说：
共享内存使用量 = 每个SM单元中线程块的数量 X 核函数中用到的共享内存大小
如果共享内存资源用尽，那么SM单元则无法启动更多的线程块，当线程块的大小不足以满足SM单元的最大线程束的限制，那么理论线程束就无法达到100%。
举个例子，假如每个SM单元有10K的共享内存大小,线程块的大小为32。而核函数中，用到了4k的共享内存，那么这个SM单元最多能启动两个线程块。
静态共享内存
__global__ void MyKernel() {
__shared__ int i; // Shared int
__shared__ float f_array[10]; // 10 shared floats
// ... Some other code
}
在核函数中使用 shared关键字声明的变量将放入共享内存中的。
静态共享内存在编译期就需要决定变量的长度。
动态共享内存
// Kernel
__global__ void SomeKernel() {
// The size of the following is set by the host
extern __shared__ char sharedbuffer[];
}
int main() {
// Other code
// Host launch configuration
SomeKernel<<<10, 23, 32>>>();
// Other code
}
动态分配共享内存时，需要在调用核函数时使用第三个调用参数，这个调用参数为共享内存变量的长度。
修改共享内存大小
如前文所说，共享内存和L1缓存属于同一块物理内存，可以通过设置为它分配大小。
__host__ cudaError_t cudaFuncSetCacheConfig ( const void* func, cudaFuncCache cacheConfig );
- func 为核函数指针
- cudaFuncCache 共享内存配置类型
- cudaFuncCachePreferNone: 默认情况
- cudaFuncCachePreferShared: 更大的共享内存和更小的L1缓存
- cudaFuncCachePreferL1: 更大的L1缓存和更小的共享内存
- cudaFuncCachePreferEqual: 相同大小的 L1 缓存和共享内存
需要注意的是，这个函数只是建议怎么去分配共享内存的大小，在实际过程中会根据核函数所需要使用到的共享内存来分配，所以并不是强制性的。
共享内存的线程安全
首先，共享内存是在线程块中共享。
如果在线程块内，不同的线程对同一个线程进行读写，或者都进行写操作的时候，并不会引发运行时错误。不过这会出现不正确的结果。
线程块内同步
__syncthreads()
这个函数没有参数也没用返回值，用于线程块内的同步。
使用这个函数会让块内都线程都执行到这个函数，然后再继续往下执行。所以使用时需要注意尽量不要讲这个函数放到分支语句中，避免因为有的线程没有进入分支而造成死锁。
为了避免读写共享内存时出现冲突，得到不正确的结果，可以根据使用场景，使用这个函数对核函数进行同步。
文章评论