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

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

共享內(nèi)存

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

線程同步

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

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

共享內(nèi)存示例

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

#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]);?

}此代碼使用共享內(nèi)存反轉(zhuǎn) 64 元素數(shù)組中的數(shù)據(jù)。這兩個內(nèi)核非常相似，只是在共享內(nèi)存數(shù)組的聲明方式和內(nèi)核的調(diào)用方式上有所不同。

靜態(tài)共享內(nèi)存

如果共享內(nèi)存數(shù)組大小在編譯時已知，就像在 staticReverse 內(nèi)核中一樣，那么我們可以顯式地聲明一個該大小的數(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];

}在這個內(nèi)核中， t 和 tr 是分別表示原始順序和反向順序的兩個索引。線程使用語句 s[t] = d[t] 將數(shù)據(jù)從全局內(nèi)存復(fù)制到共享內(nèi)存，然后在兩行之后使用語句 d[t] = s[tr] 完成反轉(zhuǎn)。但是在執(zhí)行最后一行之前，每個線程訪問共享內(nèi)存中由另一個線程寫入的數(shù)據(jù)，請記住，我們需要通過調(diào)用 __syncthreads() 來確保所有線程都已完成對共享內(nèi)存的加載。

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

動態(tài)共享內(nèi)存

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

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

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

如果在一個內(nèi)核中需要多個動態(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

在內(nèi)核中指定啟動所需的總內(nèi)存。

myKernel<<>>(...);

共享內(nèi)存庫沖突

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

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

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

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

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

配置共享內(nèi)存量

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

關(guān)于作者

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

審核編輯：郭婷