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