# GPUDirect RDMA ソース: https://docs.nvidia.com/cuda/gpudirect-rdma/ | 発行: [[NVIDIA]] Corporation | 最終更新: 2026-06-25 (CUDA 13.3時点) ## 概要 [[NVIDIA]] が公開する GPUDirect RDMA の公式リファレンスドキュメント。Kepler世代GPUとCUDA 5.0で導入された、GPUとサードパーティPCIeデバイス(NIC・映像取り込みデバイス・ストレージアダプタ等)間の直接データ交換技術の、設計原理・カーネルAPI・ユーザ空間API・移植ガイドを網羅する一次資料。[[GPUDirect RDMA]] の正典的な技術仕様書であり、[[RDMA]] 通信ライブラリ(MPI実装・NCCL等)の実装者が参照すべき基盤文書。 ## 動作原理と標準DMA転送との対比 PCI Expressデバイス間では物理アドレス空間が共有されており、各デバイスは最大6個のBAR(Base Address Register)領域を持つ。従来はCPUのMMUを介してBAR窓をユーザ/カーネルアドレス空間にマップするが、OSにはドライバ間でMMIO領域を交換する仕組みがないため、GPUDirect RDMA対応にはカーネルドライバの `get_user_pages()` 呼び出し付近への小規模な変更が必要になる。 標準DMA転送は (1) ユーザ空間プログラムが仮想アドレスとサイズを通信ライブラリへ渡す、(2) 通信ライブラリが転送準備(ピニング)をカーネルドライバに要求する、(3) カーネルドライバが仮想アドレス範囲を物理ページへ変換しDMAエンジンをプログラムする、(4) 転送完了後にアンピニングする、という流れを取る。GPUDirect RDMAはこれにユーザ空間CUDAライブラリとNVIDIAカーネルドライバの2要素を追加し、CPUメモリは既存の `get_user_pages()`/`put_page()` で、GPUメモリはNVIDIAカーネルドライバ関数(`nvidia_p2p_get_pages()` 等)でピニングを行う。 ## 設計上の要点 ### 遅延アンピニングとレジストレーションキャッシュ GPU BARへのメモリピニングはミリ秒オーダーで高コストであり、転送のたびにpin/unpinする素朴な実装は性能が出ない。高性能実装は転送完了後もメモリをピンしたまま保持し(遅延アンピニング)、同一領域の再利用を見込んでpin/unpin回数を削減する。通信ミドルウェアは典型的にレジストレーションキャッシュ(ピンダウンキャッシュ)を持ち、GPU BAR空間が64KB単位でマップされる特性を踏まえて64KB境界に丸めた領域単位でキャッシュを管理するのが資源効率上有利である。デアロケーションイベントの検知には (1) CUDAアロケーション/デアロケーションAPIの計装、または (2) `CU_POINTER_ATTRIBUTE_BUFFER_ID` によるタグチェック、の2方式があり、タグチェックはレイテンシ要求が低い場面向けで、通常はAPI計装によるコールバック方式が推奨される。 ### アンピンコールバックの制約 サードパーティデバイスドライバが `nvidia_p2p_get_pages()` でGPUページをピンする際は、アクセス取り消しが必要になった場合にNVIDIAドライバから同期的に呼ばれるコールバック関数の登録が必須。コールバックはミリ秒単位でブロックしてよいが、GPU側のアクションを待つとデッドロックしうるため注意が必要。コールバック内では `nvidia_p2p_put_pages()` ではなく `nvidia_p2p_free_page_table()` を呼んでページテーブルを解放する。コールバックは (1) ユーザ空間プログラムが対応GPUメモリを明示的に解放した場合(`cuMemFree`・`cuCtxDestroy`等)、(2) プロセスの早期終了、の2状況で発火する。 ### 同期とメモリ順序 GPUDirect RDMAはGPUの緩やかなメモリモデルの外側にサードパーティデバイスへ露出する独立したデータフロー経路を導入するため、CUDA APIとの整合が重要になる。CUDA APIメモリ操作がBARマッピングに対して呼び出し復帰前に可視になることを保証するには、`cuPointerSetAttribute()` に `CU_POINTER_ATTRIBUTE_SYNC_MEMOPS` を渡して該当バッファ全体を厳密同期モードにする必要がある(最適化を無効化するためオーバーヘッドを伴う)。また、GPUカーネルやコピー操作がGPUDirect RDMA経由の書き込みを観測する順序は、CPU起点のCUDA同期・ワーク投入APIによってのみ保証される。ネットワークRDMA書き込み完了とGPUカーネル実行が並行する場合、明示的な同期なしでは古いデータや部分書き込みをGPUカーネルが観測しうるデータ競合となる。 ### PCIトポロジとIOMMU GPUとサードパーティデバイスが同一のPCI Expressルートコンプレックス配下にあることが要件。PCIeスイッチのみの経路が最良の性能を示し、単一CPU/IOH経由は劣化するが動作し、CPU/IOH間をQPI/HT越しに跨ぐ経路は極めて性能劣化するか不安定になる(`lspci -t` でトポロジ確認可能)。GPUDirect RDMAは全PCIデバイスから見た物理アドレスの同一性に依存するため、1:1変換以外を行うIOMMUとは非互換であり、無効化またはパススルー設定が必要。 ## API層の構成 - **ユーザ空間CUDA API**: `cuPointerSetAttribute()`(`CU_POINTER_ATTRIBUTE_SYNC_MEMOPS` 設定)、`cuPointerGetAttribute()`(`CU_POINTER_ATTRIBUTE_P2P_TOKENS`[CUDA 6.0で非推奨]・`CU_POINTER_ATTRIBUTE_BUFFER_ID` 取得)、`cuPointerGetAttributes()`(複数属性の一括取得、クリティカルパスでの呼び出し推奨形態)。 - **カーネルAPI**(`nv-p2p.h` 宣言): `nvidia_p2p_get_pages()`(高コスト、GPUページのピン)、`nvidia_p2p_put_pages()`(ピン解放、コールバック内から呼んではならない)、`nvidia_p2p_free_page_table()`(コールバック内で呼ぶページテーブル解放)、`nvidia_p2p_dma_map_pages()`/`nvidia_p2p_dma_unmap_pages()`(I/Oアドレスが物理アドレスと異なるプラットフォームでの物理ページマッピング)。 ## Tegra移植とnv-p2p APIの非推奨化 GPUDirect RDMAはJetson AGX Xavier(CUDA 10.1〜)、DRIVE AGX Xavier Linux(CUDA 11.2〜)、Jetson Orin(CUDA 11.4〜)へ拡張されてきた。Tegraではアロケータを `cudaMalloc()` から `cudaHostAlloc()` へ切り替え、カーネルAPIシグネチャからは `p2p_token`/`va_space_token` 引数が落ちる(例: `nvidia_p2p_get_pages(u64 virtual_address, u64 length, ...)`)。マッピング長・ベースアドレスの整列要件もデスクトップの64KBに対しTegraでは4KBとなる。 CUDA 13.0以降、nv-p2p APIはBlackwell世代Tegra SoCから非対応となり、L4T上のOrinのみ継続サポート(ただし非推奨予定でCUDA 14.0で削除)。ThorおよびThor以降のTegraプラットフォームでは、Linuxアップストリームカーネルの `pin_user_pages()`(`FOLL_LONGTERM` 付き)・`unpin_user_page()`・`dma_map_sgtable()`/`dma_unmap_sgtable()` への移行が明示的に推奨されている。これは第三者I/Oデバイスドライバ(カメラ・センサ等)がnv-p2p依存からLinux標準のページピン留めAPIへ移行する、ベンダ固有APIからカーネル標準APIへの回帰(收斂)を示す事例といえる。 ## nvidia-peermemによるInfiniBand連携 [[nvidia-peermem]] は、Mellanox/[[NVIDIA]] InfiniBand系HCAにGPUビデオメモリへの直接P2P read/writeアクセスを提供するカーネルモジュール(CUDA 11.4で導入)。ConnectX-3 VPI以降のアダプタでInfiniBand・RoCE双方に対応する。旧来の `nv_peer_mem`(GitHubコミュニティ提供)と機能重複するため、`nvidia-peermem` 利用時は `nv_peer_mem` サービスの停止・アンインストールが必要。詳細は [[nvidia-peermem]] を参照。 ## 横断的知見 - レイテンシ/バンド幅の非対称性に対処する [[MVAPICH2]] のハイブリッド設計(→ [[@2013__ICPP__Efficient Inter-node MPI Communication using GPUDirect RDMA for InfiniBand Clusters with NVIDIA GPUs]])は、本ドキュメントが規定する「遅延アンピニング」「レジストレーションキャッシュ」の設計原則を具体的なMPIライブラリ実装として体現したものと読める。本ドキュメントはその設計原則の一次仕様にあたる。 - 「同期とメモリ順序」節が明示する制約(CPU起点のCUDA同期APIのみがGPUDirect RDMA操作の順序を保証する)は、GPUDirect RDMAを使うRDMA通信ライブラリ(NCCL・MPI等)が完了通知の伝搬に単純なポーリングではなくCUDAイベント/ストリーム同期を要する理由の一次的根拠になる。 - nv-p2p APIからLinuxアップストリームAPI(`pin_user_pages`等)への移行が推奨される流れは、ベンダ固有カーネルAPIがLinuxカーネル標準機構に収斂していく傾向の一例であり、将来のRDMA関連ソース取り込み時に同様の収斂パターン(ベンダAPI→Linux標準API)がないか確認する価値がある。 ## 出典 - NVIDIA Corporation, "GPUDirect RDMA," CUDA Documentation, version 13.3, last updated 2026-06-25. https://docs.nvidia.com/cuda/gpudirect-rdma/