> [!abstract] 概要(arXiv abstract の日本語訳) > GPU の集合通信は通常、帯域幅に最適化されているが、新興のワークロードの多くはレイテンシによって制限されるようになりつつある。長文脈・デコード主体の大規模言語モデル(LLM)推論はその代表例であり、大規模モデルの提供には複数の GPU が必要となり、多くの小さな集合通信操作がトークン生成のクリティカルパス上に直接位置する。そのため、わずか数マイクロ秒のオーバーヘッドでさえ性能とコストに影響しうる。本研究では、scale-up ネットワーク内における GPU 集合通信のハードウェアの Speed-of-Light(SoL)下限にどう近づくかを研究する。障壁(バリア)フリーの同期、および対称メモリとマルチキャストの効率的な利用を含む、ほぼ最適な設計のための重要な原理を特定する。NCCL のデバイス側 API を基盤として、カスタム集合通信カーネルを構築するための低レイテンシインタフェースを開発し、これらを用いて NCCL に新しい対称的な集合通信を実装する。マイクロベンチマークでは小〜中サイズのメッセージにおいて大幅なレイテンシ削減が示され、オーバーヘッドを絶対的な SoL 下限の 7% 以内まで抑える。実アプリケーションに統合すると、これらのカーネルは LLM 推論におけるトークン間レイテンシとスループットを改善し、cuSOLVERMp を高速化することで、AI 推論と従来の HPC ワークロードの双方に利益をもたらすことを実証する。 ## 論文情報 - タイトル: Every µs Matters: Achieving Near Speed-of-Light Latency in GPU Collectives - 著者: Siyuan Shen(ETH Zürich、NVIDIA インターンシップ中に大半を実施)、Anton Korzh・John Bachan・Arnav Goel・Ludwig Schneider・Pouya Kousha・Zhenhao He・Sylvain Jeaugey・Kamil Iskra・Nishank Chandawala・Jeff R. Hammond(NVIDIA Corporation、一部 NVIDIA Helsinki Oy)、Tiancheng Chen・Torsten Hoefler(ETH Zürich) - 媒体: arXiv(cs.DC)、投稿日 2026-07-17 - コード: https://github.com/ss16118/low-latency-nccl - 謝辞: FastTrackAI project(Singapore-ETH Centre、ETH Zürich と National Research Foundation Singapore の共同設立)、NRF Singapore・MDDI の AI Visiting Professorship、European Research Council(Project PSAP)の支援を受けた。 ## 概要 NVIDIA GB200 NVL72 の scale-up ネットワーク(単一 NVLink ドメイン)において、GPU 集合通信のハードウェア理論下限(Speed-of-Light、SoL)に迫る AllReduce 設計を研究する。既存実装が明示的な memory barrier に依存してレイテンシを浪費している点を突き止め、LL・sentinel 同期・双方向通信+double buffering・two-shot LL128 atomic という 4 つの barrier-free 技術を提案する。これらを NCCL のデバイス側 API の上に構築した低レイテンシ API 群として実装し、複数の新しい AllReduce カーネルを NCCL に組み込む。マイクロベンチマークで SoL 比 7% 以内のオーバーヘッドを達成し、vLLM 推論・cuSOLVERMp という実アプリケーションでの性能改善を実証する。 ## 問題設定 - **入力**: N 個の GPU(単一 NVLink ドメイン、scale-up ネットワーク)にまたがる AllReduce 操作。メッセージサイズは主に小〜中サイズ(長文脈・小バッチのデコード時に頻発する)。 - **出力**: 各 GPU が全 GPU の入力の縮約結果(reduction)を保持する状態。 - **前提条件**: 対称メモリ(symmetric memory)が `ncclCommWindowRegister()` で登録可能であること、GPU が NVLink で接続され Load/Store Accessible(LSA)であること、NCCL 2.28 以降のデバイス側通信 API が利用可能であること。 - **スコープ限定**: scale-out(ノード間・複数 NVLink ドメインをまたぐ)通信は対象外。理由は、(1) LLM 推論の並列グループは通常単一 scale-up ドメインに収まる、(2) マルチノードシステムは階層的集合通信でローカル/scale-out フェーズを分離するため scale-up 内最適化は上位最適化と補完的、(3) 新しい GPU システムは scale-up ドメインの規模・能力を拡大しつつある、という 3 点。 ## 提案手法 ### アーキテクチャ - **対称メモリ(symmetric memory)**: SHMEM 系 PGAS モデルに由来し、各 processing element(GPU)上で同一の型・サイズ・レイアウトを持つ symmetric object を symmetric heap 上に配置する。CUDA VMM により統一仮想アドレス空間へマッピングされ Load/Store Accessible(LSA)になる。NVSwitch + NVLink SHARP(NVLS)環境では multimem load/store 命令によりハードウェアマルチキャストとイン-ネットワーク縮約が可能になる(Fig. 2)。 - **push 型 one-shot / two-shot AllReduce**: push 型 one-shot は各 GPU がデータを全ピアへ書き込み、ローカルで縮約する(remote store は remote load の半分の RTT で済むためレイテンシに有利)。two-shot は ReduceScatter フェーズ+AllGather フェーズに分割し、通信量を O(N・M) から O(M) に削減する(N ランク、M バイトのメッセージ)。 ### アルゴリズム/手法の詳細(barrier-free 化の 4 技術) 既存実装は `ncclLsaBarrierSession` のような明示的 memory barrier に依存しており、GB200 で測定した barrier latency は GPU 数 2〜32 にわたり 0.85〜1.68 µs(unicast)/1.12〜1.25 µs(NVLS multicast)に達する(Fig. 3)。小メッセージ AllReduce が約 5 µs で完了する場合、2 回の barrier だけで総レイテンシの約 40% を占めるため、これを排除する 4 技術を提示する。 1. **LL(low latency)**: NCCL の LL プロトコルに由来。8 バイトのフラグを 8 バイトのデータとパックし 16 バイトの atomic store で送信することで、明示的な同期フラグと順序保証のステップを排除する。有効ペイロード帯域を半減させ scratch バッファ使用量を倍増させるため、非常に小さいメッセージに適する。 2. **sentinel 同期**: 受信側の scratch バッファをあらかじめ「-NaN」のような出現しにくい値(sentinel)で初期化し、値が sentinel から変化したことでデータ到着を検知する。LL に比べ有効帯域を保持し scratch 使用量も少ないが、再利用前にバッファをリセットする必要があり、送信値が sentinel と一致してはならないという制約を持つ。 3. **双方向通信 + double buffering**: LL・sentinel は単発の交換では memory barrier を排除できるが、バッファ容量の制約で複数イテレーションが必要な場合は不十分になる。2 つの scratch buffer(Buffer 0/1)を交互に使い、各ランクが同一イテレーション内で同じピアと送受信を行うことを利用して、ピアからの受信そのものが次送信への暗黙の許可(credit-based flow control 類似)として機能し、global memory barrier なしで複数イテレーションを実行できる(Fig. 4)。 4. **two-shot LL128 atomic AllReduce**: LL・sentinel と異なる同期機構を使う新しいアルゴリズム。ReduceScatter フェーズでは 8 スレッドの groups(496 レギュラー+16 エキストラスレッド構成、128 バイトのキャッシュラインに対応)が cache-line 単位の atomic addition でスクラッチバッファへ直接縮約結果を蓄積し、AllGather フェーズでは flag carrier がキャッシュラインの先頭要素が N(ランク数)に達したことをポーリングして全ランクの寄与完了を検知する(Fig. 5)。FP32 で 128 バイトあたり 4 バイト(約 3%)、FP16/BF16 で 2 バイト(約 1.5%)の追加バッファのみで済み scratch 空間効率が高い。ただし NVLink のキャッシュライン単位 atomic addition 保証を要求し、単精度・半精度浮動小数点型の加算にのみ対応し(commutativity に依存するため乗算は不可)、浮動小数点 atomic の順序非決定性により non-deterministic である。数値安定性は標準の forward-error bound に従う(FP32・64 ランクで最悪係数 γ₆₃ ≈ 3.8×10⁻⁶)。 ### 低レイテンシ API(Fig. 6, 7) - **ncclLLBuffer**: 対称メモリをラップする device-side オブジェクトで、`ncclLLSyncMode` テンプレートパラメータにより LL/sentinel モードを統一インタフェースで切り替え可能にする。`bytesPerCtaPerEpoch`・`roundRobinFactor` で CTA ごとのバッファ領域とサブバッファ数(double buffering 等)を制御し、`advanceEpoch()` でアクティブなサブバッファを切り替える。LL128 は 8 スレッド単位で動作するためこの thread-level API 設計には含まれない。 - **send / recv / recvUnrolled / recvReduce / bcast / reset**: それぞれピアスロットへの書き込み、単一スロットのポーリング受信、複数要素のコンパイル時アンロール受信、受信+縮約の統合、全ピアへの一斉書き込み(マルチキャスト対応)、スロットのリセットを担う thread-level プリミティブ。 - **host-side ユーティリティ**: `ncclCalcScratchBufferSize()` が構成に応じた最小バッファサイズを計算し、`ncclLLBufferInitSentinel()` が sentinel 値でバッファを初期化する。 - **設計上の制約**: NVSHMEM のような warp/block レベル API は提供せず thread level に限定する(最大限のユーザ制御とレイテンシ最小化を優先)。独立イテレーション通信(one-shot/two-shot)は双方向パターンにより barrier 不要だが、ring AllReduce のように前段階への依存を持つアルゴリズムは引き続き明示的 barrier を要する。 - NCCL への統合は `NCCL_SYM_KERNEL`(`AllReduce_LLBuffer` / `AllReduce_LLBuffer_Twoshot` / `AllReduce_LL128_Atomic`)と `NCCL_SYM_LLBUFFER_SYNC` の環境変数で選択できる。 ## 新規性 - 既存の一般的な one-shot/two-shot AllReduce 実装(NCCL 含む多数のライブラリ)は memory barrier に依存し、SoL 下限に対して大きなオーバーヘッドを残していた。本研究はこの barrier コストを実測(GB200 で 1 回あたり 1 µs 超)で定量化し、4 つの barrier-free 技術の組み合わせでこれを完全に排除する初の体系的な設計を提示する。 - SoL 下限の理論的な導出方法(`LSoL = 2·L_L2RTT + L_remote store`)を、`__threadfence()` のレイテンシ計測(L2 RTT の近似)と 2 GPU 間 ping-pong 計測(remote store latency の導出)から実測に基づいて確立した点も新しい(GB200 で L_L2RTT=0.306 µs、L_remote store=0.792 µs、LSoL=1.404 µs)。 - two-shot LL128 atomic は NCCL の LL128 プロトコルに着想を得つつ、atomic addition による同期という新しい機構を導入し、GPU 数増加時のスケーラビリティ(atomic 演算の L2 キャッシュ内競合コストの相対的縮小)と scratch 空間効率(D/N バイトのみ)を両立させる。 ## 実験設定 - **実験環境**: NVIDIA GB200 NVL72 システム(4 Blackwell GPU/ノード、72 GPU が単一 NVLink ドメインで接続、130 TB/s 集約帯域幅)。再現性のため NVIDIA vLLM コンテナ(v26.02、Ubuntu 24.04、CUDA 13.1、vLLM 0.15.1、PyTorch 2.11、OpenMPI 4.1.9)を使用。各実験は 10 試行の平均を報告(誤差棒は標準偏差、多くは可視化できないほど小さい)。cuSOLVERMp 実験は Alps スーパーコンピュータ(4× NVIDIA Grace Hopper Superchip/ノード、150 GB/s NVLink、NVSwitch 非搭載のため単一ノードに限定、NVIDIA PyTorch コンテナ v25.10、CUDA 12.6、OpenMPI 4.1.7、cuSOLVERMp 0.7.2)で実施。 - **比較対象**: NCCL legacy(ring・tree)・NCCL symmetric memory kernel(v2.29.1、AGxLL/RSxLD-AGxST)・NVSHMEM(v3.5.21)・MSCCL++(v0.8.0)・vLLM custom AllReduce(v0.15.1)・NCCLX CTran(ctdirect、単一ノードのみ)。MSCCL++ のマルチキャスト実装は GB200 で一貫してハングするため除外。 - **評価指標**: マイクロベンチマークではメッセージサイズ別のレイテンシ(µs)と SoL 下限からのオーバーヘッド(%)。LLM 推論では inter-token latency(ITL、ms)・出力スループット(tokens/s)・出力 100 万トークンあたりの推定コスト($、CoreWeave の GB200 オンデマンド価格 $42/時間から算出)。cuSOLVERMp では GFLOPS/GPU。 - **LLM 推論ワークロード設定**: 長入力文脈 100〜200k トークン、出力 16K トークン、バッチサイズ 8(長文脈推論を模擬)。1 ノード 4 GPU(TP=4)と 2 ノード 8 GPU(TP=8)の 2 構成。対象モデルは Llama-3.1-70B-Instruct・DeepSeek-R1-Distill-Llama-70B・Qwen3-Next-80B-A3B-Instruct(4 GPU 構成)、DeepSeek-V3・Qwen3-235B-A22B-Thinking-2507(8 GPU 構成)。各設定 5 試行、vLLM の serving benchmark を使用。 ## 実験結果 ### マイクロベンチマーク(Fig. 11) 提案の LLBuffer ベース one-shot カーネルは全 GPU 数にわたり小メッセージで最低レイテンシを達成し、2 GPU で SoL 下限比約 7% のオーバーヘッドまで到達する(競合実装はより大きく乖離)。64 GPU でもマルチキャスト one-shot 亜種は SoL 比約 70% のオーバーヘッドに収まる。LL は極小メッセージでわずかに優位、sentinel はメッセージサイズ・ランク数が増えるほど有利になるクロスオーバーが観測される。LL128 atomic は GPU 数が少ないと two-shot 標準設計に対する優位性が限定的だが、GPU 数が増えるほど atomic 演算のスケーラビリティ上の利点が顕在化する。ハードウェアマルチキャストは小規模では若干のオーバーヘッドを持つが、`multimem.ld_reduce` 命令により大規模でのスケーラビリティ改善に重要な役割を果たす。scratch バッファサイズの影響(Fig. 12)では、one-shot カーネルは 4 MiB で頭打ちになる一方、two-shot・LL128 atomic カーネルは 64 MiB まで恩恵を受けることが確認され、それぞれのデフォルトスクラッチサイズとして採用された。 ### LLM 推論(Fig. 13) 全モデルで低レイテンシカーネルが一貫して性能を改善する。最良構成は ITL を 4 GPU で 7〜13%・8 GPU で 9〜11.4% 削減する。NCCL LL 単体でもデコードステップの小さな AllReduce に対して SoL に近い性能を発揮するため有効。Symmetric memory 対応(PyTorch のバッファをシンメトリックメモリとして登録)によって two-shot・LL128 atomic カーネルが有効になり、vLLM の prefill/decode 混在実行でさらなる改善(特にスループット面)が得られる。DeepSeek-V3 のような大規模モデルの 8 GPU 構成では、出力 100 万トークンあたり $11 超のコスト削減が推定された。 ### 伝統的 HPC(cuSOLVERMp、Fig. 14) Alps スーパーコンピュータ上の 2/4 GH200 構成で GFLOPS/GPU が一貫して改善する。通信がランタイムに占める割合が大きい m=32768(GFLOPS/GPU +7.0%)で改善がより顕著であり、m=65536(+1.5%)では相対的に小さい。cuSOLVERMp はバッファを symmetric memory として登録しないため 1 MiB 未満のメッセージサイズでは one-shot カーネルのみが使用された。 ## 考察 - 低レイテンシ集合通信の最適化は AI 推論(LLM のデコードフェーズ)と伝統的 HPC(MILC・LULESH のようなタイトに同期したソルバー、cuSOLVERMp)の双方に恩恵をもたらすことが実証され、単一の技術セットが異なるドメインに横断的に効くことを示した。 - 提案の LLBuffer ベースカーネルは既存の symmetric two-shot(RSxLD-AGxST)を完全には代替できない: ポーリングオーバーヘッドが GPU 数・メッセージサイズとともに増大するため、小〜中サイズメッセージの補完として位置づけられる。 - 関連研究(SGLang・FlashInfer・TensorRT-LLM)は類似設計だが、TensorRT-LLM は sentinel 風の同期を使いつつも依然として global barrier フラグに依存しており本研究の barrier-free 設計とは異なる。DeepEP・NCCL EP のような expert-parallel 特化ライブラリはより狭い通信パターンを対象とし汎用性がない点で対比される。NIXL は集合通信のレイテンシ最適化ではなく異種バックエンドを跨ぐ転送・オーケストレーション層であるため比較対象から除外されている。 ## 強み / 弱点・課題 ### Strengths - barrier latency の実測に基づく問題の定量化から、4 つの barrier-free 技術・SoL 理論下限の導出・低レイテンシ API 設計・実アプリケーション(vLLM・cuSOLVERMp)での実証までを一貫して構築している。 - SoL 下限を GB200 実測(1.404 µs)として明確に定義し、既存実装との比較基準を確立した点は今後の低レイテンシ集合通信研究の参照点になりうる。 - コードは公開されており(https://github.com/ss16118/low-latency-nccl)、NCCL への統合により実践的に利用可能。 ### Weaknesses / Limitations - カーネル選択(メッセージサイズ・GPU 数に応じたアルゴリズム切り替え)は現状、経験的な実測に基づいており、正確な性能モデルは今後の課題として残されている(GPU アーキテクチャ・warp スケジューリング・命令レベル挙動の詳細な理解を要する)。 - API は thread-level プリミティブに限定されており、warp/block レベルの抽象化は将来拡張として位置づけられる。 - LL128 atomic は単精度・半精度浮動小数点の加算にのみ対応し、non-deterministic である(浮動小数点 atomic の順序非保証)ため、決定性を要する用途には不向き。 - scale-out(マルチノード・複数 NVLink ドメイン)は本研究のスコープ外であり、scale-up 内の最適化がマルチノード全体の性能にどの程度寄与するかは階層的集合通信の上位層設計に依存する。 - cuSOLVERMp の評価は単一ノード(NVSwitch 非搭載の Alps)に限定されており、GB200 のような大規模 scale-up ドメインでの HPC ワークロード評価は今後の課題。 ## 図表 **Figure 1: 長文脈・小バッチ TP LLM 推論における集合通信のレイテンシボトルネック** ![[_attachments/arxiv-2607.16100/fig01-motivation-tp-decode.png]] (Figure 1. 長文脈・小バッチのテンソル並列(TP)LLM 推論では、多くの小さな AllReduce 操作がトークン生成のクリティカルパス上に位置する。マイクロベンチマークは、提案の NCCL 低レイテンシカーネルが他実装に対して小メッセージ AllReduce レイテンシを削減し、SoL 下限に近づくことを示す。これは Llama-3.1-70B 推論において、より低い inter-token latency(ITL)とコスト削減につながる。Source: Figure 1.) **Figure 2: NCCL におけるデバイス起動通信と対称メモリの概要** ![[_attachments/arxiv-2607.16100/fig02-symmetric-memory-overview.png]] (Figure 2. GPU が同一ノードまたは同一 NVLink ドメイン内にある場合、LSA 操作は PCIe・NVLink 経由でサポートされ、multimem 操作は NVLink SHARP によりハードウェアアクセラレートされたマルチキャストと縮約に対応する。ノード間通信では GPU-initiated networking(GIN)が GDAKI とプロキシ支援のデータ転送を InfiniBand・RoCE 上でサポートする。Source: Figure 2.) **Figure 3: GB200 における GPU 数別のバリアレイテンシ** ![[_attachments/arxiv-2607.16100/fig03-barrier-latency.png]] (Figure 3. unicast・multicast 実装それぞれについて、GPU 数(2〜32)の関数としてのバリアレイテンシ(µs)を示す。unicast は 0.85(2 GPU)から 1.68 µs(32 GPU)、NVLS multicast は 1.12〜1.25 µs の範囲。Source: Figure 3.) **Figure 5: two-shot LL128 atomic AllReduce アルゴリズムの概要** ![[_attachments/arxiv-2607.16100/fig05-ll128-atomic-allreduce.png]] (Figure 5. CTA 内のスレッドは 496 のレギュラースレッドと 16 のエキストラスレッド(変位要素処理用)に分かれる。緑の破線ボックスは 128 バイトのキャッシュラインを操作する 8 スレッドのグループを示す。パネル A は ReduceScatter フェーズ、パネル B は AllGather フェーズを示す。出力バッファはカーネル実行前に sentinel 値で初期化されるアウトオブプレース操作を想定。半精度浮動小数点をサポートする場合、各要素は 2 バイトで 8 個のエキストラスレッドのみが必要。Source: Figure 5.) **Figure 11: メッセージサイズ・GPU 数別の AllReduce レイテンシ比較(マイクロベンチマーク)** ![[_attachments/arxiv-2607.16100/fig11-microbenchmark-latency.png]] (Figure 11. 上段の各サブプロットは GB200 上で 2〜64 GPU の out-of-place AllReduce をメッセージサイズに対してプロットしたもの(NCCL・NCCLX・NVSHMEM・MSCCL++・vLLM の各実装との比較、10 試行平均)。網掛け領域は提案カーネル(one-shot・two-shot・LL128 atomic)が最速となるメッセージサイズ範囲を示し、ラベルは当該範囲での既存最速実装に対する幾何平均speedupを報告する。下段は 128B メッセージにおける GPU 数別レイテンシで、破線は実測 SoL 下限(1.404 µs)、パーセンテージは各カーネルの SoL 下限に対するオーバーヘッドを示す。Source: Figure 11.) **Figure 13: 低レイテンシ集合通信の vLLM 推論への効果** ![[_attachments/arxiv-2607.16100/fig13-vllm-inference-results.png]] (Figure 13. 各行は平均 inter-token latency(ITL)・出力スループット・ベースライン比の推定コスト削減(100 万出力トークンあたり)を示す。パーセンテージはベースラインに対する改善率、緑ラベルは各モデルで最良の構成を示す。1 ノード 4 GB200(TP4)・2 ノード 8 GPU(TP8)の両構成を含む。Source: Figure 13.) **Figure 14: 低レイテンシカーネルの有無による mp_sygvd(cuSOLVERMp)の性能** ![[_attachments/arxiv-2607.16100/fig14-cusolvermp-results.png]] (Figure 14. 2/4 GH200 構成における GFLOPS/GPU の平均値(5 試行、誤差棒は標準偏差)。パーセンテージ注記はベースラインに対する改善率。通信がランタイムに占める割合が大きい m=32768 でより顕著な改善(+7.0%)が見られる。Source: Figure 14.)