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

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

3天內(nèi)不再提示

使用MPI來跨多個GPU縮放應用

星星科技指導員 ? 來源:NVIDIA ? 作者:NVIDIA ? 2022-04-27 17:08 ? 次閱讀

這是 標準并行編程 系列的第三篇文章,講述在標準語言中使用并行性來加速計算的優(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 端口在單精度和雙精度浮點上的性能。

poYBAGJpCC2AfDpDAAATAbRc9a4263.png

圖 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 上獲得的性能進行了比較:

pYYBAGJpCC6ASfHwAAATabMyqSc445.png

圖 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 之前,杰夫在位于橡樹嶺國家實驗室的克雷超級計算卓越中心工作。

審核編輯:郭婷

聲明:本文內(nèi)容及配圖由入駐作者撰寫或者入駐合作網(wǎng)站授權轉載。文章觀點僅代表作者本人,不代表電子發(fā)燒友網(wǎng)立場。文章及其配圖僅供工程師學習之用,如有內(nèi)容侵權或者其他違規(guī)問題,請聯(lián)系本站處理。 舉報投訴
  • cpu
    cpu
    +關注

    關注

    68

    文章

    10876

    瀏覽量

    212126
  • NVIDIA
    +關注

    關注

    14

    文章

    5002

    瀏覽量

    103236
  • gpu
    gpu
    +關注

    關注

    28

    文章

    4749

    瀏覽量

    129034
收藏 人收藏

    評論

    相關推薦

    車間有多個PLC時,不同網(wǎng)段如何實現(xiàn)網(wǎng)段訪問?

    不同的網(wǎng)段時,上位機是不能直接通信到兩臺設備的,需要修改設備IP實現(xiàn)通信,但往往耗時費力且成本較高。 對此,物通博聯(lián)提供基于PLC網(wǎng)段通信網(wǎng)關的解決方案。它能夠生產(chǎn)網(wǎng)絡對 PLC、HMI、CNC、DCS、MES等網(wǎng)絡通訊和數(shù)據(jù)采集,兼有
    的頭像 發(fā)表于 12-27 17:15 ?152次閱讀
    車間有<b class='flag-5'>多個</b>PLC時,不同網(wǎng)段如何實現(xiàn)<b class='flag-5'>跨</b>網(wǎng)段訪問?

    《CST Studio Suite 2024 GPU加速計算指南》

    - smi啟用或禁用。 - 其他:還介紹了獨占計算模式、顯示鏈接、組合MPI計算和GPU計算、服務用戶、GPU計算使用Windows遠程桌面、運行多個模擬、視頻卡驅動、操作條件、最
    發(fā)表于 12-16 14:25

    LMK61E08能否配合多個跟隨器驅動多個器件?

    你好,LMK61E08能否配合多個跟隨器驅動多個器件,比如ADC。這樣會造成時鐘性能比如jitter變差嗎?
    發(fā)表于 11-12 07:20

    GPU服務器AI網(wǎng)絡架構設計

    眾所周知,在大型模型訓練中,通常采用每臺服務器配備多個GPU的集群架構。在上一篇文章《高性能GPU服務器AI網(wǎng)絡架構(上篇)》中,我們對GPU網(wǎng)絡中的核心術語與概念進行了詳盡介紹。本文
    的頭像 發(fā)表于 11-05 16:20 ?407次閱讀
    <b class='flag-5'>GPU</b>服務器AI網(wǎng)絡架構設計

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

    不同的線程束以執(zhí)行不同的指令。指令調度單元負責從指令緩存中取出著色器程序中的操作指令,并將其分配給每個CUDA核心執(zhí)行。 線程塊的設計為 GPU編程提供了一種高效、靈活和易用的方式組織和管理并行線程
    發(fā)表于 11-03 12:55

    如何選擇適合的GPU

    。因此,選擇一個適合自己需求的GPU是一個需要綜合考慮多個因素的過程。 1. 確定用途 首先,你需要明確你的GPU將用于什么目的。不同的應用場景對GPU的需求差異很大。例如,如果你主要
    的頭像 發(fā)表于 10-27 11:07 ?343次閱讀

    GPU加速計算平臺是什么

    GPU加速計算平臺,簡而言之,是利用圖形處理器(GPU)的強大并行計算能力加速科學計算、數(shù)據(jù)分析、機器學習等復雜計算任務的軟硬件結合系統(tǒng)。
    的頭像 發(fā)表于 10-25 09:23 ?259次閱讀

    有沒有大佬知道NI vision 有沒有辦法通過gpu和cuda加速圖像處理

    有沒有大佬知道NI vision 有沒有辦法通過gpu和cuda加速圖像處理
    發(fā)表于 10-20 09:14

    使用功率縮放

    電子發(fā)燒友網(wǎng)站提供《使用功率縮放庫.pdf》資料免費下載
    發(fā)表于 10-18 10:24 ?0次下載
    使用功率<b class='flag-5'>縮放</b>庫

    DM642 EVM上的視頻縮放示例

    電子發(fā)燒友網(wǎng)站提供《DM642 EVM上的視頻縮放示例.pdf》資料免費下載
    發(fā)表于 10-16 10:52 ?0次下載
    DM642 EVM上的視頻<b class='flag-5'>縮放</b>示例

    干貨分享:宏集物聯(lián)網(wǎng)HMI通過S7 MPI協(xié)議采集西門子400PLC數(shù)據(jù)

    宏集物聯(lián)網(wǎng)HMI集成了多個驅動適配西門子200、300、400、1200、1500、LOGO等系列PLC,本文主要介紹宏集HMI通過S7 MPI協(xié)議采集西門子400PLC數(shù)據(jù)的操作步驟
    的頭像 發(fā)表于 06-13 13:39 ?607次閱讀
    干貨分享:宏集物聯(lián)網(wǎng)HMI通過S7 <b class='flag-5'>MPI</b>協(xié)議采集西門子400PLC數(shù)據(jù)

    網(wǎng)段IP耦合器是什么?網(wǎng)段IP耦合器的功能作用

    通過別的網(wǎng)段訪問,網(wǎng)絡耦合器能夠解決網(wǎng)段訪問的問題。 車間多個設備,如果具有相同的IP地址,則可以統(tǒng)一成同網(wǎng)段不同的IP地址,進行在線編程、數(shù)據(jù)采集、遠程管理等各種網(wǎng)段訪問。 車間多個
    的頭像 發(fā)表于 04-18 09:25 ?465次閱讀

    fpga和gpu的區(qū)別

    FPGA(現(xiàn)場可編程門陣列)和GPU(圖形處理器)在多個方面存在顯著的區(qū)別。
    的頭像 發(fā)表于 03-27 14:23 ?1227次閱讀

    阻結構的優(yōu)點是怎么的的?

    阻結構是指在反饋網(wǎng)絡中引入一個電阻到地的結構。這樣做的目的是為了保持輸出電平穩(wěn)定,減小輸出偏移,提高放大器的穩(wěn)定性和精度。具體來說怎么體現(xiàn)阻結構的穩(wěn)定性? 常用的網(wǎng)絡結構除了T型網(wǎng)絡,還有π型
    發(fā)表于 01-26 12:30

    為什么GPU比CPU更快?

    GPU比CPU更快的原因并行處理能力:GPU可以同時處理多個任務和數(shù)據(jù),而CPU通常只能一次處理一項任務。這是因為GPU的架構使得它可以同時處理多個
    的頭像 發(fā)表于 01-26 08:30 ?2457次閱讀
    為什么<b class='flag-5'>GPU</b>比CPU更快?