共享儲存器

share memory是MP擁有的資源，因為它是一個片上儲存器，所以訪問共享儲存器比本地儲存器和全域性儲存器要快得多。實際上共享儲存器的延遲大約比沒有快取的全域性儲存器低100倍(假設執行緒之間沒有bank衝突)。共享儲存器被分配給每個執行緒塊，所以塊內的執行緒可以訪問同一個共享儲存器。執行緒可以訪問共享記憶體中由同一執行緒塊中的其他執行緒從全域性記憶體載入的資料

執行緒同步

當線上程之間共享資料時，我們需要小心以避免競態條件(race conditions)。因為執行緒塊中的執行緒之間雖然邏輯上是並行的，但是物理上並不是同時執行的。讓我們假設執行緒A和執行緒B分別從全域性儲存器中載入了一個數據並且將它存到了共享儲存器。然後，執行緒A想要從共享儲存器中讀取B的資料，反之亦然。我們還要假設執行緒A和B位於不同的warp。如果在A嘗試讀取B的資料時，B還未寫入，這樣就會導致未定義的行為和錯誤的結果。

為了保證在並行執行緒協作時得到正確的結果，我們必須對執行緒進行同步。CUDA提供了一個簡單的柵欄同步原語，__syncthreads()。每個執行緒只能在塊中所有的執行緒執行完__syncthreads()函式後，才能繼續執行__syncthreads()的語句。因此我們可以在向共享儲存器存資料後以及在向共享儲存器載入資料前呼叫__syncthreads()，這樣就避免了上面所描述的競態條件(race conditions)。我們必須要牢記 __syncthreads()被用在分支程式碼塊中是未定義的行為 ，很可能會導致死鎖——執行緒塊中所有的執行緒必須在同一點呼叫__syncthreads()

例子：陣列逆序

在裝置程式碼中宣告共享記憶體要使用__shared__變數宣告說明符。兩種方式申請共享記憶體，不同之處在於共享記憶體陣列的宣告以及核函式的呼叫。：

1. 靜態共享記憶體， 大小在編譯時可確定
2. 動態共享記憶體， 大小在執行時確定

靜態共享記憶體

    __global__ void staticReverse(int *d, int n)
    {
      __shared__ int s[64];  //靜態共享記憶體陣列宣告，陣列長度在編譯時就確定
      int t = threadIdx.x;
      int tr = n-t-1;
      s[t] = d[t];
      __syncthreads();
      d[t] = s[tr];
    }
    
    int main(){
        ...
        staticReverse <<< 1, n >>> (d_d, n);
        ...
    }
 

靜態共享記憶體陣列宣告，陣列長度在編譯時就確定。

動態共享記憶體

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[]; //靜態共享記憶體陣列宣告，陣列長度在執行譯時確定
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

int main(){
    ...
    dynamicReverse <<<1 , n ,n * sizeof(int) >>> (d_d, n);
    ...
}

靜態共享記憶體陣列宣告時，需要加 extern 關鍵字。
當核函式被啟動時，陣列大小從第三個執行配置引數被隱式地確定。

在一個核函式中動態地申請多個數組

如果你想在一個核函式中動態地申請多個數組時該怎麼辦呢？你必須在首先申請一個單獨的未指定大小的extern陣列，然後使用指標將它分為多個數組，如下所示：

extern __shared__ int s[];
int *integerData = s;                        // nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF];      // nC chars

這樣的話，你需要在核函式中這樣指定共享記憶體的大小：

myKernel<<<gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)>>>(...);

bank conflict

參考

對不同bank的訪問可同時進行 ：為了獲得較高的記憶體頻寬，共享儲存器被劃分為多個大小相等的儲存器模組，稱為bank，可以被同時訪問。因此任何跨越b個不同bank的n個地址的讀寫操作可以被同時進行，這樣就大大提高了整體頻寬 ——可達到單獨一個bank頻寬的b倍。

同一warp中的所有執行緒訪問同一bank可同時進行（廣播） ： 然而，如果多個執行緒請求的地址對映到相同的記憶體bank，那麼訪問就會被順序執行。硬體會把衝突的記憶體請求分為儘可能多的單獨的沒有衝突的請求，這樣就會減少一定的頻寬，減少的因子與衝突的記憶體請求個數相等。當然，也有例外的情況：當一個warp中的所有執行緒訪問同一個共享記憶體地址時，就會產生一次廣播。計算能力為2.0及以上的裝置還可以多次廣播共享記憶體訪問，這意味著一個warp中任意數量的執行緒對於同一位置的多次訪問也可以同時進行。

為了儘量減少bank衝突，理解共享記憶體地址如何對映到bank是非常重要的。計算能力5.0的裝置，共享記憶體的bank是這樣組織的：連續的32-bits字被分配到連續的bank中，每個bank的頻寬是每個時鐘週期32bits。；對於計算能力3.0的裝置，每個bank的頻寬是每個時鐘週期64bits；對於計算能力2.0的裝置每個bank的頻寬是每兩個時鐘週期32bits

註釋：

• 對於計算能力1.x的裝置，warp的大小是32而bank的數量是16。一個warp中執行緒對共享記憶體的請求被劃分為兩次請求：一個請求是前半個warp的另一個請求時後半個warp的。注意如果每個bank中只有一個記憶體地址是被半個warp中的執行緒訪問的話，是不會有bank衝突的。
• 對於計算能力為2.x的裝置，warp的大小是32而bank的數量也是32。一個warp中執行緒對共享記憶體的請求不會像計算能力1.x的裝置那樣被劃分開，這就意味著同一個warp中的前半個warp中的執行緒與後半個warp中的執行緒會有可能產生bank衝突的。
• 計算能力為3.x的裝置的bank大小是可以配置的，我們可以通過函式cudaDeviceSetSharedMemConfig()來設定，要麼設定為4位元組(預設為cudaSharedMemBankSizeFourByte)，要麼設定為8位元組(cudaSharedMemBankSizeEightByte)。當訪問雙精度的資料時，將bank大小設定為8位元組可以幫助避免bank衝突。

當一個warp中的不同執行緒訪問一個bank中的不同的字地址時，就會發生bank衝突。

如果沒有bank衝突的話，共享記憶體的訪存速度將會非常的快，大約比全域性記憶體的訪問延遲低100多倍，但是速度沒有暫存器快。然而，如果在使用共享記憶體時發生了bank衝突的話，效能將會降低很多很多。在最壞的情況下，即一個warp中的所有執行緒訪問了相同bank的32個不同字地址的話，那麼這32個訪問操作將會全部被序列化，大大降低了記憶體頻寬。

注意：不同warp中的執行緒之間不存在什麼bank衝突。

share memory的地址對映方式

要解決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衝突的情況 ：

這裡我們需要注意，下面這兩種情況是兩種特殊情況：

上圖中，所有的執行緒都訪問了同一個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衝突的例子

下面我們以平行計算中的經典的歸約演算法為例來做一個簡單的練習。假設有一個大小為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衝突。

// 連續的歸約求和
__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衝突。