CUDA_共享內存、訪存機制、訪問優化

---

---

共享內存簡介

共享內存時受用戶控制的一級緩存，共享存儲器爲片內高速存儲器，是一塊能夠被同一block中的全部線程訪問的可讀寫存儲器。訪問共享存儲器的速度幾乎和訪問寄存器同樣快(相對而言，不是十分嚴謹的說法，真實狀況是，共享內存的延時極低，大約1.5T/s的帶寬，遠高於全局內存的190G/s,此速度是寄存器的1/10)，是實現線程間通訊的延遲最小的方法。共享存儲器能夠用於實現多種功能，若是用於保存共用的計數器或者block的公用結果。

計算能力1.0、1.一、1.二、1.3硬件中，每一個SM的共享存儲器的大小爲16KByte，被組織爲16個bank，對共享存儲器的動態與靜態分配與初始化

int main(int argc, char** argv) 
{
    // ...
    testKernel<<<1, 10, mem_size >>>(d_idata, d_odata);
    // ...
    CUT_EXIT(argc, argv);
}

__global__ void testKernel(float* g_idata, float* g_odata)
{
    // extern聲明，大小由主機端程序決定。動態聲明
    extern __shared__ float sdata_dynamic[];

    // 靜態聲明
    __shared__ int sdata_static[16];

    // 注意shared memory不能再定義時初始化
    sdata_static[tid] = 0;
}

注意，將共享存儲器中的變量聲明爲外部數據時，例如

extern __shared__ float shared[];

數組的大小將在kernel啓動時肯定，經過其執行參數肯定。經過這種方式定義的全部變量都開始於相同的地址，所以數組中的變量的佈局必須經過偏移量顯示管理。例如，若是但願在動態分配的共享存儲器得到與如下代碼對應的內容：

short array0[128];
float array1[64];
int array2[256];

應該按照以下的方式對應定義：

extern __shared__ char array[];
// __device__ or __global__ function
__device__ void func()
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[128];
    int* array2 = (int*)&array1[64];
}

共享內存架構

共享內存時基於存儲器切換的架構（bank-switched architecture）.爲了可以在並行訪問時得到高帶寬，共享存儲器被劃分爲大小相等，不能被同時訪問的存儲器模塊，稱爲bank。因爲不一樣的存儲器模塊能夠互不干擾的同時工做，所以對位於n個bank上的n個地址的訪問可以同時進行，此時有效帶寬就是隻有一個bank的n倍。

若是half-warp請求訪問的多個地址位於同一個bank中，就會出現bank conflict。因爲存儲器模塊在一個時刻沒法響應多個請求，所以這些請求就必須被串行的完成。硬件會將形成bank conflict的一組訪存請求劃分爲幾回不存在conflict的獨立請求，此時的有效帶寬會下降與拆分獲得的不存在conflict的請求個數相同的倍數。例外狀況：一個half-warp中的全部線程都請求訪問同一個地址時，會產生一次廣播，此時反而只須要一次就能夠響應全部線程的請求。

bank的組織方式是：每一個bank的寬度固定爲32bit，相鄰的32bit字被組織在相鄰的bank中，每一個bank在每一個時鐘週期能夠提供32bit的帶寬。

在費米架構的設備上有32個存儲體，而在G200與G80的硬件上只有16個存儲體。每一個存儲體能夠存4個字節大小的數據，足以用來存儲一個單精度浮點型數據，或者一個標準的32位的整型數。開普勒架構的設備還引入了64位寬的存儲體，使雙精度的數據無需在跨越兩個存儲體。不管有多少線程發起操做，每一個存儲體每一個週期只執行一次操做。

若是線程束中的每一個線程訪問一個存儲體，那麼全部線程的操做均可以在一個週期內同時執行。此時無須順序地訪問，由於每一個線程訪問的存儲體在共享內存中都是獨立的，互不影響。實際上，在每一個存儲體與線程之間有一個交叉開關將它們鏈接，這在字的交換中頗有用。

此外，當線程束中的全部線程同時訪問相同地址的存儲體時，使用共享內存會有很大幫助，同常量內存同樣，當全部線程訪問同一地址的存儲單元時，會觸發一個廣播機制到線程束中的每一個線程中。一般0號線程會寫一個值而後與線程束中的其餘線程進行通訊。

共享存儲訪問優化

在訪問共享存儲器的時候，須要着重關注如何減小bank conflict.產生bank conflict會形成序列化訪問，嚴重下降有效帶寬。

對於計算能力1.x設備，每一個warp大小都是32個線程，而一個SM中的shared memory被劃分爲16個bank（0-15）。一個warp中的線程對共享存儲器的訪問請求會被劃分爲2個half-warp的訪問請求，只有處於同一half-warp內的線程纔可能發生bank conflict，而一個warp中位於前half-warp的線程與位於後half-warp的線程間則不會發生bank conflict。

沒有bank conflic的共享存儲器訪問示例（線程從數組讀取32bit字場景）：

產生bank conflict的共享存儲器訪問示例（線程從數組讀取32bit字場景）：

若是每一個線程訪問的數據大小不是32bit時，也會產生bank conflict。例如如下對char數組的訪問會形成4way bank conflict：

__shared__ char shared[32];
char data = shared[BaseIndex + tid];

此時，shared[0]、shared[1]、shared[2]、shared[3]屬於同一個bank。對一樣的數組，按照下面的形式進行訪問，則能夠避免bank conflict問題：

char data = shared[BaseIndex + 4* tid];

---

對於一個結構體賦值會被編譯爲幾回訪存請求，例如：

__shared__ struct type shared[32];
struct type data = shared[BaseIndex + tid];

假如type的類型有以下幾種：

// type1
struct type {
	float x, y, z;
};

// type2
struct type {
	float x, y;
};

// type3
struct type {
	float x;
	char c;
};

若是type定義爲type1，那麼type的訪問會被編譯爲三次獨立的存儲器訪問，每一個結構體的同一成員之間有3個32bit字的間隔，因此不存在bank conflict。（沒有bank conflic的共享存儲器訪問示例中場景c）

若是type定義爲type2，那麼type的訪問會被編譯爲兩個獨立的存儲器訪問，每一個結構體成員都有2個32bit字的間隔，線程ID相隔8的線程間就會發生bank conflict。（產生bank conflict的共享存儲器訪問示例中場景b）

若是type定義爲type3，那麼type的訪問會被編譯爲兩個獨立的存儲器訪問，每一個結構體成員都是經過5byte的間隔來訪問，因此總會存在bank conflict。

---

shared memory訪存機制

shared memory採用了廣播機制，在響應一個對同一個地址的讀請求時，一個32bit能夠被讀取的同時會廣播給不一樣的線程。當half-warp有多個線程讀取同一32bit字地址中的數據時，能夠減小bank conflict的數量。而若是half-warp中的線程全都讀取同一地址中的數據時，則徹底不會發生bank conflict。不過，若是half-warp內有多個線程要對同一地址進行寫操做，此時則會產生不肯定的結果，發生這種狀況時應該使用對shared memory 的原子操做。

對不一樣地址的訪存請求，會被分爲若干個處理步，每兩個執行單元週期完成一步，每步都只處理一個conflict-free的訪存請求的子集，知道half-warp的全部線程請求均完成。在每一步中都會按照如下規則構建子集：

（1）從還沒有訪問的地址所指向的字中，選出一個做爲廣播字；

（2）繼續選取訪問其餘bank，而且不存在bank conflict的線程，再與上一步中廣播字對應的線程一塊兒構建一個子集。在每一個週期中，選擇哪一個字做爲廣播字，以及選擇哪些與其餘bank對應的線程，都是不肯定的。

參考：

《高性能運算之CUDA》

《CUDA並行程序設計 GPU編程指南》

《GPU高性能編程 CUDA實戰》

《CUDA專家手冊 GPU編程權威指南》