# GPU起動型ネットワーキング ## 定義 GPU起動型ネットワーキング(GPU-Initiated Networking、デバイス起動通信)は、CPU の介在なしに GPU カーネル自身がネットワーク越しの RDMA 操作を発行・完了確認する通信モデルである。従来の GPU 通信はホスト起動モデルに従い、CPU が全ての通信操作をオーケストレーションし、通信呼び出しごとに個別のカーネル起動と明示的なホスト・デバイス同期を要求する。デバイス起動通信はこれを排し、GPU スレッドが計算と通信を1つのカーネル内で融合できるようにする。NVIDIA NVSHMEM がこのモデルを GPUDirect Async Kernel-Initiated(GDAKI)機能として確立し、NCCL は 2.28 の Device API で GIN として自身のエコシステムに統合した。(Source: [[@2025__arXiv__GPU-Initiated Networking for NCCL]]) ## 横断的知見 - **NVSHMEM とNCCLは同じ GDAKI 技術基盤を異なるランタイム哲学で統合する**: NVSHMEM は OpenSHMEM 由来の PGAS(Partitioned Global Address Space)モデルでポインタベースアドレシングとメモリアトミック同期を採用し、スタンドアロンランタイムとして存在する。NCCL GIN はウィンドウベースアドレシングとシグナルアトミック同期(ID ベース)を採用し、NCCL の階層的通信子・耐障害性・弾力性という本番インフラストラクチャに統合される。DeepEP の NVSHMEM→NCCL GIN 移行(zero-byte put + アトミックシグナルで release-acquire 意味論を模擬)は、この2つの設計思想の翻訳作業として具体化した。(Source: [[@2025__arXiv__GPU-Initiated Networking for NCCL]]) - **性能は同等でも「エコシステム統一」という別の価値軸が生まれる**: GIN の GDAKI バックエンドは NVSHMEM IBRC に匹敵する低レイテンシ(16.7µs vs 16.0µs)を達成するが、著者らはこれを主目的としていない。既存の集団通信(all-reduce 等)とデバイス起動 one-sided 通信を単一ランタイム(NCCL)内で使い分けられる点、複数の通信ランタイムを並行運用する運用複雑性を排除できる点を主要な価値として位置づける。(Source: [[@2025__arXiv__GPU-Initiated Networking for NCCL]]) - **GPU起動型ネットワーキングは孤立した最新技術ではなく、ノード間通信の5段階分類における最終段階「Device Native」として体系上の位置を持つ**: [[@2026__CSUR__The Landscape of GPU-Centric Communication]] はノード間通信を Host Native → Pinned Host Native → GPU RDMA → GPU-Triggered → Device Native の5型に分類し、GIN の GDAKI バックエンドや NVSHMEM の IBGDA はいずれも最終段階(Device Native、API・メッセージ登録・トリガー・データパスの全てが GPU 側で完結)に位置づけられる。この taxonomy は、GIN が単発の性能改善ではなく、GPUDirect(1.0→2.0→RDMA→Async)という10年以上続くベンダー機構の進化系列(Figure 3のタイムライン)の到達点であることを示す。(Source: [[@2026__CSUR__The Landscape of GPU-Centric Communication]], [[@2025__arXiv__GPU-Initiated Networking for NCCL]]) - **ROC_SHMEM は GDAKI/GIN が抱える GPU-NIC メモリ一貫性問題を、AMD 独自のアプローチで先に解決していた**: [[@2026__CSUR__The Landscape of GPU-Centric Communication]] は、NVSHMEM の IBGDA がカーネル境界を越えた GPU-NIC メモリ一貫性を保証できない(CPU へのコールバックに依存する)のに対し、AMD の ROC_SHMEM はカーネル内で完全にこの一貫性を保証する設計を既に持つと指摘する。GPU起動型ネットワーキングの「完全自律」という目標に対し、NVIDIA(NCCL GIN/NVSHMEM IBGDA)と AMD(ROC_SHMEM)は異なる到達度にあり、ベンダーごとに残存する CPU 依存の性質が異なることを示す横断的知見である。(Source: [[@2026__CSUR__The Landscape of GPU-Centric Communication]]) - **NVSHMEMのPGASモデルは、GINが理論として整理する「Device Native」通信を実務コード水準で先に具体化していた**: 「AI Systems Performance Engineering」第12章は、`nvshmem_float_p`(put)+`nvshmem_quiet()`(完了保証)+`nvshmem_int_p`(フラグ設定)という send-and-signal パターンと、受信側の `nvshmem_int_wait_until()` によるフラグ待機という実装レベルの一方向通信コードを示す。GINの理論的位置づけ([[@2025__arXiv__GPU-Initiated Networking for NCCL]] の PGAS ランタイム哲学の記述)と、この章の具体的なAPI呼び出し列は整合しており、実務書がNVSHMEMのPGASモデルを「点対点同期を優先しバリアを避けるべき」という運用上の注意点まで含めて裏づける。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 12 Dynamic Scheduling, CUDA Graphs, and Device-Initiated Kernel Orchestration]], [[@2025__arXiv__GPU-Initiated Networking for NCCL]]) ## 未解決の問い - NCCL GIN のウィンドウ登録は非対称バッファサイズを設計上許容するが、NCCL 2.28 実装は対称サイズを強制する。prefill/decode で異なるバッファサイズを要求する disaggregated serving への完全対応はいつ、どのように実現されるか。 - GIN は「ワークキューエントリのバッチ化」「複数操作にわたるドアベルコスト償却」を将来最適化として挙げているが、これらは GDAKI のレイテンシ・スループット特性をどの程度改善しうるか。 - RoCE 環境や H100 以外の GPU 世代(Blackwell 系等)での GIN の性能特性は本論文からは読み取れない。異なるハードウェア世代・ネットワーク技術でも GDAKI/Proxy の性能比が維持されるか。 - GIN の one-sided 通信は GPU スレッド間の順序を保証しない(ユーザーが CUDA 同期プリミティブで管理する責任を負う)。この責任分界は、GIN を直接使うアプリケーション開発者にとってどの程度のバグリスクを生むか。 - パーシステントカーネルによる完全な GPU 自律実行は、cooperative launch によるハードウェアオーバーサブスクリプション不能という副作用を伴う。GPU起動型ネットワーキングの普及とともにこの制約はどう緩和されるか(コンパイラ/ランタイムによる自動分解か、ベンダーによる新API提供か)。([[@2026__CSUR__The Landscape of GPU-Centric Communication]]) ## 関連 - 概念: [[RDMA]] / [[集合通信]] / [[Mixture-of-Experts]] - エンティティ: [[NCCL]] / [[DeepEP]] / [[NVIDIA]] / [[NVSHMEM]] - ソース: [[@2025__arXiv__GPU-Initiated Networking for NCCL]] / [[@2026__CSUR__The Landscape of GPU-Centric Communication]] / [[@2025__OReilly__AI Systems Performance Engineering - Chapter 12 Dynamic Scheduling, CUDA Graphs, and Device-Initiated Kernel Orchestration]] ## 出典 - [[@2025__arXiv__GPU-Initiated Networking for NCCL]] - [[@2026__CSUR__The Landscape of GPU-Centric Communication]](ノード間通信5分類の最終段階としてのDevice Nativeの体系的位置づけ、ROC_SHMEMのカーネル内GPU-NIC一貫性保証)