NVIDIA GPU 具有強大的計算能力，通常必須以高速傳輸數(shù)據(jù)才能部署這種能力。原則上這是可能的，因為 GPU 也有很高的內(nèi)存帶寬，但有時他們需要你的幫助來飽和帶寬。

在本文中，我們將研究一種實現(xiàn)這一點的特定方法：預(yù)取。我們將解釋在什么情況下預(yù)取可以很好地工作，以及如何找出這些情況是否適用于您的工作負(fù)載。

上下文

NVIDIA GPU 從大規(guī)模并行中獲得力量。 32 個線程的許多扭曲可以放置在流式多處理器（ SM ）上，等待輪到它們執(zhí)行。當(dāng)一個 warp 因任何原因暫停時， warp 調(diào)度程序會以零開銷切換到另一個，確保 SM 始終有工作要做。

在高性能的 NVIDIA Ampere 架構(gòu) A100 GPU 上，多達(dá) 64 個活動翹板可以共享一個 SM ，每個都有自己的資源。除此之外， A100 還有 108 條短信，可以同時執(zhí)行 warp 指令。

大多數(shù)指令都必須對數(shù)據(jù)進(jìn)行操作，而這些數(shù)據(jù)幾乎總是源于連接到 GPU 的設(shè)備內(nèi)存（ DRAM ）。 SM 上大量的翹曲都可能無法工作的一個主要原因是，它們正在等待來自內(nèi)存的數(shù)據(jù)。

如果出現(xiàn)這種情況，并且內(nèi)存帶寬沒有得到充分利用，則可以重新組織程序，以改善內(nèi)存訪問并減少扭曲暫停，從而使程序更快地完成。這叫做延遲隱藏。

預(yù)取

CPU 上的硬件通常支持的一種技術(shù)稱為預(yù)取。 CPU 看到來自內(nèi)存的請求流到達(dá)，找出模式，并在實際需要數(shù)據(jù)之前開始獲取數(shù)據(jù)。當(dāng)數(shù)據(jù)傳輸?shù)?CPU 的執(zhí)行單元時，可以執(zhí)行其他指令，有效地隱藏傳輸成本（內(nèi)存延遲）。

預(yù)取是一種有用的技術(shù)，但就芯片上的硅面積而言很昂貴。相對而言， GPU 的這些成本甚至更高，因為 GPU 的執(zhí)行單元比 CPU 多得多。相反， GPU 使用多余的扭曲來隱藏內(nèi)存延遲。當(dāng)這還不夠時，可以在軟件中使用預(yù)取。它遵循與硬件支持的預(yù)取相同的原理，但需要明確的指令來獲取數(shù)據(jù)。

要確定此技術(shù)是否能幫助您的程序更快地運行，請使用 GPU 評測工具（如 NVIDIA Nsight Compute ）檢查以下內(nèi)容：

確認(rèn)沒有使用所有內(nèi)存帶寬。

確認(rèn)翹曲被阻止的主要原因是 攤位長記分牌 ，這意味著 SMs 正在等待來自 DRAM 的數(shù)據(jù)。

確認(rèn)這些暫停集中在迭代互不依賴的大型循環(huán)中。

展開

考慮這種循環(huán)的最簡單可能的優(yōu)化，稱為展開。如果循環(huán)足夠短，可以告訴編譯器完全展開循環(huán)，并顯式展開迭代。由于迭代是獨立的，編譯器可以提前發(fā)出所有數(shù)據(jù)請求（“加載”），前提是它為每個加載分配不同的寄存器。

這些請求可以相互重疊，因此整個負(fù)載集只經(jīng)歷一個內(nèi)存延遲，而不是所有單個延遲的總和。更妙的是，加載指令本身的連續(xù)性隱藏了單個延遲的一部分。這是一種接近最優(yōu)的情況，但可能需要大量寄存器才能接收加載結(jié)果。

如果循環(huán)太長，可能會部分展開。在這種情況下，成批的迭代會被擴展，然后您會遵循與之前相同的一般策略。你的工作很少（但你可能沒那么幸運）。

如果循環(huán)包含許多其他指令，這些指令的操作數(shù)需要存儲在寄存器中，那么即使只是部分展開也可能不是一個選項。在這種情況下，在您確認(rèn)滿足之前的條件后，您必須根據(jù)進(jìn)一步的信息做出一些決定。

預(yù)取意味著使數(shù)據(jù)更接近 SMs 的執(zhí)行單元。寄存器是最接近的。如果有足夠的可用空間（可以使用 Nsight Compute Occupation 視圖找到），可以直接預(yù)取到寄存器中。

考慮下面的循環(huán)，其中數(shù)組arr被存儲在全局存儲器（ DRAM ）中。它隱式地假設(shè)只使用了一個一維線程塊，而對于從中派生的激勵應(yīng)用程序來說，情況并非如此。然而，它減少了代碼混亂，并且不會改變參數(shù)。

在本文的所有代碼示例中，大寫變量都是編譯時常量。BLOCKDIMX假定預(yù)定義變量blockDim.x的值。出于某些目的，它必須是編譯時已知的常數(shù)，而出于其他目的，它有助于避免在運行時進(jìn)行計算。

for (i=threadIdx.x; i
};>

假設(shè)您有八個寄存器用于預(yù)取。這是一個調(diào)整參數(shù)。下面的代碼在每四次迭代開始時獲取四個雙精度值，占據(jù)八個 4 字節(jié)寄存器，并逐個使用它們，直到批耗盡，此時您將獲取一個新批。

為了跟蹤批處理，引入一個計數(shù)器（ctr），該計數(shù)器隨著線程執(zhí)行的每個后續(xù)迭代而遞增。為了方便起見，假設(shè)每個線程的迭代次數(shù)可以被 4 整除。

double v0, v1, v2, v3;
for (i=threadIdx.x, ctr=0; i
};>

通常，預(yù)取的值越多，該方法就越有效。雖然前面的例子并不復(fù)雜，但有點麻煩。如果預(yù)取值（PDIST或預(yù)取距離）的數(shù)量發(fā)生變化，則必須添加或刪除代碼行。

將預(yù)取值存儲在共享內(nèi)存中更容易，因為您可以使用數(shù)組表示法，無需任何努力就可以改變預(yù)取距離。然而，共享內(nèi)存并不像寄存器那樣接近執(zhí)行單元。當(dāng)數(shù)據(jù)準(zhǔn)備好使用時，它需要一條額外的指令將數(shù)據(jù)從那里移動到寄存器中。為了方便起見，我們引入宏vsmem來簡化共享內(nèi)存中數(shù)組的索引：

#define vsmem(index) v[index+PDIST*threadIdx.x]
__shared__ double v[PDIST* BLOCKDIMX];
for (i=threadIdx.x, ctr=0; i
};>

除了批量預(yù)取，還可以進(jìn)行“滾動”預(yù)取。在這種情況下，在進(jìn)入主循環(huán)之前填充預(yù)取緩沖區(qū)，然后在每次循環(huán)迭代期間從內(nèi)存中預(yù)取一個值，以便在以后的PDIST迭代中使用。下一個示例使用數(shù)組表示法和共享內(nèi)存實現(xiàn)滾動預(yù)取。

__shared__ double v[PDIST* BLOCKDIMX];
for (k=0; k
};>

與批處理方法相反，滾動預(yù)取在主循環(huán)執(zhí)行期間不會再出現(xiàn)足夠大的預(yù)取距離的內(nèi)存延遲。它還使用相同數(shù)量的共享內(nèi)存或寄存器資源，因此它似乎是首選。然而，一個微妙的問題可能會限制其有效性。

循環(huán)中的同步（例如，syncthreads）構(gòu)成了一個內(nèi)存圍欄，并迫使arr的加載在同一迭代中的該點完成，而不是在以后的 PDIST 迭代中完成。解決方法是使用異步加載到共享內(nèi)存中，最簡單的版本在 CUDA 程序員指南的 Pipeline interface 部分中解釋。這些異步加載不需要在同步點完成，只需要在顯式等待時完成。

以下是相應(yīng)的代碼：

#include 
__shared__ double v[PDIST* BLOCKDIMX];
for (k=0; k
};>

由于每一條__pipeline_wait_prior指令都必須與一條__pipeline_commit指令匹配，我們在進(jìn)入主計算循環(huán)之前，將后者放入預(yù)取緩沖區(qū)的循環(huán)中，以簡化匹配指令對的簿記。

績效結(jié)果

圖 1 顯示，對于不同的預(yù)取距離，在前面描述的五種算法變化下，從金融應(yīng)用程序中獲取的內(nèi)核的性能改進(jìn)。

分批預(yù)取到寄存器（標(biāo)量分批）

分批預(yù)取到共享內(nèi)存（ smem 分批）

將預(yù)取滾動到寄存器（標(biāo)量滾動）

將預(yù)取滾動到共享內(nèi)存（ smem 滾動）

使用異步內(nèi)存拷貝將預(yù)取滾動到共享內(nèi)存（ smem 滾動異步）

圖 1 。不同預(yù)取策略的內(nèi)核加速

顯然，將預(yù)取滾動到具有異步內(nèi)存拷貝的共享內(nèi)存中會帶來很好的好處，但隨著預(yù)取緩沖區(qū)大小的增加，這是不均勻的。

使用 Nsight Compute 對結(jié)果進(jìn)行更仔細(xì)的檢查后發(fā)現(xiàn)，共享內(nèi)存中會發(fā)生內(nèi)存組沖突，這會導(dǎo)致異步負(fù)載的扭曲被拆分為比嚴(yán)格必要的更連續(xù)的內(nèi)存請求。經(jīng)典的優(yōu)化方法是在共享內(nèi)存中填充數(shù)組大小，以避免錯誤的跨步，這種方法在這種情況下有效。PADDING的值的選擇應(yīng)確保PDIST和PADDING之和等于二加一的冪。將其應(yīng)用于所有使用共享內(nèi)存的變體：

#define vsmem(index) v[index+(PDIST+PADDING)*threadIdx.x]

這導(dǎo)致圖 2 所示的共享內(nèi)存結(jié)果得到改善。預(yù)取距離僅為 6 ，再加上以滾動方式進(jìn)行的異步內(nèi)存拷貝，就足以以比原始版本代碼近 60% 的加速比獲得最佳性能。實際上，我們可以通過更改共享內(nèi)存中數(shù)組的索引方案來實現(xiàn)這種性能改進(jìn)，而無需使用填充，這是留給讀者的練習(xí)。

圖 2 。使用共享內(nèi)存填充的不同預(yù)取策略的內(nèi)核加速

一個尚未討論的 預(yù)取的變化 將數(shù)據(jù)從全局內(nèi)存移動到二級緩存，如果共享內(nèi)存中的空間太小，無法容納所有符合預(yù)取條件的數(shù)據(jù)，這可能很有用。這種類型的預(yù)取在 CUDA 中無法直接訪問，需要在較低的 PTX 級別進(jìn)行編程。

總結(jié)

在本文中，我們向您展示了源代碼的本地化更改示例，這些更改可能會加快內(nèi)存訪問。這些不會改變從內(nèi)存移動到 SMs 的數(shù)據(jù)量，只會改變時間。通過重新安排內(nèi)存訪問，使數(shù)據(jù)在到達(dá) SM 后被多次重用，您可以進(jìn)行更多優(yōu)化。

關(guān)于作者

Rob Van der Wijngaart 是 NVIDIA 的高級高性能計算（ HPC ）架構(gòu)師。他在各種工業(yè)和政府實驗室從事 HPC 領(lǐng)域的研究超過三十年，是廣泛使用的 NAS 并行基準(zhǔn)測試的共同開發(fā)者。Ren é Peters 是 NVIDIA 的產(chǎn)品經(jīng)理，他在增強/虛擬現(xiàn)實和人工智能的交叉點指導(dǎo)產(chǎn)品開發(fā)。在科技行業(yè)任職期間，他還與物聯(lián)網(wǎng)（ IoT ）和云計算等技術(shù)合作。邁爾斯·麥克林（ Miles Macklin ）是NVIDIA 的首席工程師，致力于模擬技術(shù)。他從哥本哈根大學(xué)獲得計算機科學(xué)博士學(xué)位，從事計算機圖形學(xué)、基于物理學(xué)的動畫和機器人學(xué)的研究。他在 ACM SIGGRAPH 期刊上發(fā)表了幾篇論文，他的研究已經(jīng)被整合到許多商業(yè)產(chǎn)品中，包括NVIDIA 的 PhysX 和 ISAAC 健身房模擬器。他最近的工作旨在為 GPU 上的可微編程開發(fā)健壯高效的框架。

Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級產(chǎn)品營銷經(jīng)理。弗雷德?lián)碛屑又荽髮W(xué)戴維斯分校計算機科學(xué)和數(shù)學(xué)學(xué)士學(xué)位。他的職業(yè)生涯開始于一名 UNIX 軟件工程師，負(fù)責(zé)將內(nèi)核服務(wù)和設(shè)備驅(qū)動程序移植到 x86 體系結(jié)構(gòu)。他喜歡《星球大戰(zhàn)》、《星際迷航》和 NBA 勇士隊。

審核編輯：郭婷