共享內存是編寫優(yōu)化良好的 CUDA 代碼的一個強大功能。共享內存的訪問比全局內存訪問快得多，因為它位于芯片上。因為共享內存由線程塊中的線程共享，它為線程提供了一種協(xié)作機制。利用這種線程協(xié)作使用共享內存的一種方法是啟用全局內存合并，如本文中的數(shù)組反轉所示。通過使用 CUDA GPU 共享內存，我們可以在 GPU 上執(zhí)行所有讀操作。在下一篇文章中，我將通過使用共享內存來優(yōu)化矩陣轉置來繼續(xù)我們的討論。

在 上一篇文章 中，我研究了如何將一組線程訪問的全局內存合并到一個事務中，以及對齊和跨步如何影響 CUDA 各代硬件的合并。對于最新版本的 CUDA 硬件，未對齊的數(shù)據(jù)訪問不是一個大問題。然而，不管 CUDA 硬件是如何產(chǎn)生的，在全局內存中大步前進都是有問題的，而且在許多情況下似乎是不可避免的，例如在訪問多維數(shù)組中沿第二個和更高維的元素時。但是，在這種情況下，如果我們使用共享內存，就可以合并內存訪問。在我在下一篇文章中向您展示如何避免跨越全局內存之前，首先我需要詳細描述一下共享內存。

共享內存

因為它是片上的，共享內存比本地和全局內存快得多。實際上，共享內存延遲大約比未緩存的全局內存延遲低 100 倍（前提是線程之間沒有內存沖突，我們將在本文后面討論這個問題）。共享內存是按線程塊分配的，因此塊中的所有線程都可以訪問同一共享內存。線程可以訪問由同一線程塊中的其他線程從全局內存加載的共享內存中的數(shù)據(jù)。此功能（與線程同步結合）有許多用途，例如用戶管理的數(shù)據(jù)緩存、高性能的協(xié)作并行算法（例如并行縮減），以及在不可能實現(xiàn)全局內存合并的情況下促進全局內存合并。

線程同步

在線程之間共享數(shù)據(jù)時，我們需要小心避免爭用情況，因為雖然塊中的線程并行運行 邏輯上 ，但并非所有線程都可以同時執(zhí)行 身體上 。假設兩個線程 A 和 B 分別從全局內存加載一個數(shù)據(jù)元素并將其存儲到共享內存中。然后，線程 A 想從共享內存中讀取 B 的元素，反之亦然。我們假設 A 和 B 是兩個不同翹曲中的線。如果 B 在 A 嘗試讀取它之前還沒有完成它的元素的編寫，我們就有一個競爭條件，它可能導致未定義的行為和錯誤的結果。

為了保證并行線程協(xié)作時的正確結果，必須同步線程。 CUDA 提供了一個簡單的屏障同步原語 __syncthreads() 。一個線程的執(zhí)行只能在其塊中的所有線程都執(zhí)行了 __syncthreads() 之后通過 __syncthreads() 繼續(xù)執(zhí)行。因此，我們可以通過在存儲到共享內存之后和從共享內存加載任何線程之前調用 __syncthreads() 來避免上面描述的競爭條件。需要注意的是，在發(fā)散代碼中調用 __syncthreads() 是未定義的，并且可能導致死鎖，線程塊中的所有線程都必須在同一點調用 __syncthreads() 。

共享內存示例

使用 Clara 變量 D __shared__ 指定說明符在 CUDA C / C ++設備代碼中聲明共享內存。在內核中聲明共享內存有多種方法，這取決于內存量是在編譯時還是在運行時已知的。下面的完整代碼（ 在 GitHub 上提供 ）演示了使用共享內存的各種方法。

#include __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 元素數(shù)組中的數(shù)據(jù)。這兩個內核非常相似，只是在共享內存數(shù)組的聲明方式和內核的調用方式上有所不同。

靜態(tài)共享內存

如果共享內存數(shù)組大小在編譯時已知，就像在 staticReverse 內核中一樣，那么我們可以顯式地聲明一個該大小的數(shù)組，就像我們對數(shù)組 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] 將數(shù)據(jù)從全局內存復制到共享內存，然后在兩行之后使用語句 d[t] = s[tr] 完成反轉。但是在執(zhí)行最后一行之前，每個線程訪問共享內存中由另一個線程寫入的數(shù)據(jù)，請記住，我們需要通過調用 __syncthreads() 來確保所有線程都已完成對共享內存的加載。

在這個例子中使用共享內存的原因是為了在舊的 CUDA 設備（計算能力 1 . 1 或更早版本）上促進全局內存合并。由于全局內存總是通過線性對齊索引 t 訪問，所以讀寫都可以實現(xiàn)最佳的全局內存合并。反向索引 tr 僅用于訪問共享內存，它不具有全局內存的順序訪問限制以獲得最佳性能。共享內存的唯一性能問題是銀行沖突，我們將在后面討論。（請注意，在計算能力為 1 . 2 或更高版本的設備上，內存系統(tǒng)甚至可以將反向索引存儲完全合并到全局內存中。但是這種技術對于其他訪問模式仍然有用，我將在下一篇文章中展示。）

動態(tài)共享內存

本例中的其他三個內核使用動態(tài)分配的共享內存，當編譯時共享內存的數(shù)量未知時，可以使用該內存。在這種情況下，必須使用可選的第三個執(zhí)行配置參數(shù)指定每個線程塊的共享內存分配大?。ㄒ宰止?jié)為單位），如下面的摘錄所示。

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

動態(tài)共享內存內核 dynamicReverse() 使用未大小化的外部數(shù)組語法 extern shared int s[] 聲明共享內存數(shù)組（注意空括號和 extern 說明符的使用）。大小在內核啟動時由第三個執(zhí)行配置參數(shù)隱式確定。內核代碼的其余部分與 staticReverse() 內核相同。

如果在一個內核中需要多個動態(tài)大小的數(shù)組怎么辦？您必須像前面一樣聲明一個 extern 非大小數(shù)組，并使用指向它的指針將其劃分為多個數(shù)組，如下面的摘錄所示。

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

在內核中指定啟動所需的總內存。

myKernel<<>>(...);

共享內存庫沖突

為了實現(xiàn)并發(fā)訪問的高內存帶寬，共享內存被分成大小相等的內存模塊（庫），這些模塊可以同時訪問。因此，任何跨越 b 不同內存組的 n 地址的內存負載或存儲都可以同時進行服務，從而產(chǎn)生的有效帶寬是單個存儲庫帶寬的 b 倍。

但是，如果多個線程的請求地址映射到同一個內存庫，則訪問將被序列化。硬件根據(jù)需要將沖突內存請求拆分為多個獨立的無沖突請求，將有效帶寬減少一個與沖突內存請求數(shù)量相等的因子。一個例外情況是，一個 warp 中的所有線程都使用同一個共享內存地址，從而導致廣播。計算能力 2 . 0 及更高版本的設備具有多播共享內存訪問的額外能力，這意味著在一個 warp 中通過任意數(shù)量的線程對同一個位置的多個訪問同時進行。

為了最小化內存沖突，了解內存地址如何映射到內存庫是很重要的。共享存儲庫被組織成這樣，連續(xù)的 32 位字被分配給連續(xù)的存儲庫，帶寬是每個庫每個時鐘周期 32 位。對于計算能力為 1 . x 的設備， warp 大小為 32 個線程，庫的數(shù)量為 16 個。一個 warp 的共享內存請求被分為一個對 warp 前半部分的請求和一個對 warp 后半部分的請求。請注意，如果每個內存庫只有一個內存位置被半個線程訪問，則不會發(fā)生庫沖突。

對于計算能力為 2 . 0 的設備， warp 大小是 32 個線程，而 bank 的數(shù)量也是 32 個。 warp 的共享內存請求不會像計算能力為 1 . x 的設備那樣被拆分，這意味著 warp 前半部分的線程和同一 warp 后半部分的線程之間可能會發(fā)生庫沖突。

計算能力為 3 . x 的設備具有可配置的存儲大小，可以使用 CUDA Devicsetsharedmeconfig() 將其設置為四個字節(jié)（ CUDA SharedMemBankSizeFourByte ，默認值）或八個字節(jié)（ cudaSharedMemBankSizeEightByte) 。將存儲大小設置為 8 字節(jié)有助于避免訪問雙精度數(shù)據(jù)時的共享內存庫沖突。

配置共享內存量

在計算能力為 2 . x 和 3 . x 的設備上，每個多處理器都有 64KB 的片上內存，可以在一級緩存和共享內存之間進行分區(qū)。對于計算能力為 2 . x 的設備，有兩個設置： 48KB 共享內存/ 16KB 一級緩存和 16KB 共享內存/ 48KB 一級緩存。默認情況下，使用 48KB 共享內存設置。這可以在運行時 API 期間使用 cudaDeviceSetCacheConfig() 為所有內核配置，也可以使用 cudaFuncSetCacheConfig() 在每個內核的基礎上進行配置。它們接受以下三個選項之一： cudaFuncCachePreferNone 、 cudaFuncCachePreferShared 和 cudaFuncCachePreferL1 。驅動程序將遵循指定的首選項，除非內核每個線程塊需要比指定配置中可用的共享內存更多的共享內存。計算能力為 3 . x 的設備允許使用選項 cudaFuncCachePreferEqual 獲得 32KB 共享內存/ 32kbl1 緩存的第三個設置。

關于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過 20 年的 GPUs 軟件開發(fā)經(jīng)驗，從圖形和游戲到基于物理的模擬，到并行算法和高性能計算。當他還是北卡羅來納大學的博士生時，他意識到了一種新生的趨勢，并為此創(chuàng)造了一個名字： GPGPU （圖形處理單元上的通用計算）。

審核編輯：郭婷