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

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

共享內(nèi)存

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

線程同步

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

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

共享內(nèi)存示例

使用 Clara 變量 D __shared__ 指定說明符在 CUDA C / C ++設(shè)備代碼中聲明共享內(nèi)存。在內(nèi)核中聲明共享內(nèi)存有多種方法，這取決于內(nèi)存量是在編譯時(shí)還是在運(yùn)行時(shí)已知的。下面的完整代碼（ 在 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 元素?cái)?shù)組中的數(shù)據(jù)。這兩個(gè)內(nèi)核非常相似，只是在共享內(nèi)存數(shù)組的聲明方式和內(nèi)核的調(diào)用方式上有所不同。

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

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

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

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

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

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

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

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

如果在一個(gè)內(nèi)核中需要多個(gè)動(dòng)態(tài)大小的數(shù)組怎么辦？您必須像前面一樣聲明一個(gè) extern 非大小數(shù)組，并使用指向它的指針將其劃分為多個(gè)數(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)核中指定啟動(dòng)所需的總內(nèi)存。

myKernel<<>>(...);

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

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

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

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

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

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

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

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

關(guān)于作者

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

審核編輯：郭婷