CUDA-aware MPIとGPUDirect

ホスト経由コピーという隠れたコスト

マルチGPUのMPIプログラムで最初に書かれがちなのは、GPU上の計算結果を一度ホストメモリへcudaMemcpyで退避し、その host ポインタをMPIに渡し、受信側で逆にGPUへ書き戻す実装です。これは動きますが、1回の論理的な「GPU間送信」の裏で、デバイスからホストへのコピー、ネットワーク転送、ホストからデバイスへのコピー、という最低3段の転送が直列に走ります。GPUの演算がいくら速くても、この staging（ステージング）コピーが通信レイテンシと実効帯域を律速します。

CUDA-aware MPIとは何をする実装か

CUDA-aware MPIは、MPI_Send / MPI_Recv / MPI_Allreduce などに渡されたバッファがデバイスポインタかホストポインタかを実行時に判別し、経路を自動で切り替えるMPI実装です。判別には統合仮想アドレス（UVA, Unified Virtual Addressing）を使い、ポインタの属性を問い合わせて「これはGPUメモリだ」と認識します。デバイスバッファだと分かれば、MPI実装は内部で以下のいずれか最適な経路を選びます。

• 同一ノード内のGPU間なら GPUDirect P2P による直接コピー
• ノード間なら GPUDirect RDMA でNICからGPUメモリへ直接DMA
• 直接経路が使えない環境ではフォールバックとして内部で自動 staging（利用者コードは変えずに従来経路へ退避）

重要なのは、利用者コードはMPI_Send(device_ptr, ...)と書くだけで、ホストへのコピーコードが不要になる点です。最適化の判断と経路切り替えはMPI実装（Open MPI + UCX、MVAPICH2-GDRなど）の内部に隠蔽されます。

// 素朴な実装: 手動でホスト staging（3段転送）
cudaMemcpy(host_buf, dev_buf, n, cudaMemcpyDeviceToHost);
MPI_Send(host_buf, n, MPI_BYTE, dst, tag, comm);
// 受信側で逆コピーが必要

// CUDA-aware MPI: デバイスポインタを直接渡す（実装が最適経路を選択）
MPI_Send(dev_buf, n, MPI_BYTE, dst, tag, comm);

GPUDirect P2PとGPUDirect RDMAは別物

「GPUDirect」は複数の技術の総称で、CUDA-aware MPIが使い分ける2つの主役を混同しないことが理解の鍵です。両者は解決している問題も物理経路も異なります。

P2P（Peer-to-Peer）は1ノード内の話で、GPU AがGPU Bのメモリを自分のアドレス空間の延長として直接読み書きできるようにし、間にホストメモリを挟みません。NVLinkがあればその高帯域リンク上を、なければPCIe上を通ります。一方RDMA（Remote Direct Memory Access）はノード間の話で、NICがCPUを起こさずGPUメモリへ直接DMAする仕組みです。CUDA-aware MPIは送受信ペアの位置関係を見て、ノード内ならP2P、ノード外ならRDMAへと自動で振り分けます。ハードウェア側の階層構造（ノード内NVLink／ノード間InfiniBand）についてはGPUクラスタの相互接続（NVLink・InfiniBand）で詳しく扱っています。

マルチGPUノードのトポロジとアフィニティ

1ノードに複数のGPUと複数のNICが載る構成では、どのGPUがどのNIC・どのPCIeスイッチにぶら下がっているかが実効性能を決めます。GPUDirect RDMAが本来の帯域を出すのは、送信に使うGPUと、そのトラフィックを担うNICが同じPCIeスイッチ配下にあり、DMAがCPUのPCIeルートを跨がないときです。GPUとNICが別スイッチ・別ソケット配下だと、転送はCPU内部のインターコネクトを経由してレイテンシと帯域が悪化します。

このトポロジはnvidia-smi topo -mで行列として確認できます。出力に現れる接続種別を読み解けるかが実務の分かれ目です。

nvidia-smi topo -m の接続種別（近い順）

  NV#  : NVLinkで直結（数字はリンク本数）。最速
  PIX  : 同一PCIeスイッチ配下。GPUDirect RDMAが効きやすい
  PXB  : 複数PCIeスイッチをまたぐ（同一ホストブリッジ内）
  PHB  : PCIeホストブリッジ経由。CPUのPCIeルートを通る
  SYS  : ソケット間インターコネクト（QPI/UPI等）をまたぐ。最も遠い

MPIランクをGPUへ割り当てる際は、この距離を踏まえて「各ランクが使うGPUと、そのランクの通信を担うNICを同一スイッチ配下（PIX）に揃える」ようにアフィニティを設定します。これはノード間集団通信（AllReduce等）のアルゴリズム段数とは独立した、物理経路レベルの最適化です。集団通信そのもののアルゴリズムと計算量についてはMPIと集団通信を参照してください。

性能モデル上の意味

staging を消すことの効果は、通信時間の内訳で捉えると明確です。1回の転送時間はおおむね「固定レイテンシ」と「転送量 割る 帯域」の和で近似できますが、ホスト経由実装ではこれに加えてデバイス・ホスト間コピー2回分の時間と、CPU介入によるレイテンシが上乗せされます。

横にスクロール

転送時間の内訳（概念）

  ホスト staging あり:
    latency + copy(D2H) + (size / net_bw) + copy(H2D) + CPU介入
  GPUDirect（P2P/RDMA）:
    latency + (size / link_or_net_bw)
    （D2H/H2Dの2コピーとCPU介入の項が消える）

小さいメッセージを高頻度で送る通信パターン（領域分割の袖領域交換、勾配同期の細切れ送信）では、消えるのは主にコピー回数とCPU介入によるレイテンシで、ここが実効スループットを底上げします。大きいメッセージでは、staging の2コピー分の帯域消費が丸ごと消えることが効きます。いずれにせよ、この改善は演算律速か通信律速かの境界を通信側へ押し戻すため、Rooflineモデルや強スケーリング・弱スケーリング解析で通信コストを見積もる際の前提条件を変えます。

まとめ

• 素朴なMPI-GPUコードはデバイスからホストへの staging コピーが2回入り、これが通信レイテンシと実効帯域を律速する。
• CUDA-aware MPIはUVAでデバイスポインタを判別し、経路をGPUDirect P2P（ノード内）／GPUDirect RDMA（ノード間）へ自動切り替えして staging を消す。
• GPUDirect P2PとRDMAは別機構で、前者はGPU間のNVLink/PCIe直接コピー、後者はNICによるGPUメモリへの直接DMAである。
• 実効性能はGPU・NIC・スイッチのPCIe/NVLinkトポロジに支配され、NICと同一PCIeスイッチ配下のGPUを選ぶアフィニティ設計と実測確認が欠かせない。