共享內存之bank衝突

前面博客中咱們說到了共享內存的使用方法以及一些高級特性，並簡單說明了一下bank衝突，這裏咱們將會經過一些簡單的例子來詳細介紹一下bank衝突。
爲了得到較高的內存帶寬，共享存儲器被劃分爲多個大小相等的存儲器模塊，稱爲bank，能夠被同時訪問。所以任何跨越b個不一樣的內存bank的對n個地址進行讀取和寫入的操做能夠被同時進行，這樣就大大提升了總體帶寬 ——可達到單獨一個bank帶寬的b倍。可是不少狀況下，咱們沒法充分發揮bank的功能，以至於shared memory的帶寬很是的小，這多是由於咱們遇到了bank衝突。

bank衝突

當一個warp中的不一樣線程訪問一個bank中的不一樣的字地址時，就會發生bank衝突。
若是沒有bank衝突的話，共享內存的訪存速度將會很是的快，大約比全局內存的訪問延遲低100多倍，可是速度沒有寄存器快。然而，若是在使用共享內存時發生了bank衝突的話，性能將會下降不少不少。在最壞的狀況下，即一個warp中的全部線程訪問了相同bank的32個不一樣字地址的話，那麼這32個訪問操做將會所有被序列化，大大下降了內存帶寬。

NOTE:不一樣warp中的線程之間不存在什麼bank衝突。

共享內存的地址映射方式

要解決bank衝突，首先咱們要了解一下共享內存的地址映射方式。
在共享內存中，連續的32-bits字被分配到連續的32個bank中，這就像電影院的座位同樣：一列的座位就至關於一個bank，因此每行有32個座位，在每一個座位上能夠「坐」一個32-bits的數據(或者多個小於32-bits的數據，如4個char型的數據，2個short型的數據)；而正常狀況下，咱們是按照先坐完一行再坐下一行的順序來坐座位的，在shared memory中地址映射的方式也是這樣的。下圖中內存地址是按照箭頭的方向依次映射的：

上圖中數字爲bank編號。這樣的話，若是你將申請一個共享內存數組(假設是int類型)的話，那麼你的每一個元素所對應的bank編號就是地址偏移量(也就是數組下標)對32取餘所得的結果，好比大小爲1024的一維數組myShMem：

• myShMem[4]: 對應的bank id爲#4 (相應的行偏移量爲0)

• myShMem[31]: 對應的bank id爲#31 (相應的行偏移量爲0)

• myShMem[50]: 對應的bank id爲#18 (相應的行偏移量爲1)

• myShMem[128]: 對應的bank id爲#0 (相應的行偏移量爲4)

• myShMem[178]: 對應的bank id爲#18 (相應的行偏移量爲5)

典型的bank訪問方式

下面我介紹幾種典型的bank訪問的形式。

下面這這種訪問方式是典型的線性訪問方式(訪問步長(stride)爲1)，因爲每一個warp中的線程ID與每一個bank的ID一一對應，所以不會產生bank衝突。

下面這種訪問雖然是交叉的訪問，每一個線程並無與bank一一對應，但每一個線程都會對應一個惟一的bank，因此也不會產生bank衝突。

下面這種雖然也是線性的訪問bank，但這種訪問方式與第一種的區別在於訪問的步長(stride)變爲2，這就形成了線程0與線程28都訪問到了bank 0，線程1與線程29都訪問到了bank 2...，因而就形成了2路的bank衝突。我在後面會對以不一樣的步長(stride)訪問bank的狀況作進一步討論。

下面這種訪問形成了8路的bank衝突，

這裏咱們須要注意，下面這兩種狀況是兩種特殊狀況：

上圖中，全部的線程都訪問了同一個bank，貌似產生了32路的bank衝突，可是因爲廣播(broadcast)機制(當一個warp中的全部線程訪問一個bank中的同一個字(word)地址時，就會向全部的線程廣播這個字(word))，這種狀況並不會發生bank衝突。

一樣，這種訪問方式也不會產生bank衝突：

這就是所謂的多播機制(multicast)——當一個warp中的幾個線程訪問同一個bank中的相同字地址時，會將該字廣播給這些線程。

NOTE:這裏的多播機制(multicast)只適用於計算能力2.0及以上的設備，上篇博客中已經提到。

數據類型與bank衝突

咱們都知道，當每一個線程訪問一個32-bits大小的數據類型的數據(如int，float)時，不會發生bank衝突。

extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]

可是若是每一個線程訪問一個字節(8-bits)的數據時，會不會發生bank衝突呢？其實這種狀況是不會發生bank衝突的。當同一個字(word)中的不一樣字節被訪問時，也不會發生bank衝突，下面是這種狀況的兩個例子：

extern __shared__ char shrd[];
foo = shrd[baseIndex + threadIdx.x];

extern __shared__ short shrd[];
foo = shrd[baseIndex + threadIdx.x];

訪問步長與bank衝突

咱們一般這樣來訪問數組：每一個線程根據線程編號tid與s的乘積來訪問數組的32-bits字(word)：

extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];

若是按照上面的方式，那麼當s*n是bank的數量(即32)的整數倍時或者說n是32/d的整數倍(d是32和s的最大公約數)時，線程tid和線程tid+n會訪問相同的bank。咱們不難知道若是tid與tid+n位於同一個warp時，就會發生bank衝突，相反則不會。

仔細思考你會發現，只有warp的大小(即32)小於等於32/d時，纔不會有bank衝突，而只有當d等於1時才能知足這個條件。要想讓32和s的最大公約數d爲1，s必須爲奇數。因而，這裏有一個顯而易見的結論：當訪問步長s爲奇數時，就不會發生bank衝突。

bank衝突的例子

既然咱們已經理解了bank衝突，那咱們就小試牛刀，來練習下吧！下面咱們以並行計算中的經典的歸約算法爲例來作一個簡單的練習。

假設有一個大小爲2048的向量，咱們想用歸約算法對該向量求和。因而咱們申請了一個大小爲1024的線程塊，並聲明瞭一個大小爲2048的共享內存數組，並將數據從全局內存拷貝到了該共享內存數組。

咱們能夠有如下兩種方式實現歸約算法：

不連續的方式：

連續的方式：

下面咱們用具體的代碼來實現上述兩種方法。

// 非連續的歸約求和
__global__ void BC_addKernel(const int *a, int *r)
{
    __shared__ int cache[ThreadsPerBlock];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int cacheIndex = threadIdx.x;

    // copy data to shared memory from global memory
    cache[cacheIndex] = a[tid];
    __syncthreads();

    // add these data using reduce
    for (int i = 1; i < blockDim.x; i *= 2)
    {
        int index = 2 * i * cacheIndex;
        if (index < blockDim.x)
        {
            cache[index] += cache[index + i];
        }
        __syncthreads();
    }

    // copy the result of reduce to global memory
    if (cacheIndex == 0)
        r[blockIdx.x] = cache[cacheIndex];
}

上述代碼實現的是非連續的歸約求和，從int index = 2 * i * cacheIndex和cache[index] += cache[index + i];兩條語句，咱們能夠很容易判斷這種實現方式會產生bank衝突。當i=1時，步長s=2xi=2，會產生兩路的bank衝突；當i=2時，步長s=2xi=4，會產生四路的bank衝突...當i=n時，步長s=2xn=2n。能夠看出每一次步長都是偶數，所以這種方式會產生嚴重的bank衝突。

NOTE:在《GPU高性能運算之CUDA》這本書中對實現不連續的歸約算法有兩種代碼實現方式，但筆者發現書中的提到(p179)的兩種所謂相同計算邏輯的函數reduce0和reduce1，其實具備本質上的不一樣。前者不會發生bank衝突，然後者(即本文中所使用的)纔會產生bank衝突。因爲前者線程ID要求的條件比較「苛刻」，只有知足tid % (2 * s) == 0的線程纔會執行求和操做(sdata[tid]+=sdata[tid+i])；然後者只要知足index(2 * s * tid，即線程ID的2xs倍)小於線程塊的大小(blockDim.x)便可。總之，前者在進行求和操做(sdata[tid]+=sdata[tid+i])時，線程的使用一樣是不連續的，即當s=1時，線程編號爲0,2,4,...,1022；然後者的線程使用是連續的，即當s=1時，前512個線程(0,1,2,...,511)在進行求和操做(sdata[tid]+=sdata[tid+i])，然後512個線程是閒置的。前者不會出現多個線程訪問同一bank的不一樣字地址，然後者正如書中所說會產生嚴重的bank衝突。(書中用到的s與本文中屢次用到的步長s不是同一個變量，注意不要混淆這兩個變量)固然這些只是筆者的想法，若有不一樣，歡迎來與我討論，郵箱:<chaoyanglius@outlook.com>。

// 連續的歸約求和
__global__ void NBC_addKernel2(const int *a, int *r)
{
    __shared__ int cache[ThreadsPerBlock];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int cacheIndex = threadIdx.x;

    // copy data to shared memory from global memory
    cache[cacheIndex] = a[tid];
    __syncthreads();

    // add these data using reduce
    for (int i = blockDim.x / 2; i > 0; i /= 2)
    {
        if (cacheIndex < i)
        {
            cache[cacheIndex] += cache[cacheIndex + i];
        }
        __syncthreads();
    }

    // copy the result of reduce to global memory
    if (cacheIndex == 0)
        r[blockIdx.x] = cache[cacheIndex];
}

因爲每一個線程的ID與操做的數據編號一一對應，所以上述的代碼很明顯不會產生bank衝突。

參考資料

1. C語言程序設計現代方法，[美]K.N.King著，人民郵電出版社

2. 英偉達CUDA C programming guide v7.0

3. 威斯康星大學仿真實驗室CUDA課程講義10-07-2013:http://sbel.wisc.edu/Courses/...

4. GPU高性能運算之CUDA，張舒，褚豔利，中國水利水電出版社