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

FiBird發表於2019-05-10
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)
", 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)
", 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上實現完全地合併訪問。

相關文章