NVIDIA GPU 具有強大的計算能力,通常需要高速傳輸數據才能部署這種能力。原則上,這是可能的,因為 GPU 也有很高的內存帶寬,但有時他們需要程序員的幫助來飽和帶寬。在這篇博文中,我們研究了一種實現這一點的方法,并將其應用于金融計算中的一個示例。我們將解釋在什么情況下這種方法可以很好地工作,以及如何找出這些情況是否適用于您的工作負載。
上下文
NVIDIA GPU 的力量來自大規模并行。可以將 32 個線程的許多扭曲放置在流式多處理器( SM )上,等待輪到它們執行。當一個 warp 因任何原因暫停時, warp 調度程序將切換到另一個,開銷為零,確保 SM 始終有工作要做。在高性能 NVIDIA Ampere 100 ( A100 ) GPU 上,多達 64 個活動經線可以共享一個 SM ,每個都有自己的資源。除此之外, A100 還有許多 SMs-108 ,它們都可以同時執行 warp 指令。大多數指令都必須對數據進行操作,而這些數據幾乎總是源自連接到 GPU 的設備內存( DRAM )。 SM 上大量的翹曲也可能無法工作的一個主要原因是,它們正在等待來自內存的數據。如果發生這種情況,并且內存帶寬沒有得到充分利用,則可以重新組織程序以改進內存訪問并減少扭曲暫停,從而使程序更快完成。
第一步:寬負載
在之前的博客文章中,我們檢查了一個工作負載,該工作負載沒有充分利用 GPU 的可用計算和內存帶寬資源。我們確定,在需要之前從內存中預取數據可以大大減少內存暫停并提高性能。當預取不適用時,需要確定哪些其他因素可能會限制內存子系統的性能。一種可能性是,向該子系統發出請求的速率太高。直觀地說,我們可以通過在每個加載指令中提取多個單詞來降低請求速率。最好用一個例子來說明這一點。
在本文的所有代碼示例中,大寫變量都是編譯時常量。 BLOCKDIMX 采用預定義變量 blockDim 的值。 x 、 出于某些目的,它必須是編譯時已知的常量,而出于其他目的,它有助于避免在運行時進行計算。
原始代碼如下所示,index是計算數組索引的輔助函數。它隱式地假設只使用了一個一維線程塊,而派生它的激勵應用程序則不是這樣。但是,它減少了代碼混亂,并且不會更改參數。
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;
}
請注意,每個線程從建議命名的small_array
中加載kmax
個連續值。此陣列足夠小,完全適合一級緩存,但要求它以非常高的速率返回數據可能會出現問題。下面的更改表明,如果我們稍微重新構造代碼并引入 double2 數據類型,則每個線程可以在同一條指令中發出兩個雙精度字的請求,這在 NVIDIA GPU 上本機支持;它將兩個雙精度字存儲在相鄰的內存位置,可以使用字段選擇器“ x ”和“ y ”訪問這些位置。之所以這樣做,是因為每個線程都訪問small_array
的連續元素。我們稱這種技術為 VZX28 。請注意,索引“k
”上的內部循環現在增加了 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;
}
有幾個注意事項。首先,我們沒有檢查kmax
是否為偶數。如果沒有,修改后的k
循環將執行額外的迭代,我們需要編寫一些特殊代碼來防止這種情況發生。其次,我們沒有確認small_array
是否在 16 字節邊界上正確對齊。否則,寬荷載將失效。如果它是使用cudaMalloc
分配的,它將自動在 256 字節的邊界上對齊。但是,如果使用指針算法將其傳遞給內核,則需要執行一些檢查。
接下來,我們檢查輔助函數指數,發現它在 pt 中與系數 1 呈線性關系。因此,通過在一條指令中請求兩個雙精度值,我們可以對從 big \ U 數組獲取的值應用類似的寬負載方法。對big_array
和small_array
的訪問之間的區別在于,現在 warp 中的連續線程訪問相鄰的數組元素。下面重構的代碼將數組元素上的循環增量加倍big_array
,現在每個線程在每次迭代中處理兩個數組元素。
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;
}
與之前相同的注意事項也適用,現在應該擴展到ptmax
的奇偶校驗和big_array
的對齊。幸運的是,從中派生此示例的應用程序滿足所有要求。下圖顯示了在應用程序中重復多次的一組內核的持續時間(以納秒為單位)。對于寬負載組合,內核的平均加速比為 1.63 倍。
圖 1 :由于負載較寬,內核持續時間減少
第二步:寄存器使用
我們可能想到此為止并宣布成功,但使用 NVIDIA Nsight Compute 對程序執行的深入分析表明,即使我們將加載指令的數量減少了一半,我們也沒有從根本上改變對內存子系統的請求速率。原因是一條扭曲加載指令(即 32 個線程同時發出加載指令)會導致一個或多個扇區請求,這是硬件處理的實際內存訪問單元。每個扇區是 32 字節,因此每個線程一條 8 字節雙精度字的扭曲加載指令會導致 8 個扇區請求(訪問以單位跨距進行),而一條雙精度字的扭曲加載指令會導致 16 個扇區請求。普通負載和寬負載的扇區請求總數相同。那么,是什么導致了性能的提高呢?
為了理解代碼行為,我們需要考慮一個尚未討論的資源,即寄存器。這些用于存儲從內存加載的數據,并用作算術指令的輸入。寄存器是一種有限的資源。如果流式多處理器( SM )在 A100 GPU 上承載盡可能多的扭曲,則每個線程可以使用 32 個 4 字節寄存器,這些寄存器總共可以容納 16 個雙精度字。將代碼翻譯成機器語言的編譯器知道這一點,并將限制每個線程的寄存器數量。我們如何確定代碼的寄存器使用及其在性能中所起的作用?我們使用 Nsight Compute 中的“ source ”視圖來并排查看匯編代碼(“ SASS ”)和 C 源代碼。
代碼的最內層循環是執行次數最多的循環,因此,如果我們在導航菜單中選擇“已執行的指令”,然后要求轉到 SASS 代碼中數量最多的那一行,我們會自動進入內部循環。如果不確定,可以將 SASS 與突出顯示的相應源代碼進行比較以確認。接下來,我們在內環的 SASS 代碼中識別從內存( LDG )加載數據的所有指令。圖 2 顯示了 SASS 的一個片段,我們在其中搜索以找到內部循環的開始;在第 166 行,指令的執行次數突然跳到其最大值。
圖 2 :演示內部循環開始的 SASS 代碼段(第 166 行)
LDG 。 E 、 64 是我們所追求的指令。它從全局內存( DRAM )加載一個具有擴展地址的 64 位字。寬單詞的負載對應于 LDG 。 E 、 128 。加載指令名稱后的第一個參數(圖 2 中的 R34 )是接收該值的寄存器。由于雙精度值占用兩個相鄰寄存器,因此加載指令中隱含 R35 。接下來,我們比較三個版本的代碼( 1.基線, 2.寬負載的small_array, 3.寬負載的small_array和big_array)在內部循環中使用寄存器的方式。回想一下,編譯器試圖保持在限制范圍內,有時需要對寄存器進行處理。也就是說,如果沒有足夠的寄存器可用于從內存接收每個唯一值,它將重用以前在內部循環中使用的寄存器。
這樣做的結果是,算術指令需要使用以前的值,以便新值可以覆蓋它。此時,從內存加載需要等待該指令完成:內存延遲暴露。在所有現代計算機體系結構上,此延遲構成了一個顯著的延遲。在 GPU 上,可以通過切換到另一個扭曲來隱藏部分扭曲,但通常不是全部扭曲。因此,寄存器在內環中被重用的次數可以表示代碼的速度變慢。
有了這一見解,我們分析了代碼的三個版本,發現它們在每個內部循環中分別經歷了 8 、 6 和 3 個內存延遲,這解釋了圖 1 所示的性能差異。不同寄存器重用模式背后的主要原因是,當兩個普通加載融合為單個寬加載時,通常需要更少的地址計算,并且地址計算的結果也會進入寄存器。隨著持有地址的寄存器越來越多,剩下來充當從內存中提取的值的“著陸區”的地址越來越少,我們在 Music chairs 游戲中失去了席位;寄存器壓力增大。
第三步:啟動邊界
我們還沒有完成。現在我們知道了寄存器在程序性能中所起的關鍵作用,我們將查看三個版本的代碼使用的寄存器總數。最簡單的方法是再次檢查 Nsight Compute 報告。我們發現使用的寄存器數量分別為 40 、 36 和 44 。
編譯器確定這些數字的方法是使用復雜的啟發式算法,該算法考慮了大量因素,包括 SM 上可能存在多少活動扭曲、在忙循環中加載的唯一值的數量以及每個操作所需的寄存器數量。如果編譯器不知道 SM 上可能存在的扭曲數,它將嘗試將每個線程的寄存器數限制為 32 ,因為如果存在硬件允許的絕對最大同時扭曲數( 64 ),那么這就是可用的數字。在我們的例子中,我們沒有告訴編譯器期望的是什么,所以它盡了最大努力,但顯然確定僅使用 32 個寄存器生成的代碼效率太低。
然而,內核的 launch 語句中指定的線程塊的實際大小是 1024 個線程,因此有 32 個扭曲。這意味著,如果 SM 上只存在一個線程塊,則每個線程最多可以使用 64 個線程。在實際使用的每個線程中有 40 、 36 和 44 個寄存器時,沒有足夠的寄存器可用于支持每個 SM 的兩個或多個線程塊,因此將只啟動一個,每個線程分別保留 24 、 28 和 20 個未使用的寄存器。
通過使用 launch bounds 將我們的意圖告知編譯器,我們可以做得更好。通過告訴編譯器一個線程塊中的最大線程數( 1024 )和同時支持的最小塊數( 1 ),編譯器可以放松,并且很高興每個線程分別使用 63 、 56 和 64 個寄存器。
有趣的是,最快的代碼版本現在是基線版本,沒有任何廣泛的負載。雖然組合寬負載 without 啟動邊界的加速比為 1.64 倍,但寬負載 with 啟動邊界的加速比為 1.76 倍,而基線代碼的加速比為 1.77 倍。這意味著我們不必費心修改內核定義;在這種情況下,僅提供啟動邊界就足以獲得這種特定線程塊大小的最佳性能。
通過對 SM 上的線程塊大小和預期的最小線程塊數進行更多的實驗,我們在每個 SM 有 512 個線程的 2 個線程塊的情況下達到了 1.79 倍的加速,對于沒有寬負載的基線版本也是如此。
結論
寄存器的有效使用對于獲得良好的 GPU 內核性能至關重要。有時,一種稱為“寬負載”的技術可以帶來顯著的好處。它減少了計算并需要存儲在寄存器中的內存地址的數量,留下更多的寄存器來接收來自內存的數據。然而,向編譯器提示在應用程序中啟動內核的方式可能會帶來同樣的好處,而無需更改內核本身。
關于作者
Rob Van der Wijngaart 是 NVIDIA 的高級高性能計算( HPC )架構師。他在各種工業和政府實驗室從事 HPC 領域的研究超過三十年,是廣泛使用的 NAS 并行基準測試的共同開發者。
Fred Oh 是 CUDA 、 CUDA on WSL 和 CUDA Python 的高級產品營銷經理。弗雷德擁有加州大學戴維斯分校計算機科學和數學學士學位。他的職業生涯開始于一名 UNIX 軟件工程師,負責將內核服務和設備驅動程序移植到 x86 體系結構。
審核編輯:郭婷
-
處理器
+關注
關注
68文章
19354瀏覽量
230391 -
NVIDIA
+關注
關注
14文章
5039瀏覽量
103309 -
gpu
+關注
關注
28文章
4754瀏覽量
129100
發布評論請先 登錄
相關推薦
評論