# ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernels > [!info] Talk metadata > - **会議:** [[MLSys2026]] Day 4 (May 21 / Thu)、Grand Ballroom 2、16:30 - 16:45 PDT > - **著者:** Stuart H. Sul, Simran Arora, Benjamin F. Spector, Christopher Re（Stanford University / HazyResearch） > - **URL:** https://mlsys.org/virtual/2026/oral/3845 > - **OpenReview:** https://openreview.net/forum?id=Cv5e5uRXFb > - **関連研究:** https://github.com/HazyResearch/ThunderKittens > [!abstract] 概要 > GPU 間通信は現代の AI ワークロードにおける主要なボトルネックとなっている。モデルの規模拡大に伴い、ハードウェアの演算スループット向上がインターコネクト帯域の改善を上回っているためである。既存手法は演算と通信のオーバーラップによってこの問題を緩和しようとするが、異種ワークロードや新世代アクセラレータにおいて理論的ピーク性能を達成できないことが多い。本研究はオペレータ固有の手法に代えて、少数の単純かつ再利用可能な原則がマルチ GPU カーネル設計を体系的に導けるかを問う。提案する ParallelKittens (PK) は、オーバーラップされたマルチ GPU カーネルの開発を大幅に簡素化する最小限の CUDA フレームワークである。PK は ThunderKittens フレームワークを拡張し、8 つの中核プリミティブと統一プログラミングテンプレートを通じてマルチ GPU カーネル設計原則を体現する。Hopper および Blackwell アーキテクチャで検証した結果、50 行未満のデバイスコードで、データ並列・テンソル並列ワークロードに対し最大 2.33 倍、系列並列ワークロードに対し 4.08 倍、エキスパート並列ワークロードに対し 1.22 倍の高速化を達成した。 ## テーゼ・背景 - 数年前まで GPU 演算のボトルネックは GPU 内メモリアクセスであったが、FlashAttention やタイルベース DSL などの IO 最適化が進んだ結果、GPU 間通信が残存する最大のボトルネックとなった。 - NVLink のような高速インターコネクトを用いても、LLM ワークロードにおける通信は実行時間の 50% 以上を占めうる。Nvidia A100 から B200 にかけて、BF16 テンソルコア性能は 7.2 倍、HBM 帯域は 5.1 倍向上したが、NVLink 帯域は約 3 倍、PCIe/InfiniBand は約 2 倍の改善にとどまる。 - 既存の通信オーバーラップ手法には 3 つの選択肢があるが、いずれも課題を抱える。(1) 既製ライブラリ（NCCL 等）は手動チューニングの 4.08 倍まで遅いことがある。(2) コンパイラベースのアプローチ（Triton Distributed 等）は新アーキテクチャへの適応に失敗し、非オーバーラップベースラインを下回ることがある。(3) 低レベルプリミティブ（CUTLASS、NVSHMEM 等）は複雑で実装負荷が高い。 - 本研究の問いは、ハードウェアのトレードオフとマルチ GPU カーネル設計空間を明確にし、最小限のプログラミング抽象・プリミティブの集合に到達できるかである。 - ハードウェア側の新機軸として、インネットワーク演算（NVSwitch 上でのリダクション）、非同期かつバルクなデバイス主導ネットワーキング（TMA）、スケールアップアーキテクチャ（NVL72/NVL144/NVL576）がある。 - PK は Cursor 社の大規模社内学習（Composer 2）および Together.ai で本番採用されている。 ## コストモデル 論文ではマルチ GPU カーネル設計の分析基盤として、以下のコストモデルを定式化している。 $T_{\text{kernel}} = T_{\text{launch}} + \max(T_{\text{comp}},\; T_{\text{mem}},\; T_{\text{comm}}) + T_{\text{non-overlap}} + T_{\text{sync}}$ - $T_{\text{launch}}$: カーネル起動コスト（ホスト側レイテンシ、スレッドブロックのセットアップ、テンソルメモリの確保、パイプラインのフィル/ドレインフェーズを含む） - $T_{\text{comp}},\; T_{\text{mem}},\; T_{\text{comm}}$: 演算・メモリアクセス・通信のフルパイプライン時間。理想的にはこれら 3 つの最大値のみが全体時間に支配する - $T_{\text{non-overlap}}$: オーバーラップできない操作の時間 - $T_{\text{sync}}$: SM 間またはデバイス間の同期オーバーヘッド 各コンポーネントのコスト（例えば $T_{\text{comm}}$）は、ワークサイズ $S_{\text{comm}}$ と達成帯域 $B_{\text{comm}}$ に依存し、$T_{\text{comm}} = S_{\text{comm}} / B_{\text{comm}}$ で表される。このモデルにおいて 3 つの設計決定（転送メカニズム・スケジューリング・設計オーバーヘッド）がコストを制御する。 ## 3 つの設計トレードオフ ### トレードオフ 1: 転送メカニズム - GPU 間データ転送には 3 つのメカニズムが存在する。コピーエンジン（CE）、テンソルメモリアクセラレータ（TMA）、レジスタレベル命令である。 - コピーエンジンはホスト起動で、大規模かつ連続的な転送に最適である。H100 で 368.82 GB/s（理論最大の 82%）、B200 で 726.13 GB/s（81%）を達成するが、帯域の 80% 以上を活用するには 256 MB 以上のメッセージサイズが必要となる。 - TMA はデバイス起動で、単一スレッドから非同期に発行可能であり、わずか 2 KB のメッセージで近ピーク帯域を達成する。H100 で 350.01 GB/s（78%）、B200 で 669.12 GB/s（74%）を示す。NVLink 帯域の飽和には約 15 SM で足りる。 - レジスタ命令はインネットワークリダクションやファブリック内ブロードキャストなど最も多機能であるが、帯域飽和に 76 SM 以上を要し、同期的でレジスタプレッシャーが高い。H100 で 342.68 GB/s（76%）、B200 で 628.35 GB/s（70%）。 - 既存手法（Triton Distributed、Flux、CUTLASS）はいずれもコピーエンジンに依存しており、ファイングレイン通信で性能が劣化する。PK は TMA を採用し、細粒度転送で 0.97 倍から 7.39 倍の改善を達成する。 - 機能面では、P2P 転送とファブリック内ブロードキャストは 3 メカニズムすべてが対応する。P2P リダクションは TMA とレジスタのみ、ファブリック内リダクションと要素単位転送はレジスタのみが対応する。 - TMA の最大メッセージサイズは 227 KB であり、それを超える値のスループットは視覚比較用に一定値として描画されている（Figure 2）。 - レジスタレベル命令は NVLink 帯域の飽和に TMA の 3.2 倍から 5.1 倍の SM 数を要する。レジスタ命令が有用なのは、コピーエンジンにも TMA にもない機能（NVSwitch インネットワークリダクション: `multimem.ld_reduce`、`multimem.red`）が必要な場合に限られる。 ### トレードオフ 2: スケジューリング - 演算と通信のオーバーラップには 2 つの戦略がある。SM 内オーバーラップ（intra-SM）は同一 SM 内の異なるスレッドプールが演算・メモリ・通信を分担する。SM 間オーバーラップ（inter-SM）は SM 群をまるごと演算用と通信用に分割する。 - SM 内オーバーラップは全 SM の演算ユニット（テンソルコア）を活用でき、SM 間同期を必要としないため、演算と通信の粒度が一致する場合に最適である。GEMM + ReduceScatter では SM 間オーバーラップより 1.2 倍高い性能を示す（743.7 対 618.1 TFLOP/s、No Overlap は 510.1 TFLOP/s）。 - SM 間オーバーラップは演算と通信の粒度が不一致でも対応でき、インネットワーク演算やリモートキャッシュ再利用を活用できる。GEMM + AllReduce では SM 内方式の 172.3 TFLOP/s に対し 623.9 TFLOP/s と 3.62 倍の改善を達成する。 - SM 内オーバーラップの限界として、ピア HBM へのリモートアクセスにおける L2 キャッシュの片側性（far-sided）がある。データは送信元でのみキャッシュされ、要求元ではキャッシュされないため、Ring Attention のような KV テンソルの再利用パターンでインターコネクトが飽和しうる。 - SM 間オーバーラップでは通信専用 SM と演算用 SM の分割比を入力サイズに応じて最適化する必要がある。PK は統一プログラムテンプレートを通じて実行時に最適な SM 分割を自動探索可能とする。 - 同期コストにも差がある。SM 内同期（mbarrier）は約 64 ns であるのに対し、SM 間同期は HBM を経由するため約 832 ns かかる。 #### 通信完全隠蔽の条件（SM 内オーバーラップ） $M \times N \times K$ の GEMM+RS 融合カーネルにおいて、反復タイルサイズ $m \times n \times k$ とすると、出力タイル $m \times n$ を生成するまでに $K/k$ 回のサブ GEMM が実行される。要素あたりバイト数 $s$、テンソルコアスループット $R$（FLOP/s）、GPU あたり NVLink 帯域 $B$（バイト/s）とすると、 $T_{\text{comp\_tile}} = \frac{2mnK}{R}, \quad T_{\text{comm\_tile}} = \frac{smn}{B}$ 通信が演算に完全に隠蔽される条件は $T_{\text{comp\_tile}} \geq T_{\text{comm\_tile}}$、すなわち $K \geq \frac{sR}{2B}$ H100 GPU で BF16（$s = 2$）の場合、$R = 989 \times 10^{12}$、$B = 450 \times 10^9$ であり、$K \gtrsim 2197$ で通信が完全に隠蔽される。この理論予測は Table 3 の実測値と整合する。 | $M \times N$ | $K$ | GEMM (ms) | GEMM+RS (ms) | 通信比率 | |---|---|---|---|---| | 32768 | 512 | 2.071 | 6.483 | **68%** | | 32768 | 1024 | 2.918 | 6.613 | **56%** | | 32768 | 2048 | 5.567 | 7.531 | **26%** | | 32768 | 4096 | 11.78 | 11.828 | **<1%** | | 32768 | 8192 | 23.285 | 25.325 | **8%** | $K = 2048$ で通信比率がほぼ半減し、$K = 4096$ 以降でほぼ完全に隠蔽される。$K = 2048$ 付近の残存通信は出力タイル蓄積時のアトミック加算に起因する。 #### SM 間オーバーラップにおけるインネットワーク加速 SM 内オーバーラップでは all-reduce の際に各出力タイルに対し $N$（GPU 数）回のアトミック書き込みが発生し、完全接続 NVSwitch ファブリック（ポートあたり 450 GB/s）でも並行ピア書き込みの直列化が帯域ボトルネックとなる。SM 間オーバーラップでは通信 SM が部分結果をローカル HBM に蓄積した後、インネットワーク all-reduce を一括実行することで $T_{\text{comm}}$ をおよそ $N$ 分の 1 に削減する。GEMM+AR では SM 内方式の 172.3 TFLOP/s に対し SM 間方式が 623.9 TFLOP/s を達成し（同 450.9 TFLOP/s の No Overlap を含む Figure 4）、3.62 倍の改善に至る。 #### SM 分割比の入力依存性 SM 間オーバーラップにおける通信用 SM と演算用 SM の最適分割比は入力サイズに依存する（Figure 5）。AG+GEMM ($N \times N/8 \times N$) の場合、$N = 8192$ では通信 SM 約 10 台が最適だが、$N = 65536$ では約 40 台が最適となる。小さいワークロードほど演算 SM の割合を高く、大きいワークロードほど通信 SM を増やす必要がある。 ### トレードオフ 3: 設計オーバーヘッド - NCCL などの通信ライブラリは、簡潔性を志向する設計上の選択（同期・バッファリング）がかえって性能を劣化させる。 - 双方向同期と中間バッファリング: NCCL はすべての操作で双方向同期を強制し、事前確保された中間バッファ（通信チャネル）を用いるため、ファイングレイン通信で追加のデータ移動が発生する。PK は事前確保された宛先バッファへの直接的な片方向転送を用い、中間ステージングを排除する。 - ピアメモリアクセスと同期: NVSHMEM はリモートピアアクセスのたびにグローバルメモリロード（`__ldg`）でアドレスを取得し、グループ同期（`__syncthreads`）を強制する。PK はピアアドレスをレジスタに保持し、不要な同期を除去することで、NVLink アクセスレイテンシを最大 4.5 倍削減し、帯域利用率を約 20 GB/s 向上させる。 - All-reduce カーネル単体で、PK は NCCL に対し H100 で最大 1.32 倍、B200 で最大 1.79 倍の性能改善を達成する（Figure 6）。H100 における行列サイズ別の改善: 2048 で 1.28 倍、4096 で 1.32 倍、8192 で 1.0 倍、16384 で 1.12 倍、32768 で 1.07 倍、65536 で 1.02 倍。B200 では小規模行列での優位が顕著であり、2048 で 1.79 倍、4096 で 1.66 倍と、ファイングレイン通信での PK の設計上の優位性を示す。 - NCCL の性能損失は純粋な通信カーネルでも 1.7 倍以上に達しうる。通信レイテンシでは最大 4.5 倍の増加が確認される。 ## ParallelKittens フレームワーク ### データ構造 - PK は GPU メモリ階層の各レベルに対応するデータ構造を定義する。レジスタレベルでは 16x16 タイルの Register Tile (`rt<M,N>`)、共有メモリレベルでは Shared Tile (`st<M,N>`)、HBM レベルでは Global Layout (`gl`)、ピア GPU メモリレベルでは Parallel Global Layout (`pgl`) を導入する。 - メモリ階層のパラメータ: レジスタ 64 KB・130 TB/s、共有メモリ 227 KB・33 TB/s、L2 キャッシュ 50 MB・12 TB/s、HBM 80 GB・3 TB/s、ピアグローバルメモリ 80xN GB・0.45 TB/s。 - PGL はすべてのデバイスにまたがる同一形状・同一サイズのメモリ領域を表現し、非同期 P2P 転送、ブロードキャスト、タイルインデクスに基づくインネットワークリダクションを可能とする。 - すべてのデータ抽象はバンクコンフリクト最小化、スウィズリング、テンソルコアとのレイアウト整合性、完全なデバイス主導通信などの原則を強制する。 ### 8 つのコアプリミティブ - **P2P 通信プリミティブ**: `store_async(dst, src, coord)` は非同期 P2P 転送、`store_add_async(dst, src, coord)` は非同期 P2P アトミック加算を行う。いずれも単一スレッドから発行可能で、テンソルコア演算との融合に適する。 - **ネットワーク加速通信プリミティブ**: `reduce(dst, dst_coord, src, src_coord)` はポイント対ポイントリダクション、`all_reduce(dst_and_src, coord)` はインネットワーク all-reduce を実行する。これらはワープレベル以上の参加を要する。 - **デバイス間・SM 間同期プリミティブ**: `signal(bar, coord, dev_idx, val)` は特定デバイスへのシグナル送信、`signal_all(bar, coord, val)` は全デバイスへのブロードキャストシグナル、`wait(bar, coord, dev_idx, expected)` はシグナル待機、`barrier(bar, coord, dev_idx)` はバリア同期を行う。 - すべてのプリミティブはタイル粒度（最小 16x16 から共有メモリ上限の約 256x256 まで）で動作し、デバイス主導で `int4` 座標によりタイルインデクスを指定する。 ### プログラムテンプレート - PK は 4 つのワーカーコンポーネントからなる統一プログラムテンプレートを提供する。**Loader**（ローカルまたはピア HBM からの読み出し）、**Storer**（ローカルまたはピア HBM への書き込みとデバイス間シグナリング）、**Consumer**（テンソルコアまたは CUDA コアによるローカル演算）、**Communicator**（専用 SM 上での通信・インネットワーク all-reduce）。 - Loader と Consumer が同一 SM 上でピア HBM にアクセスする場合は SM 内オーバーラップ、Communicator が専用 SM を占有する場合は SM 間オーバーラップとなる。 - テンプレートはカーネル構成、共有メモリと TMA のセットアップ、バリア・同期管理、SM/ワープ分割の最適チューニングを自動化する。開発者はタイルごとの演算・通信ロジックのみに集中すればよい。 ### ユーティリティ - PK はプロセス間通信（IPC）と PyTorch ユーティリティを提供し、`torchrun` 経由でのマルチプロセス実行をシームレスに行う。OS ドライバとの低レベルインタラクション管理やマルチ GPU メモリの事前確保をサポートし、中間ステージングなしの直接 P2P 通信を実現する。 ## 実験・結果 ### 実験環境 - 8x Nvidia H100 80GB SXM GPU、第 4 世代 NVLink/NVSwitch（片方向 450 GB/s）、5 世代 PCIe、CUDA 12.8、PyTorch 2.8.0 で実施。行列演算は BF16 要素型、FP32 アキュムレータを使用。 - GEMM 形状は $M \times N \times K$ と記述し、第 1 オペランドが $M \times K$、第 2 オペランドが $K \times N$ である。観測値は平均演算スループット（TFLOP/s）で報告する。 - Blackwell（B200）での結果は付録 A・B に記載されており、H100 と同様の性能特性を示す。 ### データ並列・テンソル並列 - **AG+GEMM** (Figure 7): ローカル GEMM サイズ $N \times N/8 \times N$。非オーバーラップベースライン（cuBLAS + NCCL）に対し 1.06 倍から 1.68 倍の高速化。コンパイラベースアプローチ（Triton Distributed）に対し 1.07 倍から 5.63 倍。手動最適化カーネル（Flux）に対し 0.97 倍から 2.33 倍、CUTLASS に対し 0.90 倍から 7.39 倍の改善。$N = 32768$ では PK が 727 TFLOP/s に達し、次点の Flux 662 TFLOP/s、cuBLAS+NCCL 628 TFLOP/s を上回る。AG+GEMM と GEMM+RS は実際には連続して使用されるため、単一のベースラインでは両方で PK を上回ることはない。 - **GEMM+RS** (Figure 8): ローカル GEMM サイズ $N \times N \times N/8$。SM 内オーバーラップを採用。$N = 32768$ で PK が 793 TFLOP/s を達成し、全ベースラインを上回る（Flux 744、CUTLASS 431、cuBLAS+NCCL 602、Triton Distributed 510）。PK は SM 内オーバーラップ、GEMM+AR では SM 間オーバーラップを使い分ける。 - **GEMM+AR** (Figure 9): ローカル GEMM サイズ $N \times N/8$。SM 間オーバーラップでインネットワークリダクションを活用。Flux と CUTLASS は GEMM+AR カーネルを提供しないため省略。$N = 32768$ では PK が 624 TFLOP/s、cuBLAS+NCCL が 317 TFLOP/s、Triton Distributed が 451 TFLOP/s。 - 十分大きなリダクション軸では、非オーバーラップ通信時間が 1% 未満に低下する。通信部分のコードは GEMM 本体を除き 50 行未満のプリミティブで実装されている。 ### 系列並列 - **Ring Attention** (Figure 10): 評価パラメータは $B = 16$、$H = 16$、$D = 128$、系列長は 768 の倍数（元の TK アテンションカーネルの要件）で 8 デバイスに均等分割。xDiT 実装に対し 1.07 倍から 4.08 倍の高速化。xDiT は NCCL P2P 送信と FlashAttention-3 カーネルを別々の CUDA ストリームで起動する粗粒度のストリームレベルオーバーラップを行うが、PK はこれを明示的な SM 間オーバーラップによる単一カーネルに融合する。通信 SM が次ブロックの KV テンソルをローカル HBM にバルク転送し、残りの SM が演算を並行実行する。この構成はリモートキャッシュ再利用の改善にも寄与する。非オーバーラップ通信比率は 9% まで削減される。 - **DeepSpeed-Ulysses** (Figure 11): 評価パラメータは $B = 16$、$H = 128$、$D = 128$、総系列長（X 軸）を 8 デバイスに均等分割。YunChang 実装に対し 1.01 倍から 1.39 倍の高速化。主なボトルネックはファイングレイン all-to-all であり、NCCL は内部次元に沿った通信をネイティブにサポートしないため、ベースラインはテンソルの reshape を通信前後に挿入する必要がある。PK はこのオーバーヘッドを排除するファイングレイン all-to-all カーネルを 50 行未満のデバイスコードで実装している。 ### エキスパート並列 - **MoE トークンディスパッチ + GEMM** (Figure 12): 評価パラメータは TopK = 8、$N_{\text{experts}} = 256$、$H = 7168$、$H_{\text{expert}} = 2048$。総入力トークン数（X 軸）を 8 デバイスに均等分割。Comet（最先端のファイングレインオーバーラップ手法）に対し 0.92 倍から 1.22 倍の性能。MoE レイヤの前半（トークンディスパッチと最初のエキスパート MLP の融合）を評価対象とする。PK は SM 間オーバーラップを採用してトークン通信と演算を並行実行する。グループ化 GEMM に追加するデバイスコードは 40 行未満。非オーバーラップ通信比率は 15% まで削減される。 ### 非オーバーラップ通信比率のまとめ 各並列パターンにおいて PK が達成した非オーバーラップ通信時間の割合は以下のとおりである。 | 並列パターン | 非オーバーラップ通信比率 | |---|---| | データ並列・テンソル並列（十分大きな $K$） | **< 1%** | | 系列並列（Ring Attention） | **9%** | | エキスパート並列 | **15%** | ### コンパイラベースアプローチとの比較 - Triton Distributed は H800 GPU 向けに開発されており、H100 では非オーバーラップベースラインを下回ることがある。PK はハードウェアプラットフォーム間で一貫した性能を示す。手動最適化カーネルも特定の問題形状では効率が低下するケースが観測されている。 ## 関連研究 論文では関連研究を 3 カテゴリに整理している。 ### オペレータ固有カーネル TP-Async、Flux、Ring Attention、DeepEP（DeepSeek-AI）、Comet、FlashDMoE、および CUTLASS ベースの分散 GEMM カーネル群が挙げられる。これらは特定の AI オペレータに対して演算と通信をオーバーラップさせる手法であり、ホスト起動コピーとデバイスカーネルの重ね合わせから高度にチューニングされたデバイス主導通信まで多様な技法を用いる。高性能を達成するが、複雑な実装を要し再利用可能な抽象を提供しない。例えば FlashDMoE は TF32 精度のみに最適化されており、BF16/FP16 サポートはリリース後 5 ヶ月経過時点で開発中であった。 ### スケジューリングフレームワーク Megatron-LM、FlexFlow、NanoFlow がバルクコレクティブ（主に NCCL）のオーケストレーションによって並列化とスケジューリングを自動化する。これらは PK と相補的であるが、ストリームレベルのオーバーラップにとどまり、SM 内ワープ特殊化やデバイス主導タイルレベル転送を行わないため、ピークカーネル性能の達成には不十分である。 ### マルチ GPU プログラミングプリミティブ Triton Distributed、TileLink が Triton を拡張して OpenSHMEM スタイルの片方向操作によるコンパイラベースのマルチ GPU カーネル生成を可能にする。しかし、ワークロード分配の明示的制御（ワープ/SM 特殊化）を欠き、最適なオーバーラップに必要な柔軟性を提供しない。ベンチマークでは、H800 GPU 向けにチューニングされた Triton Distributed が H100 では非オーバーラップベースラインを下回ることがある。NCCL の大規模クラスタ版である NCCLX は 10 万 GPU 以上のノード間コレクティブを加速するが、デバイス主導の非同期オーバーラップやインネットワーク加速を利用しない点で PK と相補的である。AMD の Iris（ROCm）や Google の Pallas（JAX）もマルチ GPU プログラミング体験の改善を目指しているが、PK とは異なるハードウェアプラットフォームを対象とする。 ## 結論・オープン課題 - PK はマルチ GPU カーネル設計空間を 3 つの原則（転送メカニズム・スケジューリング・設計オーバーヘッド）に形式化し、少数のプリミティブで手動最適化カーネルと同等以上の性能を、実装の大幅な簡素化とともに達成することを実証した。 - 本研究はノード内（intra-node）実行に焦点を当てており、ノード間（inter-node）通信への拡張が重要な将来課題として残る。この制約は論文が明示する主要な限界事項である。 - ノード内スケールアップは Nvidia の NVL72（2026）から NVL144、NVL576（2027）へのロードマップに示されるように急速に拡大しており、効率的なノード内カーネル設計の研究の重要性は増している。 - 今後の方向として InfiniBand/イーサネットへの拡張、および AMD GPU や Apple M シリーズなどのマルチシリコン対応が示された。 - 本番採用: Cursor 社の大規模社内学習（Composer 2 テクニカルレポート arXiv:2603.24477）および Together.ai で採用されている。 - PK のフレームワークとカーネルは https://github.com/HazyResearch/ThunderKittens で公開されている。 - 受賞: MLSys 2026 において Artifacts Available、Artifacts Evaluated-Functional、Artifacts Evaluated-Reproducible の 3 バッジを獲得している。 ## GPU ネットワーキング階層（補足） 論文 Section 2.1 で整理されたマルチ GPU システムのインターコネクト階層は以下のとおりである。 | 層 | 帯域 | 用途 | |---|---|---| | PCIe（第 5 世代） | 64 GB/s | CPU-GPU 間、カーネル起動、ホスト起動転送、InfiniBand/TCP 経由のマルチノード通信 | | NVLink | 450 GB/s（片方向） | GPU 間ポイント対ポイント接続、ピア GPU HBM アクセス | | NVSwitch | 全 NVLink エンドポイントを非ブロッキングファブリックに接続 | フル GPU 間通信、マルチキャスト・リダクションのオフデバイス加速（インネットワーク演算） |