NVIDIA GPU 具有強(qiáng)大的計(jì)算能力,通常需要高速傳輸數(shù)據(jù)才能部署這種能力。原則上,這是可能的,因?yàn)?GPU 也有很高的內(nèi)存帶寬,但有時(shí)他們需要程序員的幫助來飽和帶寬。在這篇博文中,我們研究了一種實(shí)現(xiàn)這一點(diǎn)的方法,并將其應(yīng)用于金融計(jì)算中的一個(gè)示例。我們將解釋在什么情況下這種方法可以很好地工作,以及如何找出這些情況是否適用于您的工作負(fù)載。
上下文
NVIDIA GPU 的力量來自大規(guī)模并行??梢詫?32 個(gè)線程的許多扭曲放置在流式多處理器( SM )上,等待輪到它們執(zhí)行。當(dāng)一個(gè) warp 因任何原因暫停時(shí), warp 調(diào)度程序?qū)⑶袚Q到另一個(gè),開銷為零,確保 SM 始終有工作要做。在高性能 NVIDIA Ampere 100 ( A100 ) GPU 上,多達(dá) 64 個(gè)活動(dòng)經(jīng)線可以共享一個(gè) SM ,每個(gè)都有自己的資源。除此之外, A100 還有許多 SMs-108 ,它們都可以同時(shí)執(zhí)行 warp 指令。大多數(shù)指令都必須對(duì)數(shù)據(jù)進(jìn)行操作,而這些數(shù)據(jù)幾乎總是源自連接到 GPU 的設(shè)備內(nèi)存( DRAM )。 SM 上大量的翹曲也可能無法工作的一個(gè)主要原因是,它們正在等待來自內(nèi)存的數(shù)據(jù)。如果發(fā)生這種情況,并且內(nèi)存帶寬沒有得到充分利用,則可以重新組織程序以改進(jìn)內(nèi)存訪問并減少扭曲暫停,從而使程序更快完成。
第一步:寬負(fù)載
在之前的博客文章中,我們檢查了一個(gè)工作負(fù)載,該工作負(fù)載沒有充分利用 GPU 的可用計(jì)算和內(nèi)存帶寬資源。我們確定,在需要之前從內(nèi)存中預(yù)取數(shù)據(jù)可以大大減少內(nèi)存暫停并提高性能。當(dāng)預(yù)取不適用時(shí),需要確定哪些其他因素可能會(huì)限制內(nèi)存子系統(tǒng)的性能。一種可能性是,向該子系統(tǒng)發(fā)出請(qǐng)求的速率太高。直觀地說,我們可以通過在每個(gè)加載指令中提取多個(gè)單詞來降低請(qǐng)求速率。最好用一個(gè)例子來說明這一點(diǎn)。
在本文的所有代碼示例中,大寫變量都是編譯時(shí)常量。 BLOCKDIMX 采用預(yù)定義變量 blockDim 的值。 x 、 出于某些目的,它必須是編譯時(shí)已知的常量,而出于其他目的,它有助于避免在運(yùn)行時(shí)進(jìn)行計(jì)算。
原始代碼如下所示,index是計(jì)算數(shù)組索引的輔助函數(shù)。它隱式地假設(shè)只使用了一個(gè)一維線程塊,而派生它的激勵(lì)應(yīng)用程序則不是這樣。但是,它減少了代碼混亂,并且不會(huì)更改參數(shù)。
for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; ++k) { double c = big_array[index(pt, k)]; c += small_array[k] ; best = max(c, best); } final[pt] = best;
}
請(qǐng)注意,每個(gè)線程從建議命名的small_array中加載kmax個(gè)連續(xù)值。此陣列足夠小,完全適合一級(jí)緩存,但要求它以非常高的速率返回?cái)?shù)據(jù)可能會(huì)出現(xiàn)問題。下面的更改表明,如果我們稍微重新構(gòu)造代碼并引入 double2 數(shù)據(jù)類型,則每個(gè)線程可以在同一條指令中發(fā)出兩個(gè)雙精度字的請(qǐng)求,這在 NVIDIA GPU 上本機(jī)支持;它將兩個(gè)雙精度字存儲(chǔ)在相鄰的內(nèi)存位置,可以使用字段選擇器“ x ”和“ y ”訪問這些位置。之所以這樣做,是因?yàn)槊總€(gè)線程都訪問small_array的連續(xù)元素。我們稱這種技術(shù)為 VZX28 。請(qǐng)注意,索引“k”上的內(nèi)部循環(huán)現(xiàn)在增加了 2 ,而不是 1 。
for (pt = threadIdx.x; pt < ptmax ; pt += BLOCKDIMX ) { double best = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double c = big_array[index(pt, k)]; double2 val = *(double2 *) &small_array[k]; c += val.x; best = max(c, best); c = big_array[index(pt, k+1)]; c += val.y; best = max(c, best); } final[pt] = best;
}
有幾個(gè)注意事項(xiàng)。首先,我們沒有檢查kmax是否為偶數(shù)。如果沒有,修改后的k循環(huán)將執(zhí)行額外的迭代,我們需要編寫一些特殊代碼來防止這種情況發(fā)生。其次,我們沒有確認(rèn)small_array是否在 16 字節(jié)邊界上正確對(duì)齊。否則,寬荷載將失效。如果它是使用cudaMalloc分配的,它將自動(dòng)在 256 字節(jié)的邊界上對(duì)齊。但是,如果使用指針算法將其傳遞給內(nèi)核,則需要執(zhí)行一些檢查。
接下來,我們檢查輔助函數(shù)指數(shù),發(fā)現(xiàn)它在 pt 中與系數(shù) 1 呈線性關(guān)系。因此,通過在一條指令中請(qǐng)求兩個(gè)雙精度值,我們可以對(duì)從 big \ U 數(shù)組獲取的值應(yīng)用類似的寬負(fù)載方法。對(duì)big_array和small_array的訪問之間的區(qū)別在于,現(xiàn)在 warp 中的連續(xù)線程訪問相鄰的數(shù)組元素。下面重構(gòu)的代碼將數(shù)組元素上的循環(huán)增量加倍big_array,現(xiàn)在每個(gè)線程在每次迭代中處理兩個(gè)數(shù)組元素。
for (pt = 2*threadIdx.x; pt < ptmax ; pt += 2*BLOCKDIMX ) { double best1 = 0.0, best2 = 0.0; #pragma unroll for (int k = 0; k < kmax; k+=2) { double2 c1 = *(double2 *) &big_array[index(pt, k)]; double2 c2 = *(double2 *) &big_array[index(pt, k+1)]; double2 val = *(double2 *) &small_array[k]; c1.x += val.x; best1 = max(c1.x, best1); c2.x += val.y; best1 = max(c2.x, best1); c1.y += val.x; best2 = max(c1.y, best2); c2.y += val.y; best2 = max(c2.y, best2); } final[pt] = best1; final[pt+1] = best2;
}
與之前相同的注意事項(xiàng)也適用,現(xiàn)在應(yīng)該擴(kuò)展到ptmax的奇偶校驗(yàn)和big_array的對(duì)齊。幸運(yùn)的是,從中派生此示例的應(yīng)用程序滿足所有要求。下圖顯示了在應(yīng)用程序中重復(fù)多次的一組內(nèi)核的持續(xù)時(shí)間(以納秒為單位)。對(duì)于寬負(fù)載組合,內(nèi)核的平均加速比為 1.63 倍。

圖 1 :由于負(fù)載較寬,內(nèi)核持續(xù)時(shí)間減少
第二步:寄存器使用
我們可能想到此為止并宣布成功,但使用 NVIDIA Nsight Compute 對(duì)程序執(zhí)行的深入分析表明,即使我們將加載指令的數(shù)量減少了一半,我們也沒有從根本上改變對(duì)內(nèi)存子系統(tǒng)的請(qǐng)求速率。原因是一條扭曲加載指令(即 32 個(gè)線程同時(shí)發(fā)出加載指令)會(huì)導(dǎo)致一個(gè)或多個(gè)扇區(qū)請(qǐng)求,這是硬件處理的實(shí)際內(nèi)存訪問單元。每個(gè)扇區(qū)是 32 字節(jié),因此每個(gè)線程一條 8 字節(jié)雙精度字的扭曲加載指令會(huì)導(dǎo)致 8 個(gè)扇區(qū)請(qǐng)求(訪問以單位跨距進(jìn)行),而一條雙精度字的扭曲加載指令會(huì)導(dǎo)致 16 個(gè)扇區(qū)請(qǐng)求。普通負(fù)載和寬負(fù)載的扇區(qū)請(qǐng)求總數(shù)相同。那么,是什么導(dǎo)致了性能的提高呢?
為了理解代碼行為,我們需要考慮一個(gè)尚未討論的資源,即寄存器。這些用于存儲(chǔ)從內(nèi)存加載的數(shù)據(jù),并用作算術(shù)指令的輸入。寄存器是一種有限的資源。如果流式多處理器( SM )在 A100 GPU 上承載盡可能多的扭曲,則每個(gè)線程可以使用 32 個(gè) 4 字節(jié)寄存器,這些寄存器總共可以容納 16 個(gè)雙精度字。將代碼翻譯成機(jī)器語言的編譯器知道這一點(diǎn),并將限制每個(gè)線程的寄存器數(shù)量。我們?nèi)绾未_定代碼的寄存器使用及其在性能中所起的作用?我們使用 Nsight Compute 中的“ source ”視圖來并排查看匯編代碼(“ SASS ”)和 C 源代碼。
代碼的最內(nèi)層循環(huán)是執(zhí)行次數(shù)最多的循環(huán),因此,如果我們?cè)趯?dǎo)航菜單中選擇“已執(zhí)行的指令”,然后要求轉(zhuǎn)到 SASS 代碼中數(shù)量最多的那一行,我們會(huì)自動(dòng)進(jìn)入內(nèi)部循環(huán)。如果不確定,可以將 SASS 與突出顯示的相應(yīng)源代碼進(jìn)行比較以確認(rèn)。接下來,我們?cè)趦?nèi)環(huán)的 SASS 代碼中識(shí)別從內(nèi)存( LDG )加載數(shù)據(jù)的所有指令。圖 2 顯示了 SASS 的一個(gè)片段,我們?cè)谄渲兴阉饕哉业絻?nèi)部循環(huán)的開始;在第 166 行,指令的執(zhí)行次數(shù)突然跳到其最大值。

圖 2 :演示內(nèi)部循環(huán)開始的 SASS 代碼段(第 166 行)
LDG 。 E 、 64 是我們所追求的指令。它從全局內(nèi)存( DRAM )加載一個(gè)具有擴(kuò)展地址的 64 位字。寬單詞的負(fù)載對(duì)應(yīng)于 LDG 。 E 、 128 。加載指令名稱后的第一個(gè)參數(shù)(圖 2 中的 R34 )是接收該值的寄存器。由于雙精度值占用兩個(gè)相鄰寄存器,因此加載指令中隱含 R35 。接下來,我們比較三個(gè)版本的代碼( 1.基線, 2.寬負(fù)載的small_array, 3.寬負(fù)載的small_array和big_array)在內(nèi)部循環(huán)中使用寄存器的方式。回想一下,編譯器試圖保持在限制范圍內(nèi),有時(shí)需要對(duì)寄存器進(jìn)行處理。也就是說,如果沒有足夠的寄存器可用于從內(nèi)存接收每個(gè)唯一值,它將重用以前在內(nèi)部循環(huán)中使用的寄存器。
這樣做的結(jié)果是,算術(shù)指令需要使用以前的值,以便新值可以覆蓋它。此時(shí),從內(nèi)存加載需要等待該指令完成:內(nèi)存延遲暴露。在所有現(xiàn)代計(jì)算機(jī)體系結(jié)構(gòu)上,此延遲構(gòu)成了一個(gè)顯著的延遲。在 GPU 上,可以通過切換到另一個(gè)扭曲來隱藏部分扭曲,但通常不是全部扭曲。因此,寄存器在內(nèi)環(huán)中被重用的次數(shù)可以表示代碼的速度變慢。
有了這一見解,我們分析了代碼的三個(gè)版本,發(fā)現(xiàn)它們?cè)诿總€(gè)內(nèi)部循環(huán)中分別經(jīng)歷了 8 、 6 和 3 個(gè)內(nèi)存延遲,這解釋了圖 1 所示的性能差異。不同寄存器重用模式背后的主要原因是,當(dāng)兩個(gè)普通加載融合為單個(gè)寬加載時(shí),通常需要更少的地址計(jì)算,并且地址計(jì)算的結(jié)果也會(huì)進(jìn)入寄存器。隨著持有地址的寄存器越來越多,剩下來充當(dāng)從內(nèi)存中提取的值的“著陸區(qū)”的地址越來越少,我們?cè)?Music chairs 游戲中失去了席位;寄存器壓力增大。
第三步:?jiǎn)?dòng)邊界
我們還沒有完成?,F(xiàn)在我們知道了寄存器在程序性能中所起的關(guān)鍵作用,我們將查看三個(gè)版本的代碼使用的寄存器總數(shù)。最簡(jiǎn)單的方法是再次檢查 Nsight Compute 報(bào)告。我們發(fā)現(xiàn)使用的寄存器數(shù)量分別為 40 、 36 和 44 。
編譯器確定這些數(shù)字的方法是使用復(fù)雜的啟發(fā)式算法,該算法考慮了大量因素,包括 SM 上可能存在多少活動(dòng)扭曲、在忙循環(huán)中加載的唯一值的數(shù)量以及每個(gè)操作所需的寄存器數(shù)量。如果編譯器不知道 SM 上可能存在的扭曲數(shù),它將嘗試將每個(gè)線程的寄存器數(shù)限制為 32 ,因?yàn)槿绻嬖谟布试S的絕對(duì)最大同時(shí)扭曲數(shù)( 64 ),那么這就是可用的數(shù)字。在我們的例子中,我們沒有告訴編譯器期望的是什么,所以它盡了最大努力,但顯然確定僅使用 32 個(gè)寄存器生成的代碼效率太低。
然而,內(nèi)核的 launch 語句中指定的線程塊的實(shí)際大小是 1024 個(gè)線程,因此有 32 個(gè)扭曲。這意味著,如果 SM 上只存在一個(gè)線程塊,則每個(gè)線程最多可以使用 64 個(gè)線程。在實(shí)際使用的每個(gè)線程中有 40 、 36 和 44 個(gè)寄存器時(shí),沒有足夠的寄存器可用于支持每個(gè) SM 的兩個(gè)或多個(gè)線程塊,因此將只啟動(dòng)一個(gè),每個(gè)線程分別保留 24 、 28 和 20 個(gè)未使用的寄存器。
通過使用 launch bounds 將我們的意圖告知編譯器,我們可以做得更好。通過告訴編譯器一個(gè)線程塊中的最大線程數(shù)( 1024 )和同時(shí)支持的最小塊數(shù)( 1 ),編譯器可以放松,并且很高興每個(gè)線程分別使用 63 、 56 和 64 個(gè)寄存器。
有趣的是,最快的代碼版本現(xiàn)在是基線版本,沒有任何廣泛的負(fù)載。雖然組合寬負(fù)載 without 啟動(dòng)邊界的加速比為 1.64 倍,但寬負(fù)載 with 啟動(dòng)邊界的加速比為 1.76 倍,而基線代碼的加速比為 1.77 倍。這意味著我們不必費(fèi)心修改內(nèi)核定義;在這種情況下,僅提供啟動(dòng)邊界就足以獲得這種特定線程塊大小的最佳性能。
通過對(duì) SM 上的線程塊大小和預(yù)期的最小線程塊數(shù)進(jìn)行更多的實(shí)驗(yàn),我們?cè)诿總€(gè) SM 有 512 個(gè)線程的 2 個(gè)線程塊的情況下達(dá)到了 1.79 倍的加速,對(duì)于沒有寬負(fù)載的基線版本也是如此。
結(jié)論
寄存器的有效使用對(duì)于獲得良好的 GPU 內(nèi)核性能至關(guān)重要。有時(shí),一種稱為“寬負(fù)載”的技術(shù)可以帶來顯著的好處。它減少了計(jì)算并需要存儲(chǔ)在寄存器中的內(nèi)存地址的數(shù)量,留下更多的寄存器來接收來自內(nèi)存的數(shù)據(jù)。然而,向編譯器提示在應(yīng)用程序中啟動(dòng)內(nèi)核的方式可能會(huì)帶來同樣的好處,而無需更改內(nèi)核本身。
關(guān)于作者
Rob Van der Wijngaart 是 NVIDIA 的高級(jí)高性能計(jì)算( HPC )架構(gòu)師。他在各種工業(yè)和政府實(shí)驗(yàn)室從事 HPC 領(lǐng)域的研究超過三十年,是廣泛使用的 NAS 并行基準(zhǔn)測(cè)試的共同開發(fā)者。
Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級(jí)產(chǎn)品營(yíng)銷經(jīng)理。弗雷德?lián)碛屑又荽髮W(xué)戴維斯分校計(jì)算機(jī)科學(xué)和數(shù)學(xué)學(xué)士學(xué)位。他的職業(yè)生涯開始于一名 UNIX 軟件工程師,負(fù)責(zé)將內(nèi)核服務(wù)和設(shè)備驅(qū)動(dòng)程序移植到 x86 體系結(jié)構(gòu)。
審核編輯:郭婷
-
處理器
+關(guān)注
關(guān)注
68文章
20264瀏覽量
252671 -
NVIDIA
+關(guān)注
關(guān)注
14文章
5602瀏覽量
109832 -
gpu
+關(guān)注
關(guān)注
28文章
5204瀏覽量
135547
發(fā)布評(píng)論請(qǐng)先 登錄
C語言訪問某特定內(nèi)存位置
內(nèi)存與數(shù)據(jù)處理優(yōu)化藝術(shù)
通過sysmem接口擴(kuò)展內(nèi)存空間
提高RISC-V在Drystone測(cè)試中得分的方法
蜂鳥E203內(nèi)核優(yōu)化方法
學(xué)生適合使用的SOLIDWORKS 云應(yīng)用程序
樹莓派5超頻指南:安全高效地提升性能!
Linux系統(tǒng)性能指南
【「算力芯片 | 高性能 CPU/GPU/NPU 微架構(gòu)分析」閱讀體驗(yàn)】+NVlink技術(shù)從應(yīng)用到原理
HarmonyOS優(yōu)化應(yīng)用內(nèi)存占用問題性能優(yōu)化四
HarmonyOS優(yōu)化應(yīng)用內(nèi)存占用問題性能優(yōu)化一
如何使用USB中斷傳輸方法訪問FPGA?
如何使用CYUSB3KIT-003使用GPIO訪問SRAM的應(yīng)用程序?
可以手動(dòng)構(gòu)建imx-gpu-viv嗎?
MCP:連接AI與應(yīng)用程序的開放標(biāo)準(zhǔn)!
通過GPU內(nèi)存訪問調(diào)整提高應(yīng)用程序性能
評(píng)論