cuda GPU 編程之共享內存的使用
阿新 • • 發佈：2017-10-02
ret 指定大小 最新 宏定義 編程 int 重要 core 申請
原理上來說，共享內存是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[];
並且，在調用內核函數的時候，需要在<<<>>>內加上第三個參數來指明所需分配的共享內存的字節大小。
const size_t smemSize=(TPB+ 2*RAD)*sizeof(float); ddkernel<<<Grids, Blocks,smemSize>>>(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<RAD){
s_in[s_idx-RAD]=d_in[i-RAD];
s_in[s_idx+blockDim.x]=d_in[i+blockDim.x];
}
__syncthread();
}
cuda GPU 編程之共享內存的使用