@2014__TPDS__GPU-Aware MPI on RDMA-Enabled Clusters - Design, Implementation and Evaluation

# GPU-Aware MPI on RDMA-Enabled Clusters: Design, Implementation and Evaluation > [!abstract] 概要 > GPUクラスタ上で高性能かつスケーラブルなアプリケーションを設計するには、いくつかの課題に取り組む必要がある。鍵となる課題は、ホストメモリとデバイスメモリが分離していることであり、これは異なるメモリ空間上のデータを操作するために、プログラマがCUDAとMPIなど複数のプログラミングモデルを使うことを要求する。多次元構造の非連続データが実世界のアプリケーションで使われる場合、この課題への対処はさらに難しくなる。これらの課題はプログラミング生産性とアプリケーション性能を制限する。本論文では、標準MPIを用いてGPU間データ通信をサポートするGPU-Aware MPIを提案する。これは分離したメモリ空間を統一し、明示的なCPU-GPUデータ移動とCPU/GPUバッファ管理を回避する。GPU-Aware MPIは、GPUデータ型ベクトル化アルゴリズムとベクトルベースGPUカーネルデータpack/unpackアルゴリズムという2つのアルゴリズムにより、デバイスメモリ上のすべてのMPIデータ型をサポートする。非連続データのpack/unpack、PCIe上のデータ移動、ネットワーク上のRDMAデータ転送を重ね合わせるためのパイプラインが設計されている。本設計をオープンソースMPIライブラリMVAPICH2に組み込み、実運用アプリケーションであるmultiphase 3D LBMを最適化した。プログラミング生産性の向上に加え、Oakleyスーパーコンピュータの64GPU上で、アプリケーションレベル性能が最大19.9%改善することを確認した。 ## 論文情報 - タイトル: GPU-Aware MPI on RDMA-Enabled Clusters: Design, Implementation and Evaluation - 著者: Hao Wang, Sreeram Potluri, Devendar Bureddy(以上 Ohio State University)、Carlos Rosales(Texas Advanced Computing Center, UT Austin)、Dhabaleswar K. (DK) Panda(Ohio State University) - 媒体: IEEE Transactions on Parallel and Distributed Systems (TPDS), Vol. 25, No. 10, October 2014, pp. 2595-2605 - DOI: 10.1109/TPDS.2013.222(受理2013-08-03、オンライン先行公開2013-08-28、正式版2014-09-17) - 助成: US NSF grants #OCI-0926691, #OCI-1148371, #CCF-1213084 ## 概要 GPUクラスタではホストメモリとデバイスメモリが分離しており、プログラマはCUDAとMPIを併用して両メモリ空間を扱う必要がある。本論文は、この分離を隠蔽してGPUデータを標準MPI APIだけで送受信できるGPU-Aware MPIを提案し、非連続データ型(vector・subarray・struct等)への対応と、パック/転送を重ねる非同期パイプラインを設計する。これをMVAPICH2 1.8に実装し、multiphase Lattice Boltzmann Method(LBM)を最適化して評価する。著者らの先行研究([[MVAPICH2]]-GPU、2011年)が連続データのGPU間通信をサポートしていたのに対し、本論文はベクトル化アルゴリズムとGPUカーネルによるpack/unpackで非連続データ型対応を拡張したものである。 ## 問題設定 - **入力**: ホストメモリまたはデバイスメモリ上に存在するMPI送受信バッファ(連続データ・非連続データいずれも可)。 - **出力**: 標準MPI API(`MPI_Send`/`MPI_Recv`/`MPI_Isend`/`MPI_Irecv`等)でGPUデバイスメモリ上のデータを直接送受信できること。プログラマ側の明示的なCPU-GPUデータ移動コードとバッファ管理コードは不要になる。 - **前提**: 対象環境ではGPUDirect RDMAが利用できず(本文Section 2.1で明言)、ホストメモリを経由するステージングが必須。CUDA 4.0以降のUVA(Unified Virtual Addressing)でポインタのメモリ種別(ホスト/デバイス)を判定できることが前提。 - **必要なデータ**: MPI標準3.0で定義される派生データ型(vector、hvector、indexed、hindexed、blockindexed、struct、subarray、darray)のサポート範囲(Fig. 2)。 **Figure 1: ヘテロジニアスGPUクラスタのアーキテクチャ** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig01-architecture.png]] (Fig. 1. Architecture for Heterogeneous Cluster with GPGPUs. 各ノードはGPGPU・ホスト・NIC(Network Card)がPCIeで接続され、GDDRを積んだGPUデバイスメモリとDDRのホストメモリが分離している。ノード間はPCIe HubとNICを介してネットワークで接続される。この分離構造こそが、CUDAとMPIを併用させる根本原因である。) **Figure 2: MPI標準が定義する派生データ型の階層** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig02-mpi-datatypes.png]] (Fig. 2. MPI Datatypes defined in the MPI standard and supported by our designs. contiguous・vector・indexed・struct・subarrayを頂点に、vectorの子としてhvector、indexedの子としてhindexed・blockindexed、subarrayの子としてdarrayが定義される。3D LBMアプリケーションで実際に使うのはvector・struct・subarrayである。) ## 提案手法 ### アーキテクチャ - **メモリ種別検出**: MPIランタイムは各MPI呼び出しの実行時に、UVAを用いて送受信バッファがホストメモリかデバイスメモリかを判定し、その結果をMPIリクエストごとのデータ構造にキャッシュして、UVA問い合わせの重複呼び出しを避ける。 - **連続データ用バッファプール**: `vbuf`という名のホストメモリ上のバッファプールを`MPI_Init()`/`MPI_Finalize()`で確保・解放し、ブロック単位でプロセスごとに1つ維持する(プロセスペアごとの $O(N^2)$ ではなくプロセスごとの $O(N)$)。送信側はソースアドレスがデバイスメモリと判定されると`vbuf`を1つ取得してCUDAメモリコピーでデータを転送し、RTS(Request To Send)/CTS(Clear To Send)ハンドシェイクの後にRDMA writeを発行する。 ### アルゴリズム/手法の詳細 **Algorithm 1: GPUデータ型ベクトル化** IOV(Input/Output Vector、各連続セグメントのポインタと長さの組)から、ブロック数(`v_count`)・ブロック長(`v_blen`)・ブロック間ストライド(`v_stride`)の三つ組として定義されるベクトルへ変換する。`load_iov()`でIOVを順に読み込み、連続する2セグメントの長さとストライドが一致すれば同一ベクトルへ統合し(Line 12)、不一致になった時点で`put_vector()`によりグローバル配列へ確定したベクトルを積む。この処理はCPU上で実行される(GPU上での条件分岐処理のオーバーヘッドが高いため)。MPI subarrayデータ型は多次元ベクトルとして扱えるため、ベクトル化ステップを経ずに直接GPUカーネルpack/unpack関数を呼び出す。 **Algorithm 2/3: 非同期パイプラインとGPUカーネルpack/unpack** データを`cuda_block`という単位に分割し、GPUデバイスメモリ上の`tbuf`バッファプールを介して、非連続データのGPU内pack/unpack、PCIe上のデータ移動、ネットワーク上のデータ転送を異なるブロックについて重ね合わせる。送信側では、送信元アドレスがデバイスメモリかつ非連続データ型であれば、GPUカーネルで`tbuf`へ非同期にpackし、受信側ではホスト→デバイス転送完了後に`tbuf`から宛先へ非同期にunpackする。MPI受信呼び出しは、全てのunpack完了イベントが検出された後に完了する。Algorithm 3(vector_pack_gpu/カーネル本体)は、送信元アドレス`src`(ピッチ`spitch`)から宛先アドレス`dst`(ピッチ`dpitch`)へ、複数スレッドが元データ型サイズ単位でコピーする単純な2次元グリッドカーネルである。ベクトルのブロック長が`cuda_block`サイズ以上であれば、pack/unpackステップを省略し`tbuf`のオフセットを調整して`vbuf`へ直接非同期コピーする。 **Figure 4: 非同期パイプラインの模式図** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig04-async-pipeline.png]] (Fig. 4. Asynchronous pipeline design in GPU-aware MPI. 送信側はGPUメモリ(src)→GPUメモリ(tbuf)→ホストメモリ(vbuf)の順にvector_pack_gpu()とcudaEventQuery()で確認しながらパイプライン化し、RTS/CTSハンドシェイク後にRDMA writeで転送、受信側は逆順にcudaMemcpyAsync()とvector_unpack_gpu()でホスト経由デバイスメモリへ書き戻す。) ### 実装上の工夫 - **ブロックサイズのチューニング**: パイプラインの各段(Pack, D2H, RDMA, H2D, Unpack)の実行時間が均一になる最適ブロックサイズ $n$ が存在し(式(2))、総転送時間は $T_{total}(N) = (N/n + 2) \cdot T_{block}(n)$ で近似できる(式(3))。実測ではOakley環境で16KBが最適ブロックサイズであり、この閾値を超えるデータでのみパイプラインが自動的に有効化される(16KB未満では有意差なし)。ブロックサイズはクラスタ固有の特性としてマイクロベンチマークスイート(著者らの別論文「OMB-GPU」)でシステム管理者がインストール時に1回チューニングし、設定ファイルに配置することで、エンドユーザーには透過的になる。 - **GPUDirect RDMA環境への拡張性**: 本設計はGPUDirect RDMA利用時にも有効であり、その場合CPU-GPU間のデータ移動のみがバイパスされ、pack/unpackはGPUDirect RDMA経由のデータ転送と重ね合わせられる。関連研究([[GPUDirect RDMA]]を用いた先行研究)では、中〜大メッセージについてパイプライン設計がGPUDirect RDMA単体より優れる場合があると報告されている。 ### 対象アプリケーション: multiphase 3D LBM 評価に用いるmultiphase Lattice Boltzmann Method([[格子ボルツマン法]])は、Zhengらが提案したFree-energyモデルに基づく実装で、運動量分布にD3Q19離散化、相の分布にD3Q7離散化を用いる(D3Q7はD3Q19の部分方向集合)。各反復は次の6ステップからなる: (1) $f$・$g$配列からρ・velocity配列を更新、(2) 近傍とφを交換、(3) $f$・$g$のcollisionルーチンを実行、(4) 近傍と$f$を交換、(5) $f$・$g$のstreamingルーチンを実行、(6) 近傍と$f$・$g$を交換。ステップ(2)(4)の交換はD3Q7に基づき、ステップ(6)の$g$交換はD3Q19に基づくため総データ量は他ステップの6倍になる。 **Figure 3: D3Q19離散化(赤矢印がD3Q7の部分集合)** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig03-d3q19.png]] (Fig. 3. D3Q19. Red arrows are the subset for D3Q7. 19方向の速度分布(D3Q19)のうち、赤矢印で示された7方向(6面方向+中心)がD3Q7のサブセットとして使われる。) X・Y・Zの3次元それぞれで近傍と境界データを交換する必要があり、送信元・宛先アドレスは共にメモリ上で非連続になる。実務上最良の方法は、非連続データをデバイスメモリ上で連続バッファへpackし、デバイスからホストへ移動してからMPIで送信することである。 **Figure 5: X方向の境界交換** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig05-boundary-exchange-x.png]] (Fig. 5. Boundary exchange in X dimension of 3D LBM. 送信境界`e_snd`は$i = NX-2$、受信境界`e_rcv`は$i = NX-1$に固定され、$j$は$1$から$NY-2$、$k$は$1$から$NZ-2$まで変化する。図はドメインを$(j,k)$平面のスライスとして示し、斜線部が実際に交換される境界データを表す。) ## 新規性既存のGPUクラスタ向けデータ通信手法には次の限界があった。 - **ユーザーレベル最適化**(Bernaschiら, Jacobsenら等): MPI通信・CPU-GPU通信・CPU-GPU計算の重ね合わせをプログラマが手動で行う必要があり、複雑で煩雑な明示的データ処理を要する。 - **コンパイラディレクティブ方式**(ADSM、Mint、CGCM): 単一ノード内のCPU-GPU通信最適化に限定され、ノード間通信には別途分散メモリプログラミングモデルが必要。 - **アプリケーション固有ライブラリ**(Ma+、Maruyama+、Shimokawabe+): 新たなインターフェースの学習コストが生じ、ライブラリ内部実装ではやはりMPIとCUDAを併用する。 - **PGASモデル**(UPC、OpenSHMEM拡張等): 統一メモリ空間は提供するが、非連続GPUデータの送受信に十分なデータ型サポートが無く、pack/unpack関数をプログラマが供給する必要がある。 - **既存のGPU向けMPIライブラリ拡張**(CudaMPI、Jenkinsら[11]の GPUカーネルベース機構): CUDAメモリコピーベースやGPUカーネルベースのpack/unpack機構は提案されていたが、著者ら(Jenkinsら[11])の先行研究はMPIデータ型をベクトルへ変換してGPUカーネルを最大限活用する設計ではなく、かつ本論文で提案するパイプライン設計を欠くため、データpack/unpackを他の通信段と重ね合わせられない。本論文の新規性は、(1) MPIデータ型を汎用的にベクトルへ変換するアルゴリズム、(2) それをGPUカーネルで一括pack/unpackする機構、(3) pack/unpack・PCIe転送・RDMA転送を重ね合わせる非同期パイプラインの3点の組み合わせにより、非連続GPUデータについても標準MPI APIだけで高性能な通信を透過的に実現した点にある。 ## 実験設定 - **実験環境**: Oakleyスーパーコンピュータ(Ohio Supercomputing Center)。各ノード: デュアルIntel Xeon 6コアWestmere CPU、48GBホストメモリ、NVIDIA Tesla M2070 GPU、Mellanox QDR InfiniBand HCA。Red Hat Linux 6.1、OFED 1.5.1、CUDA Toolkit 4.1.28(ドライバ285.05.33)。設計をMVAPICH2 1.8に統合し、最大64GPUで実験。 - **データセット**: マイクロベンチマークでは16KB〜4MBの連続データ、および3D LBMを模した非連続データ4種(per-GPU: 128×128×128、256×128×128、256×256×128、256×256×256)。実アプリケーション評価では、512×512×512の全体ドメインサイズを8〜64GPUに分割。 - **比較対象**: 4つの実装バリアント。 - Default: 明示的なCPU-GPUデータ移動+ホストメモリ経由MPI通信(パック/アンパックはGPUカーネルで実施済みの高度にチューニング済み実装)。 - Contig: pack/unpack手動実装のまま、GPU-Aware MPIで連続データを直接送受信。 - Datatype: MPI派生データ型を定義し、GPU-Aware MPIで非連続データを直接送受信。 - Overlap: Datatypeに加え、通信と計算の重ね合わせを実施。 - **評価指標**: レイテンシ(μs)、正規化実行時間(Defaultを1として)、実行時間改善率(%)、MLUPS(Millions of Lattice site Updates per Second)、strong/weak scaling効率。 ## 実験結果 ### マイクロベンチマーク: 連続データのレイテンシ **Figure 8: GPUから遠隔ノードGPUへのレイテンシ** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig08-latency-gpu-to-gpu.png]] (Fig. 8. Latency from GPU to GPU on remote node. Defaultはデバイス→ホスト→MPI送受信→デバイスの3段、GPU-GPU MPIはGPU-Aware MPIで直接送受信。256KBデータサイズで最大42%のレイテンシ改善。改善はパイプライン設計によるPCIe/ネットワーク転送の重ね合わせに由来し、16KBが最適ブロックサイズで、16KB超のデータで自動的にパイプラインが起動する。) ### マイクロベンチマーク: 非連続データ(境界交換)のレイテンシ **Figure 9: X/Y/Z境界のレイテンシ** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig09-boundary-latency-xyz.png]] (Fig. 9. Boundary exchange latency for f. (a) X boundary MPI subarray Datatype. (b) Y boundary MPI vector datatype. (c) Z boundary MPI vector Datatype. GPU-GPU MPIはDefault比、X方向で16.1〜23.4%、Y方向で19.2〜23.6%、Z方向で19.1〜27.3%改善。X方向の境界サイズは$(NY-2)\times(NZ-2)$で64KB〜256KB。X方向のレイテンシが相対的に高いのは、デバイスメモリアクセスのストライドの違いによりYZ平面のpack/unpackでのcoalesced memory accessの効率が低下するため。) ### 3D LBM実行時間 **Figure 10: 8GPUでの正規化実行時間** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig10-normexec-8gpu.png]] (Fig. 10. Normalized execution time on 8 GPUs. per-GPUドメインサイズ128×128×128でDefault比、Contigは6.0%、Datatypeは13.4%、Overlapは18.2%の改善。256×256×256ではper-GPUデータサイズが大きいほど計算に占める時間割合が増えるため改善幅は縮小する。) **Figure 11: 64GPUでの正規化実行時間** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig11-normexec-64gpu.png]] (Fig. 11. Normalized execution time on 64 GPUs. 128×128×128でDefault比、Contigは10.1%、Datatypeは15.6%、Overlapは19.9%の改善。8GPU時より改善幅が大きいのは、per-GPU問題サイズが一定のままGPU数が増えると通信に費やす時間割合が増加するため。) ### 3D LBMスケーラビリティ **Figure 12: 実行時間改善率(512×512×512固定)** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig12-exectime-improvement.png]] (Fig. 12. Execution time improvement for 512×512×512 total domain size. 全体ドメインサイズを固定して8〜64GPUに分割した場合、GPU数が増えるほど改善率(Overlapで約20%まで)が増加する。) **Figure 13: strong scalingのMLUPS** ![[_attachments/GPU-Aware-MPI-on-RDMA-Enabled-Clusters--Design-Implementation-and-Evaluation/fig13-strong-scaling-mlups.png]] (Fig. 13. Strong scaling of 512×512×512 total domain size. 64GPUでのstrong scaling効率はDefault 76.4%、Contig 79.6%、Datatype 85.5%、Overlap 89.1%。全体ドメインサイズ固定でGPU数を増やすとper-GPUの通信割合が増えMLUPSが低下するが、計算との重ね合わせによりOverlapのscaling効率が明確に改善する。weak scaling効率はDefault 95.6%、Contig 96%、Datatype 96.6%、Overlap 97.2%と、3D LBMの近傍間P2P通信という性質により全実装で高い。) ### アプリケーションコードの簡略化(Fig. 6 / Fig. 7) Default実装(Fig. 6)は、Zディメンションのデータpack用GPUカーネル起動、ホスト-デバイス間`cudaMemcpy`、`MPI_Irecv`/`MPI_Isend`/`MPI_Waitall`によるホストメモリ経由の近傍通信、デバイス-ホスト間`cudaMemcpy`、Zディメンションのunpack用GPUカーネル起動を、X/Y/Z各方向・上下各面についてすべて明示的に記述する必要がある。 ```text DATA EXCHANGE FOR f IN DEFAULT IMPLEMENTATION // launch GPU kernel to pack data in z dimension pack_f_z() ... repeat pack data in y and x dimensions // copy data from host to device cudaMemcpy(topF_snd_h, topF_snd_d, ...) ... repeat for bot, north, south, east, and west // send and receive data with neighbours MPI_Irecv(topF_rcv_h, z_size, MPI_Float, top, ...) MPI_Isend(topF_snd_h, z_size, MPI_Float, top, ...) ... repeat for bot, north, south, east, and west MPI_Waitall() // copy data from device to host cudaMemcpy(topF_rcv_d, topF_rcv_h, ...) ... repeat for bot, north, south, east, and west // launch GPU kernel to unpack data in z dimension unpack_f_z() ... repeat unpack data in y and x dimension ``` (Fig. 6. Data Exchange for f in default implementation.) GPU-Aware MPI利用時(Fig. 7)は、`vector_z`のようなMPI派生データ型を定義しておけば、pack関数・データコピー・unpack関数がすべて不要になり、`MPI_Irecv`/`MPI_Isend`にデバイスポインタとデータ型を直接渡すだけでよい。 ```text DATA EXCHANGE FOR f USING DATATYPES // don't need pack functions and data copy from device to host topF_snd_d = f_6_d + gridID(i, j, NZ-2) + dcol topF_rcv_d = f_5_d + gridID(i, j, NZ-1) + dcol MPI_Irecv(topF_rcv_d, 1, vector_z, top, ...) MPI_Isend(topF_snd_d, 1, vector_z, top, ...) ... repeat for bot, north, south, east, and west MPI_Waitall() // don't need unpack functions and data copy from host to device ``` (Fig. 7. Data exchange for f using GPU-aware MPI.) ## 考察 - 非連続データのGPU-Aware MPI最適化(Datatype)は、連続データのみの最適化(Contig)より一貫して大きな改善をもたらす(8GPUで6.0%→13.4%、64GPUで10.1%→15.6%)。これは3D LBMの通信の大半が非連続な境界データ交換であるためである。 - 通信と計算の重ね合わせ(Overlap)は、MPI非ブロッキングプリミティブによる従来の重ね合わせと同様の枠組みで実装できる。GPU-Aware MPIがデータ型のpack/unpackとCPU-GPU移動を隠蔽するため、プログラマはMPI通信とGPU計算の重ね合わせだけを考慮すればよく、実装が大幅に簡略化される(datatypeに変位を伴う場合はZ方向で上下2つのデータ型に分割することで部分的な重ね合わせを実現)。 - strong scalingでの改善がweak scalingでの改善より顕著なのは、全体ドメインサイズ固定でGPU数を増やすとper-GPUあたりの通信割合が増加し、通信最適化の効果が相対的に大きくなるためである。 ## 強み / 弱点・課題 **Strengths** - 標準MPI APIのみで完結し、新たなAPIを追加しない設計方針により、既存MPIアプリケーションへの移植コストを最小化している。 - ベクトル化とパイプラインという2つの独立した最適化を組み合わせ、連続データ・非連続データの両方、かつGPUDirect RDMA利用/非利用の両環境に適用可能な一般性を持つ。 - 実運用アプリケーション(3D LBM)での評価を通じ、マイクロベンチマークの改善がアプリケーションレベル性能改善(最大19.9%)に結びつくことを示した。 **Weaknesses/Limitations** - 評価環境ではGPUDirect RDMAが利用できておらず(著者ら自身が明記)、ホストメモリ経由のステージングを前提とした設計である。GPUDirect RDMA環境での定量評価は将来課題として残されている。 - ブロックサイズ(16KBが最適との報告)はハードウェア構成(GPU・PCIe・InfiniBandの組み合わせ)に依存するチューニングパラメータであり、著者らもマイクロベンチマークによるインストール時チューニングを前提としている。異なるハードウェア世代への一般化は自明ではない。 - 評価はNVIDIA Tesla M2070・QDR InfiniBand・MVAPICH2 1.8という2013年当時の環境に限定されており、後継技術(GPUDirect RDMA本格導入以降の世代)での再評価は本論文の範囲外である。