# NVSHMEM NVIDIA SHMEM。各GPUをPGAS(Partitioned Global Address Space)における1つのPE(Processing Element)として扱い、GPUカーネルから直接、他GPUのメモリへ一方向(one-sided)にput/getできるようにするライブラリ。CPUを介さずGPU間でデータ転送と同期を完結できる点が特徴で、動的タスクキューや細粒度イベント通知を要する不規則ワークロードに向く。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 12 Dynamic Scheduling, CUDA Graphs, and Device-Initiated Kernel Orchestration]], §Fine-Grained GPU-to-GPU Memory Sharing with NVSHMEM) ## 仕組みと主要API - **一方向put/get + フラグ同期**: 送信側は `nvshmem_float_p` でリモートPEのメモリへペイロードを書き込み、`nvshmem_quiet()` で完了を保証したのち `nvshmem_int_p` でフラグを立てる。受信側は `nvshmem_int_wait_until()` でフラグをスピン待機し、成立後にペイロードを読む。ホストの介在は不要で、NVLink上のハードウェア加速転送のみで完結する。 - **細粒度同期プリミティブ**: `nvshmem_wait_until`・`nvshmem_signal_fetch`・`nvshmem_signal_wait_until` 等の点対点同期を優先し、`nvshmem_barrier_all()` の多用は避けるべきとされる(全体バリアは最も遅いピアに全GPUを引きずられさせるため)。 - **アトミック操作によるワークスティーリング**: `nvshmem_int_atomic_inc` でグローバルカウンタをインクリメントし、各GPU(PE)がホスト協調なしに次のタスクを動的に取得できる。 - **協調カーネル起動**: `nvshmemx_collective_launch()` で複数GPUにまたがるカーネルを同時起動し、`nvshmem_barrier_all()` によるロックステップ実行を可能にする。NVSHMEMのデバイス側同期・集団操作を使うカーネルは必ずこのAPIで起動する必要がある。 - **CUDA Graphsとの関係**: NVSHMEM呼び出しを含むカーネルはCUDA Graphにキャプチャできるが、NVSHMEM操作自体は個別のグラフノードにはならない(カーネル内部の動作として扱われる)。NVSHMEMの強みは、グラフのような固定された通信スクリプトなしにカーネルが実行時に適応・調整できる点にある。 ## 位置づけ NCCLの集団通信(all-reduce等)がバルクな定型パターンに適するのに対し、NVSHMEMは不規則・データ依存のワークロード(グラフアルゴリズム、動的負荷分散、離散イベントシミュレーション)や、GPU間の超低遅延イベント駆動協調(2段トランスフォーマー推論パイプラインでのアテンション→MLPステージ間ハンドオフ等)に向く。GPU起動型ネットワーキングの文脈では、NVSHMEMのGPUDirect Async Kernel-Initiated(GDAKI)機能がPGASモデルによるデバイス起動通信を確立したライブラリとして位置づけられ、NCCL 2.28のDevice API(GIN)はこのモデルをNCCLエコシステムに統合したものとされる。(Source: [[GPU起動型ネットワーキング]]) ## 関連 - ソース: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 12 Dynamic Scheduling, CUDA Graphs, and Device-Initiated Kernel Orchestration]] - 概念: [[GPU起動型ネットワーキング]] - エンティティ: [[NCCL]] / [[DeepEP]] ## 出典 - Chris Fregly, *AI Systems Performance Engineering*, O'Reilly Media, 2025, Chapter 12.