這是 標準并行編程 系列的第三篇文章,講述在標準語言中使用并行性來加速計算的優(yōu)點。
用標準語言并行性開發(fā)加速代碼
多個 GPU 標準 C ++并行編程,第 1 部分
在第 1 部分中,我們解釋了:
C ++并行編程的基礎
格子玻爾茲曼方法( LBM )
采取了第一步來重構 PalabOS 庫,以使用標準 C ++高效地運行 GPU 。
在這篇文章中,我們繼續(xù)優(yōu)化 ISOC ++算法的性能,然后使用 MPI 來跨多個 GPU 來縮放應用。
爭取最佳性能
期望 CPU 到 GPU 端口的性能低于專用 HPC 代碼的性能似乎很自然。畢竟,您受到軟件體系結構、已建立的 API 的限制,以及考慮用戶群期望的復雜額外功能的需要。不僅如此, C ++標準并行化的簡單編程模型允許比專用語言(如 CUDA )更少的手動微調。
在現(xiàn)實中,通??梢詫⑦@種性能損失控制和限制到可以忽略不計的程度。關鍵是分析各個代碼部分的性能指標,消除不能反映軟件框架實際需求的性能瓶頸。
一個好的做法是為數(shù)值算法的核心組件維護一個單獨的原理證明代碼。這種方法的性能可以更自由地優(yōu)化,并與完整、復雜的軟件框架(如 Palabos 中的 STLBM library )進行比較。此外,像nvprof這樣支持 GPU 的探查器可以有效地突出性能瓶頸的根源。
以下建議重點介紹了典型的性能問題及其解決方案:
不要觸摸 CPU 上的數(shù)據(jù)
了解你的算法
建立績效模型
不要觸摸 CPU 上的數(shù)據(jù)
性能損失的一個常見原因是 CPU 和 GPU 內(nèi)存之間的隱藏數(shù)據(jù)傳輸,這可能非常慢。在 CUDA 統(tǒng)一內(nèi)存模型中,每當您從 CPU 訪問 GPU 數(shù)據(jù)時,就會發(fā)生這種類型的傳輸。觸摸單個字節(jié)的數(shù)據(jù)可能會導致災難性的性能損失,因為整個內(nèi)存頁都是一次性傳輸?shù)摹?/p>
顯而易見的解決方案是盡可能只在 GPU 上操作數(shù)據(jù)。這需要仔細搜索代碼中所有對數(shù)據(jù)的訪問,然后將它們包裝成并行算法調用。雖然這有點健壯,但即使是最簡單的操作也需要這個過程。
顯而易見的地方是數(shù)據(jù)統(tǒng)計的后處理操作或中間評估。另一個經(jīng)典的性能瓶頸出現(xiàn)在 MPI 通信層,因為您必須記住在 GPU 上執(zhí)行數(shù)據(jù)打包和解包操作。
在 GPU 上表達算法說起來容易做起來難,因為for_each和transform_reduce的形式主義主要適用于結構均勻的內(nèi)存訪問。
在不規(guī)則數(shù)據(jù)結構的情況下,使用這兩種算法避免競爭條件并保證合并的內(nèi)存訪問是痛苦的。在這種情況下,您應該繼續(xù)執(zhí)行下一個建議,熟悉 C ++中提供的并行算法的家族。
了解你的算法
到目前為止,并行 STL 似乎只不過是一種用奇特的函數(shù)語法表達parallel for loops的方式。實際上, STL 提供了for_each和transform_reduce之外的大量算法,這些算法對表達數(shù)值方法非常有用,包括排序和搜索算法。
exclusive_scan算法計算累積和,值得特別提及,因為它被證明通常對非結構化數(shù)據(jù)的重新索引操作非常有用。例如,考慮 MPI 通信的打包算法,其中預先由每個網(wǎng)格節(jié)點貢獻給通信緩沖器的變量的數(shù)目是未知的。在這種情況下,需要線程之間的全局通信來確定每個網(wǎng)格節(jié)點寫入緩沖區(qū)的索引。
下面的代碼示例顯示了如何使用并行算法在 GPU 上以良好的并行效率解決此類問題:
// Step 1: compute the number of variables contributed by every node.
int* numValuesPtr = allocateMemory(numberOfCells);
for_each(execution::par_unseq, numValuesPtr, numValuesPtrl + numberOfCells, [=](int& numValues)
{ int i = &numValues - numValuesPtr; // Compute number of variables contributed by current node. numValues = computeNumValues(i);
} );
// 2. Compute the buffer index for every node.
int* indexPtr = allocateMemory(numberOfCells);
exclusive_scan(execution::par_unseq, numValuesPtr, numValuesPtr + numberOfCells, indexPtr, 0);
// 3. Pack the data into the buffer.
for_each(execution::par_unseq, indexPtr, indexPtr + numberOfCells, [=](int& index)
{ int i = &index - indexPtr; packCellData(i, index);
} );
這個例子讓你享受到基于算法的 GPU 編程方法的表達能力:代碼不需要同步指令或任何其他低級構造。
建立績效模型
性能模型通過瓶頸分析為算法的性能建立上限。這通常將峰值處理器性能(以觸發(fā)器為單位)和峰值內(nèi)存帶寬視為限制硬件特性的主要因素。
正如在上一篇文章的示例:Lattice Boltzmann 軟件 Palabos 部分中所討論的,LBM 代碼的計算與內(nèi)存訪問的比率較低,并且在現(xiàn)代 GPU 上完全受內(nèi)存限制。也就是說,至少如果您使用單精度算術或為雙精度算術優(yōu)化的 GPU。
峰值性能簡單地表示為 GPU 的內(nèi)存帶寬與代碼中執(zhí)行的內(nèi)存訪問次數(shù)之間的比率。直接的結果是,將 LBM 代碼從雙精度算術轉換為單精度算術將使性能加倍。
圖 1 顯示了在 NVIDIA A100 ( 40 GB ) GPU 上獲得的 Palabos GPU 端口在單精度和雙精度浮點上的性能。
圖 1 。 3D 蓋驅動腔體的 Palabos 性能( 6003網(wǎng)格節(jié)點)在 A100 ( 40 GB ) GPU 上以單精度和雙精度運行。型號: TRT , D3Q19
執(zhí)行的測試用例是湍流狀態(tài)下蓋驅動腔中的流動,具有簡單的立方幾何結構。然而,這種情況包括邊界條件,并表現(xiàn)出復雜的流動模式。性能以每秒百萬次晶格節(jié)點更新( MLUPS ,越多越好)來衡量,并與假設 GPU 內(nèi)存在峰值容量下被利用的理論峰值進行比較。
該代碼在雙精度下達到 73% 的峰值性能,在單精度下達到 74% 。這種性能指標在 LB 模型的最新實現(xiàn)中很常見,與所使用的語言或庫無關。
盡管一些實現(xiàn)可能會增加幾個百分點,并達到接近 80% 的值,但很明顯,我們正在接近性能模型隱含的硬限制。從大的角度來看,代碼的單 – GPU 性能是最好的。
重用現(xiàn)有的 MPI 后端以獲得多 GPU 代碼
當 C ++并行算法無縫地集成到現(xiàn)有的軟件項目中以加速關鍵代碼部分時,沒有什么能阻止您重用項目的通信后端以達到多 GPU 性能。但是,您需要密切關注通信緩沖區(qū),確保它不會繞過 CPU 內(nèi)存,這將導致代價高昂的頁面錯誤。
我們首次嘗試在多個 GPU 上運行帶有 GPU 端口的 Palabos 版本,雖然產(chǎn)生了技術上正確的結果,但沒有表現(xiàn)出可接受的性能。不是加速,而是從 1 切換到 2 GPU 將速度降低了一個數(shù)量級。這個問題可以追溯到通信數(shù)據(jù)的打包和解包。在最初的后端,這是在 CPU 上執(zhí)行的,并且是在 CPU 內(nèi)存中的其他不必要數(shù)據(jù)訪問實例上執(zhí)行的,例如調整通信緩沖區(qū)的大小。
這些問題可以在探查器的幫助下發(fā)現(xiàn)。分析器會突出顯示統(tǒng)一內(nèi)存中出現(xiàn)的所有頁面錯誤,并通過將相應的代碼部分移動到并行算法來修復?!傲私饽愕乃惴ā辈糠纸忉屃巳绻麛?shù)據(jù)遵循不規(guī)則模式,如何打包和解包通信緩沖區(qū)。
在這一點上,使用標準的 C ++,除了 MPI 以外沒有任何擴展,您可以獲得一個混合 CPU / GPU 軟件項目,具有最先進的性能,在單 G GPU 和多 GPU 上的并行性能。
不幸的是,由于語言規(guī)范和相應的 GPU 實現(xiàn)的當前限制,多 GPU 性能仍然低于預期。在未來的 C ++技術標準并行化技術的改進中,我們將基于 C ++標準之外的技術提供一些解決方案。
協(xié)調多 CPU 和多 GPU 代碼的執(zhí)行
雖然這篇文章主要關注混合 CPU 和 GPU 編程,但我們無法避免在 CPU 部分討論混合并行性( MPI 或多線程)問題。
例如, Palabos 的原始版本是非混合的,它使用 MPI 通信層在 CPU 的核心之間以及整個網(wǎng)絡中分配工作。移植到 GPU 后,生成的多 CPU 和多 GPU 代碼會自動將單個 CPU 內(nèi)核與每個 MPI 任務中的完整 GPU 進行分組,使 CPU 的動力相對不足。
每當需要或方便將計算密集型任務保留在 CPU 上時,這會導致性能瓶頸。在流體動力學中,在預處理階段(如幾何體處理或網(wǎng)格生成)通常會出現(xiàn)這種情況。
顯而易見的解決方案是使用多線程從 MPI 任務中訪問多個 CPU 內(nèi)核。這些線程的共享內(nèi)存空間可以通過 CUDA 統(tǒng)一內(nèi)存形式直接與 GPU 共享。
然而, C ++并行算法不能被重用以服務于 GPU 和多核 CPU 執(zhí)行的兩個目的。這是因為 C ++不允許從語言內(nèi)選擇并行算法的目標平臺。
雖然 C ++線程確實提供了一種解決這個問題的方法,但我們發(fā)現(xiàn) OpenMP 提供了最方便和最不受干擾的解決方案。在這種情況下,for loop的 OpenMP 注釋足以將分配給當前 MPI 任務的網(wǎng)格部分分布到多個線程上。
通過固定內(nèi)存進行通信
在當前版本的 HPC SDK 中, CUDA 統(tǒng)一內(nèi)存模型與 MPI 相結合,表現(xiàn)出另一個性能問題。
由于 MPI 通信層希望數(shù)據(jù)具有固定的硬件地址(所謂的pinned memory),因此托管內(nèi)存區(qū)域中的任何緩沖區(qū)都會首先隱式復制到主機 CPU 上的固定內(nèi)存緩沖區(qū)中。由于 GPU 和 CPU 之間的傳輸,此操作最終可能會非常昂貴。
因此,通信緩沖區(qū)應明確固定到 GPU 內(nèi)存地址。對于nvc++ compiler,這是通過使用cudaMalloc分配通信緩沖區(qū)來實現(xiàn)的:
// Allocate the communication buffer
// vector《double》 buffer(N);
// double* buffer = buffer.data();
double* buffer; cudaMalloc((void**)&buffer, N * sizeof(double));
for_each(buffer, buffer + N, … // Proceed with data packing
另一種解決方案是用推力庫中的thrust::device_vector替換 STL 向量,默認情況下,推力庫使用固定 GPU 內(nèi)存。
在不久的將來, HPC SDK 將為用戶更高效、更自動地處理這些情況。這樣他們就不必伸手去拿cudaMalloc或thrust::device_vector。所以,請繼續(xù)關注!
在本文列出的各種改進之后, Palabos 庫在一個帶有四個 GPU 的 DGX A100 ( 40-GB )工作站上進行了測試,再次用于蓋驅動型腔的基準情況。獲得的性能如圖 2 所示,并與 48 核 Xeon Gold 6240R CPU 上獲得的性能進行了比較:
圖 2 。 3D 蓋驅動腔體的 Palabos 性能( 6003網(wǎng)格節(jié)點)在 48 核 Xeon Gold 6240R CPU 和 DGX A100 ( 40 GB )工作站上,一次使用一個 GPU ,一次使用四個 GPU 。型號: TRT , D3Q19 ,單精度
對于 Xeon Gold , Palabos 的原始實現(xiàn)被證明更高效,并用于 48 個 MPI 任務,而單 GPU 和四 GPU 執(zhí)行使用并行算法后端,并使用nvc++編譯。
性能數(shù)據(jù)顯示,與單次執(zhí)行 GPU 相比, 4- GPU 執(zhí)行的速度提高了 3.27 倍。這相當于一個非常令人滿意的并行效率 82% ,在一個強大的擴展機制,在兩個執(zhí)行相同的總域大小。在弱擴展情況下,使用 4 倍于 4- GPU 執(zhí)行的問題大小,加速比增加到 3.72 (效率 93% )。
圖 2 還顯示,當使用未固定的通信緩沖區(qū)時,例如當 MPI 通信緩沖區(qū)未分配cudaMalloc時,并行效率從 82% 下降到 61% 。
最終,四 GPU DGX 工作站的運行速度比 Xeon Gold CPU 快 55 倍。雖然由于兩臺機器的范圍不同,直接比較可能不公平,但它提供了通過將代碼移植到 GPU 獲得的加速度感。 DGX 是一個連接到公共電源插頭的臺式工作站,但它提供的性能在 CPU 群集上只能通過數(shù)千個 CPU 內(nèi)核獲得。
結論
您已經(jīng)看到 C ++標準語言并行性可以用來把像 PalabOS 這樣的庫移植到 GPU ,代碼性能驚人地提高。
對于 Palabos 庫的最終用戶來說,這種性能提升是通過一行更改來實現(xiàn)的,即從 CPU 后端切換到 GPU 后端。
對于 Palabos 庫開發(fā)人員來說,開發(fā)相應的 GPU 后端需要做一些工作。
然而,這項工作不需要學習新的領域特定語言,也不依賴于 GPU 體系結構的詳細知識。
關于作者
Jonas Latt 是瑞士日內(nèi)瓦大學計算機科學系的副教授。他從事高性能計算和計算流體力學的研究,并在包括地球物理、生物醫(yī)學和航空航天領域在內(nèi)的跨學科領域進行應用。他是 lattice Boltzmann 復雜流動模擬開源軟件 Palabos 的最初開發(fā)者和當前共同維護者。他以前在日內(nèi)瓦大學獲得物理學和計算機科學博士學位,并通過塔夫斯大學(波士頓,美國)和綜合理工學校 F.E.EdRaelde 洛桑 EPFL (瑞士)的研究,并作為 CFD 公司 FuluKIT 的聯(lián)合創(chuàng)始人,對流體力學感興趣。
Christophe Guy Coreixas 是一名航空工程師, 2014 年畢業(yè)于 ISAE-SUPAERO (法國圖盧茲)。 2018 年,他在 CERFACS 從事面向行業(yè)應用的可壓縮晶格玻爾茲曼方法研究時獲得了博士學位(流體動力學)。作為日內(nèi)瓦大學計算機科學系的博士后,克里斯多夫現(xiàn)在開發(fā)格子玻爾茲曼模型來模擬航空、多物理和生物醫(yī)學流程。
Gonzalo Brito 是 NVIDIA 計算性能與 HPC 團隊的高級開發(fā)技術工程師,工作于硬件和軟件的交叉點。他熱衷于讓加速計算變得更容易實現(xiàn)。在加入NVIDIA 之前,岡薩洛在 RWTH 亞琛大學空氣動力學研究所開發(fā)了多物理方法,用于顆粒流。
Jeff Larkin 是 NVIDIA HPC 軟件團隊的首席 HPC 應用程序架構師。他熱衷于高性能計算并行編程模型的發(fā)展和采用。他曾是 NVIDIA 開發(fā)人員技術小組的成員,專門從事高性能計算應用程序的性能分析和優(yōu)化。 Jeff 還是 OpenACC 技術委員會主席,曾在 OpenACC 和 OpenMP 標準機構工作。在加入NVIDIA 之前,杰夫在位于橡樹嶺國家實驗室的克雷超級計算卓越中心工作。
審核編輯:郭婷
-
cpu
+關注
關注
68文章
10876瀏覽量
212126 -
NVIDIA
+關注
關注
14文章
5002瀏覽量
103236 -
gpu
+關注
關注
28文章
4749瀏覽量
129034
發(fā)布評論請先 登錄
相關推薦
評論