# CUDA Stream ## 定義 CUDAストリーム(CUDA stream)は、カーネル起動・メモリコピー・メモリ確保といったGPU操作を発行順に実行する操作列である。単一ストリーム内の操作は逐次実行されるが、複数のストリームに分けて発行した操作はハードウェア資源(SM・DMAエンジンの空き)が許す限り互いに重なり合って実行される。CPUはストリームへの操作をエンキューすると即座に制御を取り戻し、ホスト側コードを継続できる(非同期発行)。これによりカーネル計算・ホスト↔デバイスのメモリ転送・メモリ確保/解放を並行させ、GPUの複数エンジン(SM、Host→Device DMAエンジン、Device→Host DMAエンジン)を同時に稼働させることが、モダンGPUで実効性能を引き出す基盤になる。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Overlapping Kernel Execution with CUDA Streams) ## ストリームの種類と同期範囲 - **レガシーデフォルトストリーム(legacy default stream, stream 0)**: 明示的にストリームを指定しない操作が流れ込む先。自分自身のコマンドを逐次化するだけでなく、他の全てのストリームとの間で暗黙的なグローバルバリアとして働く。stream 0への操作発行は他ストリームの並列性を完全に破壊するため、パフォーマンスが重要な経路では避けるべきとされる。 - **per-thread default stream(PTDS)**: `nvcc --default-stream per-thread`コンパイルまたは`CUDA_API_PER_THREAD_DEFAULT_STREAM=1`環境変数で有効化する。各ホストスレッドが独立した「デフォルトストリーム」を持ち、他スレッドのデフォルトストリームと暗黙同期しない。レガシーストリームと混在させた場合、PTDSストリームはレガシーデフォルトストリームとは引き続き同期する。 - **明示的な非デフォルトストリーム(explicit stream)**: `cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking)`で生成する独立キュー。`cudaStreamWaitEvent()`などで明示的に依存関係を挿入しない限り他ストリームと同期しない。ベストプラクティスとして、性能が重要なカーネル・コピーは常にこの形で作成したストリームへ発行する。 ## Stream-Ordered Memory Allocator `cudaMallocAsync()`/`cudaFreeAsync()`は、メモリ確保・解放の要求をそれを使用するストリームのキューに記録するだけで、他のストリームをブロックしない非同期・ストリーム順序保証つきのアロケータである。対照的に、レガシーの`cudaMalloc()`/`cudaFree()`はブロッキングかつデバイス全体を同期させる操作であり、呼び出しのたびに全ストリームの作業が停止する。PyTorchでは環境変数`PYTORCH_ALLOC_CONF=backend:cudaMallocAsync`でこのアロケータを有効化でき、`cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold)`でメモリプールの解放閾値をチューニングしてOSへの返却頻度と再利用効率のトレードオフを調整できる。LLMのアテンションKV/中間活性化などミニバッチごとに可変長のスクラッチバッファを確保するワークロードでは、このアロケータの使用が並列性維持に必須になる。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Stream-Ordered Memory Allocator, §Using CUDA Streams and Stream-Ordered Memory Allocator with LLMs) ## 細粒度同期: CUDAイベント `cudaStreamSynchronize()`や`cudaDeviceSynchronize()`はホスト側でストリーム全体・デバイス全体の完了を待つ粗粒度な同期であり、不要な停止を招く。CUDAイベント(`cudaEventRecord`/`cudaStreamWaitEvent`)は、あるストリームの特定時点を他のストリームだけが待つ細粒度同期を可能にする。深層学習フレームワークでは、勾配計算完了イベントを通信ストリームが待つことでall-reduce通信を残りの計算とオーバーラップさせる用途で広く使われる。`cudaLaunchHostFunc()`によるホストコールバック登録はGPU→CPU方向の非ポーリング通知を実現するが、コールバック内でCUDA APIを呼ぶとデッドロックしうるため、CPU側処理(メモリ回収等)に限定すべきとされる。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Fine-Grained Synchronization with Events and Callbacks, §Using CUDA Events for Cross-Stream Synchronization) ## マルチGPUでの利用 複数GPU構成では、ローカル計算・P2P転送(`cudaMemcpyPeerAsync`、コピーエンジンのみ使用しSMを消費しない)・[[NCCL]]やNIXLの集団通信(低占有度カーネル)・メモリ確保/解放をそれぞれ別ストリームに分離し、`cudaStreamWaitEvent`で依存関係だけを繋ぐことで、四者を同時にオーバーラップさせられる。1イテレーション全体(フォワード・バックワード・P2Pコピー・NCCL通信・確保/解放・イベント待機)は`cudaStreamBeginCapture`/`cudaStreamEndCapture`でCUDA Graphとしてキャプチャでき、`cudaGraphLaunch()`で再生時のCPU起動オーバーヘッドをほぼゼロにできる([[CUDAGraph]]、詳細は第12章)。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Multi-GPU Compute and Data Transfer Overlap with CUDA Streams) ## 横断的知見 - **[[CUDA]]概念ページが解説するCUDA APIの非同期性(§主要なCUDAランタイムAPI)は、CUDAストリームというオブジェクトを軸に具体化される**: CUDA概念ページは`cudaMallocAsync`/`cudaFreeAsync`を「頻繁な細粒度確保・解放に推奨されるAPI」として第6章から要約しているが、第11章はこれを一歩進め、なぜ`cudaMalloc`が他ストリームを止めるのか(デバイス全体同期)、`cudaMallocAsync`がなぜ止めないのか(ストリームキューへの記録のみ)という機構レベルの説明を与える。CUDA APIトレース研究([[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]])が「非同期性はAPIトレースを難しくする」と観測課題として捉えるのに対し、第11章はストリームという単位そのものを積極的な設計資源として扱う点で一貫している。(Source: [[CUDA]], [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]]) - **NCCL集団通信は「専用の高優先度ストリーム」というCUDAストリームの応用例として位置づけられる**: [[NCCL]]エンティティページ(第4章由来)はNCCLのRing/Tree/CollNet/PATアルゴリズムと運用上の落とし穴を扱うが、第11章はNCCLの集団通信カーネルが低占有度でSMをわずかしか使わず、専用の高優先度ストリームに載せることで計算ストリームをブロックしないという、ストリーム設計の観点からの補足情報を与える。両者を合わせると、NCCLチューニング(第4章)とストリーム設計(第11章)は同じ分散学習パイプラインを異なるレイヤーから最適化していることが分かる。(Source: [[NCCL]], [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]]) ## 未解決の問い - PyTorchのデフォルトキャッシングアロケータ(stream-aware)と`cudaMallocAsync`バックエンドとで、実際のLLM訓練ワークロードにおけるフラグメンテーション削減効果・スループット差はどの程度定量化されているか。 - per-thread default stream(PTDS)を有効化した場合、PyTorchやNCCLのような既存フレームワークの内部ストリーム管理と衝突する可能性はあるか。 - CUDA Graphキャプチャ(`cudaStreamBeginCapture`)と`cudaMallocAsync`によるストリーム順序メモリ確保を組み合わせた際、動的な形状変化(可変長バッチ)にどこまで対応できるか(第12章のCUDA Graphs章で扱われる可能性が高い)。 ## 関連 - 概念: [[CUDA]] / [[CUDAGraph]] / [[Programmatic Dependent Launch]] - エンティティ: [[NVIDIA]] / [[NCCL]] - ソース: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] ## 出典 - Chris Fregly, *AI Systems Performance Engineering*, O'Reilly Media, 2025, Chapter 11.