CUDA中的共享内存（Shared Memory）

2024年3月14日 4386点热度 0人点赞 1条评论

共享内存

共享内存和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 缓存和共享内存

需要注意的是，这个函数只是建议怎么去分配共享内存的大小，在实际过程中会根据核函数所需要使用到的共享内存来分配，所以并不是强制性的。