CUDA——性能优化之共享内存

一、共享内存的结构

1)什么是共享内存？

共享内存是GPU的一种稀缺资源，它位于芯片上，所以共享内存空间要比本地和全局内存空间快得多。对于warp里的所有线程，只要线程之间没有任何存储体冲突(bank conflict)，访问共享内存就与访问寄存器一样快。

2)什么是存储体（bank）？

共享内存被划分为同样大小的、可以同时访问的内存块，名为存储体。在计算能力为1.x的设备上，存储体数为16，在2.0及以上的设备，存储体数为32。

存储体的存在可使共享内存获得高内存带宽，假设有n个存储体，此时访问n个分别位于不同bank的地址时是同时进行的，最后获得的有效带宽就是单个模块的带宽的N倍，因为只需要发射一次指令。

3）共享内存是如何映射到存储体的？

共享内存空间的存储体组织为：连续的32位（4个字节）分配到连续的存储体中，每个存储体的带宽为 32位/2个时钟周期

一维共享内存

如下的方式申请一个一维的共享内存 (32个存储体)

__shared__ float sData[64];
// sData[0]-sData[31]分别对应bank[0]-bank[31];
// sData[32]-sData[63]分别对应bank[0]-bank[31];

此时，
sData[0]与sData[32]位于同一个存储体bank[0]
sData[1]与sData[33]位于同一个存储体bank[1]
。。。。。。
sData[31]与sData[63]位于同一个存储体bank[31]

二维共享内存

二维共享内存其实可以展开成一维内存，其映射方式跟一维一样。

__shared__ float sData[32][32];

由于上面共享内存的每一行大小为32，刚好等于存储体个数，那么此时共享内存的每一列就是一个bank。

二、避免共享存储体冲突（bank conflict）

同一个warp里的线程访问同一个bank里不同的地址时，会出现bank conflict，如果访问的不同的地址的个数为n，那么此种情况称为n路存储器冲突（n-way bank conflict）。

1）同一个warp访问共享内存的同一个bank的不同地址，所产生的bank conflict。

a.通过改变数据在共享内存中的排列方式，使其映射到bank时，原本在同一个bank的不同地址的数据分开到不同的bank。

二维共享内存例子：

size_t ix=threadIdx.x+blockIdx.x*blockDim.x;//0-31
size_t iy=threadIdx.y+blockIdx.y*blockDim.y;//0-31
__shared__ float sData[32][32+1];
//此处的共享内存改变了数据的排列方式，通过映射到bank上，使得原本在同一个bank的地址偏移到了不同的bank上。
if(ix<width&&iy<height)
{
	sData[ix][iy]=Input_data[ix+iy*width];
	__syncthreads();
	Output[ix+iy*width]=sData[ix][iy];
}

一维共享内存例子：

__global__ void kernel1D(float *Input_data, float *Output_data,unsigned int length)
{
	size_t tid = threadIdx.x;//0-63
	size_t iy = tid >> (int)log2(32);//0-1
	size_t ix = tid & (32 - 1);//0-31

	__shared__ float sData[64 + 2];
	float data;
	if (tid < length)
	{
		sData[ix + iy * 33] = Input_data[tid];
		__syncthreads();
		data = sData[iy * 33];
	}
}

b.如想对共享内存的列进行访问，可让列方向上的数据存入行方向上。例如：你想对矩阵的列进行操作，但是如果直接将矩阵一一对应映射到共享内存上，此时访问列方向的数据时，会产生bank conflict。尝试将矩阵转置再进入核函数。

2）同一个warp的多条线程同时访问同一个bank的同一地址，所产生的bank conflict

共享内存具有广播机制，当处理一个内存读取请求时，可以读取一个32-位字并同时广播到多个线程。当warp的多个线程从含有同一个32-位字的地址读取时，这将减少存储体冲突的数目。但前提是 该地址得为广播字

选择哪个字作为广播字，以及在每个周期选择哪个存储体地址都不是特定的。所以，当只有一部分的线程去读取同一个32-位字的地址时，需要我们自己选中广播字

以下代码会产生的bank conflict

//tid 为0-63;下面代码存在bank conflict
if ((tid&(32-1))<32)
{
	Output_data[tid]=sData[tid];
	__syncthreads();
}
else
{
	Output_data[tid]=sData[63];
	__syncthreads();
}

解决方法：先把sData[63]进行广播

Output_data[tid]=sData[63];
if ((tid&(32-1))>32)
{
	Output_data[tid]=sData[tid];
	__syncthreads();
}

参考：

https://blog.csdn.net/endlch/article/details/47043069
https://blog.csdn.net/smsmn/article/details/6336060