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]

二維共享記憶體

二維共享記憶體其實可以展開成一維記憶體，其映射方式跟一維一樣。

由于上面共享記憶體的每一行大小為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