原理上来说,共享内存是GPU上可受用户控制的一级缓存。在一个SM中,存在着若干cuda core + DP(双精度计算单元) + SFU(特殊函数计算单元)+共享内存+常量内存+纹理内存。相对于全局内存,共享内存的方寸延迟较低,可以达到惊人的1.5TB/s。而全局内存大约只有150GB/s。(最新的NVLINK技术没有考虑在内)。因而共享内存的使用时性能提高的一个重要的因素。但是注意到,将数据拷贝到共享内存中也消耗了部分时间。因而,共享内存仅仅适合存在着数据的重复利用,全局的内存合并或者是线程之间有共享数据的时候,否则直接使用全局内存会更好一些。
下面介绍两种使用共享内存的方法。
1. 创建固定大小的共享内存。(在kernel函数内存定义)
__shared__ float a_in[34];
注意这里的34必须在编译之前指定大小。可以使用宏定义的方式进行。下面的方式是一种错误的示范。
__shared__ float s_in[blockDim.x+2*RAD];
2. 动态申请共享内存数组,声明时需要加上 extern 前缀。
extern __shared__ float a[];
<