共享内存是GPU上可受用户控制的一级缓存

原则上，共享内存是 GPU 上可由用户控制的一级缓存。在一个SM中，有几个cuda核心+DP（双精度计算单元）+SFU（特殊功能计算单元）+共享内存+常量内存+纹理内存。与全局内存相比，共享内存的平方英寸延迟更低，可以达到惊人的1.5TB/s。全局内存只有150GB/s左右。 （不考虑最新的 NVLINK 技术）。因此，共享内存的使用是性能提升的重要因素。但是请注意，将数据复制到共享内存也需要一些时间。因此，共享内存只适用于数据重用、全局内存整合或线程间有共享数据时，否则直接使用全局内存会更好。

下面介绍两种使用共享内存的方法。

1. 创建一个固定大小的共享内存。 （定义在内核函数内存中）

__shared__ float a_in[34];

注意这里的34必须在编译前指定大小。它可以通过宏定义的方式来完成。以下方法为错误示范。

__shared__ float s_in[blockDim.x+2*RAD];

2. 动态申请一个共享内存数组，声明时需要加上extern前缀。

并且，在调用内核函数时，需要添加第三个参数，以字节为单位指定要分配的共享内存的大小。

const size_t smemSize=(TPB+ 2*RAD)*sizeof(float);
ddkernel<<>>(paramenter);

共享内存分配后，可以将全局内存复制到共享内存中。基本方案是每个线程从全局索引位置读取元素并存储在共享内存中。使用共享内存的时候也要注意数据交叉的存在，把边界上的数据拷贝进去。

__global__ 
void ddkernel(paramenter)
{
    const int i=threadIdx.x+blockDim.x*BlockIdx.x;
    if(i.size)return;
    const int s_idx=threadIdx.x+RAD;
    extern __shared__ float s_in[];
    
    s_in[s_idx]=d_in[i];
    if(threadIdx.x

					

本文来自电脑杂谈，转载请注明本文网址：
http://www.pc-fly.com/a/shoujiruanjian/article-380145-1.html