CUDA中的共享内存(Shared Memory)

2024年3月14日 510点热度 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 );
  1. func 为核函数指针
  2. cudaFuncCache 共享内存配置类型

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

共享内存的线程安全

首先,共享内存是在线程块中共享。

如果在线程块内,不同的线程对同一个线程进行读写,或者都进行写操作的时候,并不会引发运行时错误。不过这会出现不正确的结果。

线程块内同步

__syncthreads()

这个函数没有参数也没用返回值,用于线程块内的同步。

使用这个函数会让块内都线程都执行到这个函数,然后再继续往下执行。所以使用时需要注意尽量不要讲这个函数放到分支语句中,避免因为有的线程没有进入分支而造成死锁。

为了避免读写共享内存时出现冲突,得到不正确的结果,可以根据使用场景,使用这个函数对核函数进行同步。

参考文献:https://www.syncfusion.com/succinctly-free-ebooks/cuda/shared-memory#:~:text=Shared%20memory%20is%20a%20small,the%20configuration%20of%20the%20kernel.

大脸猫

这个人虽然很勤快,但什么也没有留下!

文章评论

  • jiyouzhan

    这篇文章写得深入浅出,让我这个小白也看懂了!

    2024年5月21日