> [!abstract] 概要(arXiv abstract の日本語訳) > 近年、GPU はその並列性と高いメモリ帯域幅により、HPC・ML アプリケーションで好まれるアクセラレータとなった。GPU は計算を高速化する一方、GPU 間通信はスケーラビリティのボトルネックを生みうる。特にノードあたり・クラスタあたりの GPU 数が増えるほどこの問題は深刻になる。従来、マルチ GPU 通信は CPU が管理していたが、GPU 中心通信(GPU-centric communication)の進展は、CPU の関与を減らし GPU に通信タスクのより大きな自律性を与え、マルチ GPU 通信と計算のミスマッチに対処することで、この CPU の優位性に挑戦している。 > 本稿はベンダー機構とユーザーレベルのライブラリサポートに焦点を当てて GPU 中心通信のランドスケープを提供する。本稿の目的は、この分野の複雑さと多様な選択肢を明確化し、用語を定義し、ノード内・ノード間の既存アプローチを分類することである。本稿はマルチ GPU 実行における通信・メモリ管理のためのベンダー提供機構を議論し、主要な通信ライブラリとその利点・課題・性能に関する知見をレビューする。続いて、主要な研究パラダイム・将来展望・未解決の研究課題を探る。ソフトウェア・ハードウェアスタック全体にわたる GPU 中心通信技術を広範に記述することで、本稿は研究者・プログラマ・エンジニア・ライブラリ設計者に、マルチ GPU システムを最大限活用する方法についての知見を提供する。 ## 論文情報 - **タイトル**: The Landscape of GPU-Centric Communication - **著者**: [[Didem Unat]](責任著者)・[[Ilyas Turimbetov]]・[[Mohammed Issa]]・[[Dogan Sagbili]]・[[Ismayil Ismayilov]]([[Koç University]])、[[Flavio Vella]]([[University of Trento]])、[[Daniele De Sensi]]([[Sapienza University of Rome]]) - **媒体**: ACM Computing Surveys (CSUR), Vol. 58, No. 12, Article 322 - **発表**: 2026 年 6 月(日付不明、DOI 発行 2026 年) - **DOI**: 10.1145/3813799 - **ページ数**: 36 - **ライセンス**: CC-BY-NC-ND 4.0(本文図表は個別に CC-BY と注記) - **資金**: European Research Council(ERC)Horizon 2020 grant No 949587、Horizon Europe grant 101175702(NET4EXA)、Sapienza University Grants ADAGIO・D2QNeT ## 概要 本論文は、CPU がマルチ GPU 通信を管理する従来モデルから、GPU 自身が通信をより自律的に制御する GPU 中心通信(GPU-centric communication)への移行を、ベンダー提供機構とユーザーレベル通信ライブラリの両面から包括的に整理したサーベイである。ノード内・ノード間の通信をそれぞれ 4 型・5 型に分類する統一的な用語体系を提示し、NVIDIA・AMD・Intel の3ベンダーにわたるメモリ管理機構・GPUDirect 技術・インターコネクトを解説した上で、GPU-Aware MPI・GPU 中心コレクティブライブラリ(GPUCCL: NCCL/RCCL/oneCCL)・GPU 中心 OpenSHMEM(NVSHMEM/ROC_SHMEM/Intel SHMEM)の3系統のライブラリを比較する。最後に CPU-free networking・コレクティブアルゴリズム設計・デバッグ/プロファイリングツール・圧縮加速通信という5つの研究パラダイムを展望する。 **Figure 4: RDMA指向ソフトウェアスタック** ![[_attachments/2026_Unknown_The_Landscape_GPU_Centric_Communication/fig04-rdma-stack.png]] (Figure 4. アプリケーション/フレームワーク層(PyTorch・JAX・TensorFlow、PETSc・Trilinos・GROMACS)から、通信 API/ライブラリ層(MPI・NCCL/*CCL・SHMEM 系)、通信ミドルウェア/トランスポート抽象層(UCX・libfabric)、RDMA ソフトウェアスタック(RDMA-core・libibverbs・provider libraries)、カーネル/デバイスドライバ、インターコネクト/ネットワークファブリック(InfiniBand・RoCE・EFA・iWARP・Slingshot)までの全層を1枚で示す構図。本論文全体の見取り図として機能する。Source: Figure is available under CC-BY [148].) ## 問題設定 論文は GPU-centric communication を「マルチ GPU 実行のクリティカルパスにおける CPU の関与を減らす機構群」と広く定義し、その多様性を解消するために通信を構成する要素に基づいて分類する。 **ノード内通信(4型、Table 1・Figure 1)**: API(通信呼び出しがホスト/デバイスどちらで発行されるか)とデータパス(実際のデータ移動がホスト経由か直接か)の2軸で、Host Native(ホスト API・ホスト経由、P2P 無効)・Host-Controlled(ホスト API・デバイス直接、P2P 有効)・Device Native(デバイス API・デバイス直接)・Host Fallback(デバイス API だが P2P 無効時にホスト経由へフォールバック)の4型に分類する。 **ノード間通信(5型、Table 2・Figure 2)**: API・データパスに加え、Register/construct messages(NIC 上でのパケット構築・登録を誰が行うか)・Trigger communication(NIC のドアベルを誰が鳴らすか)の2軸を追加した4軸で、Host Native → Pinned Host Native → GPU RDMA → GPU-Triggered → Device Native の順に、より多くの構成要素が GPU 側へ移っていく5型に分類する。この過程でデータコピー段数は減少し、GPU の自律性は増大する。 **Figure 1: ノード内通信のデータパスとAPI呼び出し** ![[_attachments/2026_Unknown_The_Landscape_GPU_Centric_Communication/fig01-intranode-datapaths.png]] (Figure 1. Host Native(①)・Host-Controlled(②)・Device Native(③)・Host Fallback(④)の4型を、CPU/GPU間のデータパス(緑矢印)とAPI呼び出し位置(緑丸)で図示。③のみCPUが灰色(関与なし)で描かれる。Source: Figure is available under CC-BY [148].) **Figure 2: ノード間通信のデータ・制御パス** ![[_attachments/2026_Unknown_The_Landscape_GPU_Centric_Communication/fig02-internode-datapaths.png]] (Figure 2. Host Native(①)からDevice Native(⑤)へ至る5型を、データパス(緑)・メッセージ構築(青)・トリガー(紫T)・API呼び出し(緑丸)の4要素で図示。⑤のみCPUが灰色になり、NICへの全操作がGPU側で完結する。Source: Figure is available under CC-BY [148].) ## 提案手法:ベンダー機構と通信ライブラリの整理体系 論文はSection 3でベンダー提供の下位機構、Section 4でその上に構築される上位ライブラリを解説する2層構成を取る。 ### ベンダー機構(Section 3) - **メモリ管理機構**: Page-Locked/Pinned Memory(`cudaMallocHost()`によるゼロコピーメモリ)・UVA(Unified Virtual Addressing、CUDA 4.0〜、ポインタから物理位置を推論可能に)・IPC(プロセス間でのデバイスバッファ共有)・UVM(Unified Virtual Memory、`cudaMallocManaged()`によるページフォルト駆動の自動マイグレーションとメモリオーバーサブスクリプション)。 - **GPUDirect 技術**: GPUDirect 1.0(NIC、GPU-NIC間の共有ピンメモリ)・GPUDirect 2.0/P2P(CUDA 4.0〜、GPU間の直接load/store)・GPUDirect RDMA(CUDA 5.0〜、NICがGPUメモリのBARリージョンを直接読み書き)・GPUDirect Async(CUDA 8.0〜、CPUが事前登録したメッセージをGPUがトリガー)。GPUNetIO(NVIDIA DOCA)はDOCA v2.7以降、カーネル境界を越えた任意タイミングでのRDMA送受信をGPUに許可する。 - **モダンなGPU中心インターコネクト**: NVLink/NVSwitch(第5世代で1.8TB/s、最大576GPU接続)・AMD Infinity Fabric/xGMI(8GPUで896GB/s、フラットメッシュでスイッチファブリック非搭載)・UALink(オープン規格の新興コンソーシアム)・Intel Xe-Link(Aurora向け、最大8-way全対全)。 **Figure 3: NVIDIA技術のタイムライン** ![[_attachments/2026_Unknown_The_Landscape_GPU_Centric_Communication/fig03-nvidia-timeline.png]] (Figure 3. CUDA 2.2(2006)からCUDA 13.0/Rubin(2026)まで、メモリ管理機構(緑)・GPUDirect技術(青)・ハードウェア(紫)・ユーザーレベルライブラリ(黄)・開発者ライブラリ(橙)の5系列をGPUアーキテクチャ世代(Tesla〜Rubin)に対応づけて年表化。GPU-aware MPI→NCCL→NVSHMEMの登場順序と、各技術が依存する下位機構(点線矢印、例: GPUDirect RDMA→GDRCopy)が視覚化されている。Source: Figure is available under CC-BY [148].) ### 通信ライブラリ(Section 4) - **GPU-Aware MPI**: MVAPICH2(GPU-awareness統合の先駆、GDRCopy・GPUDirect RDMA/Async対応)・Open MPI(UCX経由でCUDA/ROCm対応、UCC統合)・MPICH(CH4層でCUDA/ROCm対応)。 - **GPUCCL(GPU中心コレクティブ)**: NCCL・RCCL・oneCCLをTable 3で7軸(コレクティブプリミティブ・APIモデル・実行エンジン・アクセラレータ統合・イントラノードファブリック・トランスポートバックエンド・進行/開始・トポロジ認識)で比較。NCCLは完全にGPU自律的な実行エンジンを持つが、oneCCLはCPUワーカースレッドが集団操作の進行を担う点で質的に異なる。NCCL 2.28ではデバイス側API(Device API)が導入され、GPUカーネルから直接データ移動を開始できるようになった。 - **GPU中心OpenSHMEM(GPUSHMEM)**: NVSHMEM(PGAS、symmetric heap、IBGDAトランスポート)・ROC_SHMEM(GPU-IBとReverse Offloadの2バックエンド、カーネル内でGPU-NICメモリ一貫性を保証する唯一の実装)・Intel SHMEM(SYCLベース、ホストプロキシスレッド経由のノード間通信)。 Listing 1〜3(GPU-Aware MPI/NCCL/デバイス側NVSHMEMの帯域幅ベンチマーク疑似コード)は、MPI/NCCLの二者間send/recv型セマンティクスと、NVSHMEMのone-sided put/get型セマンティクスの意味論的差異を具体的に示す。 **Figure 5: 通信ライブラリ間の相互作用** ![[_attachments/2026_Unknown_The_Landscape_GPU_Centric_Communication/fig05-library-interactions.png]] (Figure 5. Application層の下にMPI・GPUSHMEM・GPUCCLが並び、MPIはUCC経由・GPUSHMEMは直接・GPUCCLは両経路でUCXへ降りる階層を図示。最下層はlibfabric/libibverbs/sockets等の共通トランスポートとNVLink/InfinityFabric/NVSwitch/UALink等のハードウェアに収束する。Source: Figure is available under CC-BY [148].) ## 新規性 先行研究の多くはNVIDIA一社のエコシステムまたは単一ライブラリ(NCCLやMPI)に閉じた解説であり、GPU中心通信という用語自体もベンダー・論文間で不統一だった。本論文は(1)API・データパス・メッセージ構築・トリガーの4軸によるノード内/ノード間通信の統一的分類体系を初めて提示し、(2)NVIDIA中心の記述に留まらずAMD・Intelの対応機構を可能な限り明示的に併記して3ベンダーの均衡した比較を実現し、(3)ベンダー機構(Section 3)とユーザーレベルライブラリ(Section 4)を明確に分離しつつ両者の依存関係(Figure 4・5)を接続することで、GPU中心通信の複雑な設計空間を一貫した用語で読み解けるようにした点で先行研究を超える。 ## 実験設定 本論文は独自の実験を行わないサーベイであり、既存文献が報告した性能評価を引用・整理する形で知見を提供する。引用される主要な実測環境には、CG(Conjugate Gradient)ソルバをMPI/NCCL・RCCL/NVSHMEMでNVIDIA・AMD双方のGPU上に実装した比較実験、GROMACS上でのNVSHMEM評価、3台のスーパーコンピュータ上でのNCCL/RCCL対GPU-Aware MPI比較、Frontierスーパーコンピュータ上でのCray MPICH対RCCL比較などが含まれる。 ## 実験結果 - CG ソルバでCPU-GPU間ストリーム同期を回避することで5〜15%の性能改善が確認され、NVIDIAシステムでは小メッセージAllReduceにNCCL、点対点通信にMPIを使うことが推奨される一方、AMD GPUではベンダーソリューションの未成熟さからMPIがNCCL相当ライブラリを上回った。 - GROMACSでのNVSHMEM評価はMPI比最大2倍の性能改善を示し、PGASモデルがCPUをクリティカルパスから排除できることに起因すると分析される。ただし最近のNCCL Device APIの登場により、point-to-pointとall-to-allの両方でNVSHMEMと同等の性能が得られるようになったと報告される。 - 3スーパーコンピュータでの比較では、点対点操作でMPIが優位、コレクティブ操作で*CCLが優位という傾向が示されたが、システムやノード数・ベクトルサイズに強く依存する。高ノード数でのRCCLの不安定性がFrontierでのCray MPICH比較でも指摘されている。 - Open MPI・Cray MPICHはaccelerator-offloadコンポーネントを明示的に設定しない限り、`MPI_Allreduce`等の縮約系コレクティブでホストステージング(GPU→CPU→GPU)にフォールバックし性能が大きく劣化する。MFDn核構成相互作用コードの最適化では、NCCL/CUDAアプローチがベースラインMPI/OpenACC比最大4.9倍の高速化を達成した。 ## 考察 Section 5では5つの研究パラダイムが議論される。 - **CPU からの移行(§5.1)**: パーシステントカーネルとGPU中心通信の組み合わせは強スケーリング環境でのレイテンシ障壁を解消できる一方、cooperative launchによるハードウェアオーバーサブスクリプション不能・レジスタ/共有メモリ圧迫によるオキュパンシー低下という新たな課題を生む。NVSHMEM/ROC_SHMEMの症状メモリ割り当て要件が既存ランタイムへの統合を難しくすることも指摘される。 - **UCXという経路(§5.2)**: UCXはCUDAポインタを検出して`rc_mlx5`や`cuda_ipc`など適切なトランスポートを動的選択でき、MPI実装がベンダー固有APIを重複実装せずにGPU対応を得る手段として今後さらに重要になると予測される。ただしUCXは常にホスト側で動作し通信カーネルを持たないという限界がある。 - **CPU-free networking(§5.3)**: GGAS・GPUrdma・GPUNet等の系譜は、GPU側でInfiniBand Verbsを完全実装する試みだが、単一スレッドでのワークリクエスト生成性能の低さやパーシステントカーネルとのGPU-NIC一貫性問題により、実用化には追加のハードウェア支援(GPU-TN・ComP-Net)を要する。 - **より広いGPU自律性(§5.4)**: GPUfs・BaM(GPUによる直接ストレージアクセス)・CPU-free device-side task graph実行など、通信以外の領域(ファイルI/O・ストレージ)でもCPUからGPUへの権限移譲が進んでいる。ただしデバッグ・プロファイリング・監視・ソフトウェア開発はCPUの成熟したツールエコシステムに依然として依存する。 - **コレクティブアルゴリズム設計(§5.5)**: GPU間帯域幅の非対称性(同ノード内で最大4倍、ノード内外で最大10倍)がコレクティブアルゴリズム設計を複雑化する。線形プログラミングによる最適解探索はNP困難でありノード数増加に対して指数的にスケールする(128ノードで最大11時間)ため、MSCCL等の高レベル言語によるコンパイルアプローチが実用上重要になる。 - **デバッグ/プロファイリング/ベンチマーキング(§5.6)**: NSight Systemsはデバイス側ネイティブ通信を捕捉できず、ComScribe・Snoopie・ucTrace等の専用ツールが補完する。マルチGPUコンテキストでのレースハザード検知ツールは本論文執筆時点で存在せず(Compute Sanitizer RacecheckとHiRaceはいずれも単一GPUコンテキスト限定)、重要な空白領域として指摘される。 - **圧縮加速GPU通信(§5.7)**: 誤差有界損失圧縮は通信量を削減するが、Decompression-Operation-Compression(DOC)ワークフローの伸長・再圧縮オーバーヘッドが課題となる。計算を圧縮領域のまま直接実行する同型圧縮(homomorphic compression)がCPU・GPU双方で最近実装され、DOCオーバーヘッドとデータ量の両方を同時に削減する最有望アプローチとして位置づけられる。 ## 強み / 弱点・課題 **強み**: - API・データパス・メッセージ構築・トリガーの4軸taxonomyにより、ベンダー間・世代間で乱立する用語を統一的に整理した点 - NVIDIA・AMD・Intelの3ベンダーを可能な限り均等に扱い、単一ベンダー視点のサーベイに陥っていない点 - ベンダー機構(Section 3)とライブラリ(Section 4)の対応関係を明示し、下位機構の理解なしにライブラリの設計判断を評価できない構造を示した点 **弱点・課題**: - サーベイであるため独自のベンダーニュートラルな実測は行われず、性能比較は既存文献の引用に留まる(実験設定・実験結果の節が示すとおり、比較条件は論文間で統一されていない) - Intel oneCCL/Intel SHMEMの記述は、NVIDIA/AMD系に比べて実測エビデンスが薄く、公開情報への依存度が高い - マルチGPUコンテキストでのレースハザード検知ツールの欠如を指摘する一方、その解決に向けた具体的な設計方針は提示されない