[譯]在CUDA C/C++中使用共享存儲器

本文翻譯自NVIDIA官方博客Parallel Forall，內容僅供參考，若有疑問請訪問原網站:https://devblogs.nvidia.com/p...。

在之前發佈的文章中，咱們學習了被一組線程訪問的全局內存如何被合併爲一次事務以及對於不一樣的CUDA硬件，對齊和步長如何影響合併訪問。對於最近的CUDA硬件，沒有對齊的數據訪問並非什麼大問題。然而不管是哪一代的CUDA硬件，跨越全局存儲器都是個大問題，並且在不少狀況下也是很難避免的，例如沿着第二和更高維度訪問多維陣列中的元素時。可是，若是咱們使用共享存儲器的話，也是有可能進行合併訪問的。在我向你說明如何避免直接跨越全局存儲器以前，我首先須要詳細地介紹一下共享存儲器。

共享存儲器

由於它是一個片上存儲器，因此共享存儲器比本地存儲器和全局存儲器要快得多。實際上共享存儲器的延遲大約比沒有緩存的全局存儲器低100倍(假設線程之間沒有bank衝突，在以後的文章中咱們會介紹)。共享存儲器被分配給每一個線程塊，因此塊內的線程能夠訪問同一個共享存儲器。線程能夠訪問共享內存中由同一線程塊中的其餘線程從全局內存加載的數據。這種能力（與線程同步相結合）具備許多用途，例如用戶管理的數據高速緩存，高性能並行協做算法（例如並行歸約），而且在其它狀況不可能的狀況下促進全局存儲器的合併訪問 。

線程同步

當在線程之間共享數據時，咱們須要當心以免競態條件(race conditions)，由於線程塊中的線程之間雖然邏輯上是並行的，可是物理上並非同時執行的。讓咱們假設線程A和線程B分別從全局存儲器中加載了一個數據而且將它存到了共享存儲器。而後，線程A想要從共享存儲器中讀取B的數據，反之亦然。咱們還要假設線程A和B位於不一樣的warp。若是在A嘗試讀取B的數據時，B還未寫入，這樣就會致使未定義的行爲和錯誤的結果。

爲了保證在並行線程協做時獲得正確的結果，咱們必須對線程進行同步。CUDA提供了一個簡單的柵欄同步原語，__syncthreads()。每一個線程只能在塊中全部的線程執行完__syncthreads()函數後，才能繼續執行__syncthreads()的語句。所以咱們能夠在向共享存儲器存數據後以及在向共享存儲器加載數據前調用__syncthreads()，這樣就避免了上面所描述的競態條件(race conditions)。咱們必需要牢記__syncthreads()被用在分支代碼塊中是未定義的行爲，極可能會致使死鎖——線程塊中全部的線程必須在同一點調用__syncthreads()

共享內存的例子

在設備代碼中聲明共享內存要使用__shared__變量聲明說明符。在覈函數中有多種方式聲明共享內存，這取決於你要申請的內存大小是在編譯時肯定仍是在運行時肯定。下面完整的代碼(能夠在Github上下載)展現了使用共享內存的兩種方法。

#include <stdio.h>

__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];
}

__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(void)
{
  const int n = 64;
  int a[n], r[n], d[n];

  for (int i = 0; i < n; i++) {
    a[i] = i;
    r[i] = n-i-1;
    d[i] = 0;
  }

  int *d_d;
  cudaMalloc(&d_d, n * sizeof(int));

  // run version with static shared memory
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  staticReverse<<<1,n>>>(d_d, n);
  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++)
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);

  // run dynamic shared memory version
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
  cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++)
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}

上面的代碼使用共享存儲器對大小爲64的數組進行逆序處理。這兩個核函數十分類似，不一樣之處在於共享內存數組的聲明以及核函數的調用。

靜態的共享內存

若是共享內存數組的大小在編譯時就能夠肯定，就像在上節代碼中staticReverse核函數中寫的那樣，咱們就能夠顯式地聲明固定大小的數組，下面是咱們聲明的s數組：

__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];
}

在這個核函數中，t和tr分別表明了原始和倒序以後數組的下標索引。每一個線程使用語句s[t] = d[t]將全局內存的數據拷貝到共享內存，反向工做是經過語句d[t] = s[tr]來完成的。可是在執行線程訪問共享內存中被線程寫入的數據前，記住要使用__syncthreads()來確保全部的線程都已經徹底將數據加載到共享內存。

在這個例子中，使用共享內存是用於促進全局內存合併訪問(在舊的CUDA設備上，計算能力1.1或更低)。對於讀取和寫入都實現了最優的全局存儲器合併，由於全局內存老是經過線性對齊的索引t來訪問的。反向索引tr僅用於訪問共享存儲器，其不具備全局存儲器的順序訪問限制，所以不能得到最佳性能。共享內存的惟一性能問題是bank衝突，咱們以後會作討論。

NOTE:注意在計算能力爲1.2或更高版本的設備上，內存系統仍然能夠徹底地合併訪問，即便是反向的保存在全局存儲器中。這一技術在其餘訪問模式下也是頗有用的，我會在下一篇博客中介紹。

動態的共享內存

另外一個核函數使用了動態分配共享內存的方式，這主要用於共享內存的大小在編譯時不能肯定的狀況。在這種狀況下，每一個線程塊中共享內存的大小必須在覈函數第三個執行配置參數中指定(以字節爲單位)，以下所示：

dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);

該動態共享內存的核函數dynamicReverse()使用了未指定大小的extern數組語法(extern __shared__ int s[])來聲明共享內存數組。

NOTE:注意中括號與extern說明符。

當核函數被啓動時，數組大小從第三個執行配置參數被隱式地肯定。該核函數其他部分的代碼與staticReverse()核函數相同。

而若是你想在一個核函數中動態地申請多個數組時該怎麼辦呢？你必須在首先申請一個單獨的未指定大小的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衝突

爲了得到較高的內存帶寬，共享存儲器被劃分爲多個大小相等的存儲器模塊，稱爲bank，能夠被同時訪問。所以任何跨越b個不一樣bank的n個地址的讀寫操做能夠被同時進行，這樣就大大提升了總體帶寬 ——可達到單獨一個bank帶寬的b倍。

然而，若是多個線程請求的地址映射到相同的內存bank，那麼訪問就會被順序執行。硬件會把衝突的內存請求分爲儘量多的單獨的沒有衝突的請求，這樣就會減小必定的帶寬，減小的因子與衝突的內存請求個數相等。固然，也有例外的狀況：當一個warp中的全部線程訪問同一個共享內存地址時，就會產生一次廣播。計算能力爲2.0及以上的設備還能夠屢次廣播共享內存訪問，這意味着一個warp中任意數量的線程對於同一位置的屢次訪問也能夠同時進行。

譯者注:這裏關於warp的多播與bank衝突原文並未詳細介紹，詳細內容及例子能夠參考CUDA programming guide。我在後續的博客中也會詳細介紹這部分。

爲了儘可能減小bank衝突，理解共享內存地址如何映射到bank是很是重要的。共享內存的bank是這樣組織的：連續的32-bits字被分配到連續的bank中，每一個bank的帶寬是每一個時鐘週期32bits。

譯者注:這裏不一樣計算能力的bank的帶寬是不一樣的，原文提到的帶寬大小是計算能力5.0的設備，對於計算能力2.0的設備每一個bank的帶寬是每兩個時鐘週期32bits；對於計算能力3.0的設備，每一個bank的帶寬是每一個時鐘週期64bits。詳情請參考CUDA C programming guide。

對於計算能力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衝突。

配置共享內存的數量

在計算能力爲2.x和3.x的設備上，每一個多處理器有64KB的片上內存，它能夠被劃分爲L1高速緩存和共享內存。對於計算能力爲2.x的設備，總共有兩種設置：48KB的共享內存/16KBL1高速緩存和16KB的共享內存/16KB的L1高速緩存。咱們能夠在運行時使用cudaDeviceSetCacheConfig()在主機端爲全部的核函數配置或者使用cudaFuncSetCacheConfig()爲單個的核函數配置。它們有三個選項能夠設置：cudaFuncCachePreferNone(在共享內存和L1中不設置首選項，即便用默認設置), cudaFuncCachePreferShared(共享內存大於L1), 和cudaFuncCachePreferL1(L1大於共享內存)。驅動程序將按照指定的首選項，除非核函數中每一個線程塊須要比指定配置中更多的共享內存。在計算能力3.x的設備上容許有第三種設置選項——32KB的共享內存/32KB的L1高速緩存，能夠經過cudaFuncCachePreferEqual選項設置。

總結

對於寫出高性能的CUDA代碼，共享內存的確是一個十分強大的特性。因爲共享內存位於片上，因此訪問共享內存比訪問全局內存快不少。因爲共享內存在線程塊中能夠被線程共享，因此才提供了相應的機制來保證線程的正常協做。使用共享內存來利用這種線程協做的一種方法是啓用全局內存的合併訪問，正如如本文中的數組逆序所演示的。在使用共享內存來使數組逆序的例子中，咱們可使用單位步長執行全部全局內存讀取和寫入，從而在任何CUDA GPU上實現徹底地合併訪問。