在线观看www成人影院-在线观看www日本免费网站-在线观看www视频-在线观看操-欧美18在线-欧美1级

0
  • 聊天消息
  • 系統消息
  • 評論與回復
登錄后你可以
  • 下載海量資料
  • 學習在線課程
  • 觀看技術視頻
  • 寫文章/發帖/加入社區
會員中心
創作中心

完善資料讓更多小伙伴認識你,還能領取20積分哦,立即完善>

3天內不再提示

通過GPU內存訪問調整提高應用程序性能

星星科技指導員 ? 來源:NVIDIA ? 作者:NVIDIA ? 2022-08-15 16:24 ? 次閱讀

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_arraysmall_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
    gpu
    +關注

    關注

    28

    文章

    4754

    瀏覽量

    129100
收藏 人收藏

    評論

    相關推薦

    Chart FX-調整設置和功能

    體驗。即雖然它是一個分布式應用程序,可以同時為數百甚至數千個用戶提供服務,但要讓其像傳統的桌面應用程序一樣運行。Chart FX 對服務器性能的影響體現在,每次訪問包含圖表的網頁時,C
    的頭像 發表于 01-08 16:49 ?75次閱讀
    Chart FX-<b class='flag-5'>調整</b>設置和功能

    通過Skyvia Connect SQL終端節點訪問任何數據

    作為網關 ADO.NET 通過最知名和最廣泛使用的 .NET 數據訪問接口將不同的數據庫和云應用程序連接到 .NET 數據相關程序和技術。 .NET Framework 支持 我們的
    的頭像 發表于 01-02 09:31 ?82次閱讀
    <b class='flag-5'>通過</b>Skyvia Connect SQL終端節點<b class='flag-5'>訪問</b>任何數據

    虛擬內存和云計算的關系

    虛擬內存是一種計算機系統內存管理技術,它通過將物理內存與磁盤空間結合起來,使得應用程序可以訪問
    的頭像 發表于 12-04 09:50 ?145次閱讀

    虛擬內存的作用和原理 如何調整虛擬內存設置

    虛擬內存,也稱為虛擬內存管理或頁面文件,是計算機操作系統中的一種內存管理技術。它允許系統使用硬盤空間作為額外的RAM(隨機存取存儲器),以彌補物理內存(RAM)的不足。虛擬
    的頭像 發表于 12-04 09:13 ?574次閱讀

    DDR內存頻率對性能的影響

    的整體性能,特別是在處理大量數據或運行復雜程序時,如視頻編輯、3D渲染和高端游戲等場景。快速的內存可以縮短CPU和內存之間的通信時間,從而提升整體系統的響應效率。 二、多任務處理能力
    的頭像 發表于 11-20 14:25 ?1104次閱讀

    什么是RAM內存 RAM內存對電腦性能的影響

    所有存儲的數據。RAM的主要作用是提供快速的數據訪問,以便CPU(中央處理器)能夠高效地執行程序和處理任務。 RAM內存對電腦性能的影響 RAM對電腦
    的頭像 發表于 11-11 09:38 ?1727次閱讀

    【「算力芯片 | 高性能 CPU/GPU/NPU 微架構分析」閱讀體驗】--了解算力芯片GPU

    Directx是由微軟設計和維護的圖形應用程序接口。。DirectX提供了一個統的、跨硬件的接口,使得開發者只需針對一個 API進行編程即可。DirectX還定義了一系列圖形和聲音功能的標準,這使得硬件制造商知道
    發表于 11-03 12:55

    如何提高GPU性能

    學習和機器學習等領域發揮著重要作用。 1. 硬件升級 a. 更換高性能GPU 最直接的提高GPU性能的方法是升級到更高
    的頭像 發表于 10-27 11:21 ?636次閱讀

    GPU超頻設置技巧

    超頻GPU(圖形處理單元)可以提高顯卡的性能,但同時也可能增加熱量和功耗,甚至可能縮短硬件的壽命。在進行GPU超頻之前,確保你了解可能的風險,并且愿意承擔這些風險。以下是一些基本的
    的頭像 發表于 10-27 11:09 ?528次閱讀

    通過DaVinci TMS320DM644x的串行接口加載基本應用程序

    電子發燒友網站提供《通過DaVinci TMS320DM644x的串行接口加載基本應用程序.pdf》資料免費下載
    發表于 10-16 11:52 ?0次下載
    <b class='flag-5'>通過</b>DaVinci TMS320DM644x的串行接口加載基本<b class='flag-5'>應用程序</b>

    如何檢測內存泄漏

    檢測內存泄漏是軟件開發過程中一項至關重要的任務,它有助于識別和解決那些導致程序占用過多內存資源,從而影響程序性能甚至導致程序崩潰的問題。以下
    的頭像 發表于 07-30 11:50 ?2073次閱讀

    能否建議通過內存映射提高aur性能

    我們使用英飛凌 Aurix TC364 作為平臺。 我們試圖通過將 ISR 映射到不同的內存單元來提高 ISR 函數的性能,具體方法是使用 #pragma section code
    發表于 05-29 07:39

    如何去提高EtherCAT IO的性能呢?

    進行EtherCAT IO性能優化涉及多個方面,包括硬件選擇、網絡配置、軟件優化和應用程序設計。
    的頭像 發表于 03-07 09:28 ?529次閱讀

    如何在ModusToolbox?中檢查和設置應用程序內存地址?

    如何在ModusToolbox?中檢查和設置應用程序內存地址?
    發表于 03-01 10:16

    LeftoverLocals漏洞致多家知名品牌GPU內存關鍵數據被盜

    此漏洞不限于消費者應用程序,更嚴重的是,它能通過破壞大型語言模型(LLM)及機器學習(ML)方案所依賴的 GPU 以獲取關鍵數據。由于模型訓練涉及敏感數據,故此時捕獲數據的風險極大。
    的頭像 發表于 01-19 10:09 ?444次閱讀
    主站蜘蛛池模板: 丰满放荡岳乱妇91www| 国产美女视频黄a视频免费全过程 国产美女视频黄a视频全免费网站 | 轻点灬大ji巴太粗太长了爽文| 99pao强力打造免费高清色| 亚洲卡5卡6卡7国色天香| 免费黄视频网站| 一级a爱片久久毛片| 四虎国产精品高清在线观看| 色噜噜狠狠色综合欧洲| 欧美二区三区| 国产在线精彩视频二区| 91精品国产91久久久久久青草| 天天射天天舔| 插插插天天| 亚洲国产一区二区三区a毛片| 中文字幕不卡在线播放| 婷婷六月天激情| 欧美乱论视频| www.黄视频| 久久久久久亚洲精品| 黄篇网站在线观看| 中文字幕11页| 奇米影色777四色在线首页| 午夜aa| 34pao强力打造免费永久视频| 男男憋尿play按小腹| 亚洲va久久久噜噜噜久久狠狠| 日韩毛片高清在线看| 国内黄色一级片| 天天做天天爱天天做天天吃中| 国产情侣自拍小视频| 欧美日韩一区二区三区毛片| 五月亭亭免费高清在线| 久久www免费人成高清| 天天综合在线视频| 日本xxxxxx69| 曰本福利写真片视频在线| 日本xxxx色视频在线观看免费| 复古毛片| 精品国产一区二区三区成人| 亚洲一区毛片|