CUDA Graphsとは｜カーネル起動オーバーヘッドを解消する仕組みと実装

GPUの演算能力を出し切れない原因が、カーネルそのものではなくCPU側の起動処理にあるケースは多くあります。CUDA Graphsは、複数のカーネル起動やメモリ転送を1つの実行グラフとして事前に定義し、まとめてGPUへ投入することで、この起動オーバーヘッドを構造的に排除する仕組みです。CUDA Toolkit 10.0で導入され、推論パイプラインや反復計算のように短いカーネルを大量に呼び出すワークロードで効果が大きくなります。ここでは仕組みの要点から、ストリームキャプチャとExplicit APIの実装コード、効果が出る条件、PyTorch・TensorRT連携までを一次情報に基づいて整理します。

まとめ

• CUDA Graphsは、カーネル起動ごとに発生するCPU側の逐次ディスパッチ（1回あたり数μs〜10μs程度）を、グラフの一括投入で排除する。
• 導入の最短経路はストリームキャプチャ。既存のストリームコードをcudaStreamBeginCaptureとcudaStreamEndCaptureで挟むだけでグラフ化できる。
• 効果が最大化するのは、実行時間100μs未満の軽量カーネルを多数連続実行する場合。カーネルが十分に重いワークロードでは効果は小さい。
• 形状固定の推論は最適。動的な制御フローや可変形状を含む学習ループはグラフ化が難しく、見送るか部分適用に切り替える。
• PyTorchではtorch.cuda.make_graphed_callablesで既存モデルをグラフ化でき、warmupや静的入出力の管理を内部で肩代わりする。

以下、仕組みの理解から実装・運用までを順に見ていきます。

CUDA Graphsが起動オーバーヘッドを解消する仕組み

まず押さえるべきは「なぜ起動処理がボトルネックになるのか」と「グラフがそれをどう回避するのか」の2点です。CUDAの基礎的な用語はCUDAとは？定義と基本的な仕組みについての解説で補完できます。

カーネル起動の数μsが累積するCPU側ディスパッチの構造的限界

カーネルを起動するたび、CPUはドライバを介してGPUへコマンドを送信します。この1回あたりのオーバーヘッドは一般に数μs（目安として5〜10μs程度）で、単発なら無視できます。問題は、数百から数千のカーネルを連続実行するワークロードです。起動コストだけで数ミリ秒から数十ミリ秒に達し、GPUが演算を終えても次のカーネルがディスパッチされるまでアイドルします。カーネルが短いほど影響は大きく、マイクロ秒単位の軽量カーネルを多数実行するパターンではGPU利用率が50%を下回ることもあります。ボトルネックはGPUの演算能力ではなく、ホスト側の逐次処理にあります。

GPU操作をDAGとして事前定義するノードとエッジの設計

CUDA Graphsは、カーネル起動・メモリコピー・ホスト関数呼び出しといったGPU操作を有向非巡回グラフ（DAG）として表現します。各操作がノード、実行順序の制約がエッジです。実行前にワークロード全体の構造をドライバへ渡せる点が本質的な利点で、ドライバは独立ノードの並列実行やスケジューリングを事前に計画できます。従来のストリーム実行では次に来る操作を事前に知れなかったため、ランタイムのスケジューリングコストが残っていました。

キャプチャとリプレイの2段階で実行するEager実行との違い

グラフの実行は「キャプチャ」と「リプレイ」に分かれます。キャプチャ中はAPI呼び出しがGPUで実行されず、順序と依存関係だけがグラフとして記録されます。リプレイでは記録済みグラフをインスタンス化し、繰り返し実行します。API呼び出しのたびにGPUが即時実行するEager実行と違い、CPUの関与は最初の起動に限定されます。キャプチャのコストは初回のみで、2回目以降は極めて低いオーバーヘッドで再生できるため、同じワークロードを何度も回す反復処理で差が出ます。

ストリームキャプチャとExplicit APIによるグラフ構築

グラフの作り方は2通りです。既存コードをそのまま記録するストリームキャプチャと、ノードと依存関係を手で組むExplicit APIで、既存コードの構造と必要な柔軟性で選びます。実務ではまずキャプチャ方式を試し、それで表現できない構造だけExplicit APIに落とすのが定石です。

cudaStreamBeginCaptureで既存コードを最小変更でグラフ化する手順

ストリームキャプチャは、既存のストリームベースコードを3ステップでグラフ化します。cudaStreamBeginCaptureでキャプチャを開始し、通常どおりカーネルと転送を発行し、cudaStreamEndCaptureで締めるとグラフオブジェクトが返ります。カーネル起動コード自体は変更不要で、前後にAPI呼び出しを足すだけで済むため、導入の最小コスト経路になります。

cudaStream_t stream;
cudaStreamCreate(&stream);

cudaGraph_t graph;
cudaGraphExec_t graphExec;

// 1. キャプチャ開始（この区間の発行は実行されず記録される）
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// 2. 通常どおりカーネルと転送を発行
kernelA<<<grid, block, 0, stream>>>(d_in, d_tmp);
kernelB<<<grid, block, 0, stream>>>(d_tmp, d_out);

// 3. キャプチャ終了 → graph が得られる
cudaStreamEndCapture(stream, &graph);

// インスタンス化して繰り返しリプレイ
cudaGraphInstantiate(&graphExec, graph, 0);  // CUDA 12.x の署名（11.x は引数が異なる）
for (int i = 0; i < 1000; ++i) {
    cudaGraphLaunch(graphExec, stream);
}
cudaStreamSynchronize(stream);

cudaGraphCreateで依存関係を明示するExplicit APIとの使い分け

Explicit APIはcudaGraphCreateで空のグラフを作り、cudaGraphAddKernelNodeなどでノードを1つずつ追加します。追加時に依存ノードを明示するため構造を完全に制御できますが、カーネル1つでもパラメータ構造体と依存配列の指定が要り、記述量はキャプチャ方式より大幅に増えます。既存コードの高速化はキャプチャ方式、複雑な依存や条件付きノードを精密に組む新規設計はExplicit API、と切り分けるのが実務的です。

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);

cudaGraphNode_t kernelNode;
cudaKernelNodeParams params = {0};
params.func         = (void *)myKernel;
params.gridDim      = grid;
params.blockDim     = block;
params.kernelParams = args;

// 依存ノードを明示（ここでは先行ノードなし）
cudaGraphAddKernelNode(&kernelNode, graph, NULL, 0, ¶ms);

cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, 0);  // CUDA 12.x の署名（11.x は引数が異なる）

用途別に選ぶ主要グラフノード型

グラフには操作の種類ごとにノード型が用意され、CUDAのバージョン更新で段階的に追加されてきました。CUDA 12.x時点の代表的なノード型は次のとおりです。

異なる種類のノードを同一グラフ内に混在させられるため、転送・初期化・演算をまとめて1回のリプレイで流せます。

キャプチャ中のcudaMalloc制約と回避策

ストリームキャプチャ中は、キャプチャ対象ストリーム以外への同期や、キャプチャに対応しないAPI呼び出しが制約されます。代表例が同期的なcudaMallocで、キャプチャ区間内で呼ぶとエラーになります。回避策は、メモリ確保をキャプチャ前に済ませて確保済みバッファを使い回すことです。動的確保が避けられない場合は、ストリーム順序を保つcudaMallocAsyncやCUDA 11.4以降のメモリ割り当てノード（cudaGraphAddMemAllocNode）を使い、キャプチャ可能な形にします。

効果が出る条件と導入を見送るべきワークロード

CUDA Graphsは万能ではありません。効果はワークロードの性質に強く依存し、条件を外すと導入コストだけがかかります。ここで採用可否を判断する基準を明確にします。

実行時間100μs未満の軽量カーネル連続実行で効果が最大化する閾値

効果が大きいのは、1カーネルの実行時間が起動オーバーヘッドと同程度、目安として100μsを下回る軽量カーネルを多数連続で呼ぶ場合です。起動コスト5〜10μsに対しカーネル実行が100μs未満なら、消せるオーバーヘッドが全体の1割前後を占める計算になり、グラフ化の効き目が顕著になります（100μsは絶対的な境界ではなく、実測で見極める目安です）。起動コストの割合が高いほど、それを消せるグラフ化の効き目が増します。逆にカーネルあたりの実行時間が長ければ、起動コストは相対的に小さく、グラフ化しても改善幅は限定的です。まず対象のカーネルが「短くて多い」かを確認するのが最初の判断です。この構造的な視点はGPUだけでは解決できない低レイテンシ推論の構造的ボトルネックとも重なります。

固定形状の推論ワークロードがグラフ化に最適な理由

入出力テンソルの形状が固定された推論は、CUDA Graphsに最も向くケースです。グラフは記録時のメモリアドレスと形状を前提にリプレイするため、毎回同じ形状・同じバッファで実行できる推論では、一度キャプチャすれば安定して再生できます。バッチサイズが可変でも、パディングで形状を固定化すれば適用できます。反復回数が多いほど、初回キャプチャのコストが回収され費用対効果が高まります。

動的な制御フロー・可変形状でグラフ化を見送るべき失敗パターン

データ依存の条件分岐や、イテレーションごとに変わる可変形状、CPU同期を挟む処理は、グラフ化に適しません。記録した実行順序と前提が崩れるためで、無理にグラフ化すると再インスタンス化が頻発し、かえって遅くなります。学習ループも、勾配累積のタイミングや動的な計算グラフが障壁になりやすい領域です。実行回数が数回しかない小規模ワークロードも、キャプチャコストを回収できず見送るべきです。この場合はグラフ化せず、後述の部分適用に切り替えます。

損益分岐点とNsight Systemsでの効果測定

採用可否は感覚ではなく計測で決めます。Nsight Systemsのタイムラインでカーネル間のCPU待機時間（ギャップ）を見れば、起動オーバーヘッドがボトルネックかどうかが可視化できます。ギャップが目立つならグラフ化の余地が大きく、ギャップがほとんどなければ効果は薄いと判断できます。グラフ化前後で同じ計測を取り、CPUアイドルの消失とエンドツーエンドのレイテンシ短縮を数値で比較してから本番適用します。

グラフ構築からデバッグ・バージョン互換までの実装

実装で詰まりやすいのは、インスタンス化のエラー処理、引数更新、そしてCUDAバージョン差による互換性です。順に押さえます。

cudaGraphLaunchまでの実装手順とインスタンス化エラーの切り分け

基本の流れは、グラフ構築 → cudaGraphInstantiateでインスタンス化 → cudaGraphLaunchで実行 → 不要になったらcudaGraphDestroyとcudaGraphExecDestroyで解放、です。cudaGraphInstantiateが失敗する典型は、キャプチャ中に非対応APIを呼んだ、依存関係が循環している、キャプチャが正しく閉じていない、の3系統です。まずcudaStreamEndCaptureの戻り値でキャプチャ自体の成否を確認し、次にインスタンス化の戻り値コードで原因を切り分けると、デバッグが短くなります。解放は確保順と逆順で行い、実行インスタンスを先に破棄してからグラフ本体を破棄します。

cudaGraphExecUpdateで再インスタンス化を避ける引数更新

入力バッファのアドレスやカーネル引数だけが変わり、グラフの構造（ノードとエッジ）が同じなら、インスタンス化をやり直す必要はありません。cudaGraphExecUpdateで既存の実行インスタンスを差分更新すれば、インスタンス化のコストを避けられます。構造が変わって更新できなかった場合のみ、戻り値を見てインスタンス化し直します。

// カーネル引数だけ変えて再キャプチャ
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernelA<<<grid, block, 0, stream>>>(d_newInput, d_out);
cudaStreamEndCapture(stream, &newGraph);

// 既存 graphExec を作り直さず差分更新（CUDA 12.x の署名）
cudaGraphExecUpdateResultInfo updateInfo;
cudaGraphExecUpdate(graphExec, newGraph, &updateInfo);
if (updateInfo.result != cudaGraphExecUpdateSuccess) {
    // 構造が変わった場合のみインスタンス化し直す
    cudaGraphExecDestroy(graphExec);
    cudaGraphInstantiate(&graphExec, newGraph, 0);
}

CUDAバージョン差でグラフ互換性が崩れる移行リスク

CUDA Graphsの機能セットはToolkitのバージョンで異なります。CUDA 11.0でグラフ更新機能が強化され、11.4でメモリ割り当てノードが加わり、12.0ではcudaGraphInstantiateWithParamsが追加されてインスタンス化時のパラメータ指定が柔軟になりました。バージョンを上げると挙動やAPIが変わり、以前のグラフがそのまま動かないことがあります。実際、cudaGraphInstantiateはCUDA 12.0で引数が5つから3つ（フラグ指定）へ、cudaGraphExecUpdateも結果をcudaGraphExecUpdateResultInfo構造体で受け取る形へ変更され、11.x向けに書いたコードは12.xでコンパイルが通りません（本記事のコード例はいずれもCUDA 12.x基準です）。「cuda バージョンが合わない」トラブルの多くは、こうしたToolkitとドライバ、あるいはビルド時と実行時のバージョン不整合です。移行時はNVIDIA公式の対応表でToolkitとドライバの組み合わせを確認し、グラフのインスタンス化と1リプレイを通す最小検証を先に走らせてから本番に反映します。

PyTorch・TensorRTでの本番展開

フレームワーク側にCUDA Graphs連携が組み込まれているため、C++で低レベルAPIを書かずに恩恵を受けられます。PyTorchとTensorRT／Tritonでの適用パターンを示します。

torch.cuda.make_graphed_callablesで既存モデルをグラフ化する実装

PyTorchではtorch.cuda.make_graphed_callablesにモジュールとサンプル入力を渡すと、グラフ化された呼び出し可能オブジェクトが返ります。内部でwarmupを回し、静的な入出力バッファを維持し、順伝播と逆伝播をCUDA Graphとして再生するため、低レベルのtorch.cuda.graphのように手動で静的テンソルを管理する必要がありません。サンプル入力のrequires_gradは実運用の入力と一致させます。加えて、モデルが持つbufferはrequires_grad=Falseである必要があり、グラフ化した後はパラメータやbufferの増減ができない制約があるため、レイヤ構成が確定してから適用します。

import torch

model = MyModel().cuda()
opt = torch.optim.SGD(model.parameters(), lr=0.1)

// 形状固定のサンプル入力でグラフ化（requires_grad を実運用と一致させる）
sample = torch.randn(N, C, device="cuda")
model = torch.cuda.make_graphed_callables(model, (sample,))

for data, target in loader:
    opt.zero_grad(set_to_none=True)
    out = model(data)          // 順伝播・逆伝播が CUDA Graph で再生される
    loss = loss_fn(out, target)
    loss.backward()
    opt.step()

ネットワークの一部だけが動的でグラフ化できない場合は、その部分をEager実行のまま残し、キャプチャ可能なサブモジュールにだけmake_graphed_callablesを適用する部分グラフ化が使えます。PyTorchの学習基盤の設計はPyTorch Lightningとは何か？概要と誕生の背景を解説も参考になります。

TensorRT・Triton連携で推論スループットを引き上げる構成

TensorRTはビルド設定でCUDA Graphsの統合を有効化でき、生成した推論エンジンの実行をグラフ化してカーネル起動コストを削れます。Triton Inference Serverでは、リクエストのバッチ処理にグラフを適用し、形状を固定したモデルで安定したスループットを得られます。動的バッチサイズには、パディングで入力形状を固定化してからグラフ化する実務上の工夫が有効です。いずれも、形状が固定され反復回数が多い本番推論ほど効果が出ます。適用後はプロファイリングでスループットとレイテンシを計測し、効果を確認してから展開します。

よくある質問（FAQ）

cuda graph apiにはどんな種類がありますか？

大きく2系統あります。既存ストリームを記録するcudaStreamBeginCapture／cudaStreamEndCaptureによるキャプチャ系と、cudaGraphCreateとcudaGraphAddKernelNodeなどでノードを手動構築するExplicit API系です。共通して、実行にはcudaGraphInstantiateでインスタンス化しcudaGraphLaunchで起動します。まずキャプチャ系で試し、複雑な依存構造が必要なときにExplicit APIを併用するのが実務的です。

cudaStreamBeginCaptureを使うと必ず速くなりますか？

速くなるとは限りません。効果は、起動オーバーヘッドがボトルネックになっているワークロードに限られます。実行時間100μs未満の軽量カーネルを多数連続で呼ぶ場合は効果が大きく、重いカーネルが少数のワークロードでは改善はわずかです。導入前にNsight Systemsでカーネル間のCPU待機時間を確認し、ギャップが目立つかどうかで判断してください。

「cuda バージョンが合わない」でグラフが動かないときは？

多くはCUDA Toolkitとドライバ、またはビルド時と実行時のバージョン不整合が原因です。グラフの機能セットはバージョンで異なり、新しいAPIで作ったグラフは古い環境で動きません。NVIDIA公式のToolkitとドライバの対応表で組み合わせを確認し、インスタンス化と1回のリプレイを通す最小検証を先に実行して切り分けてください。

学習ループでもCUDA Graphsは使えますか？

使える場合もありますが、条件が厳しくなります。勾配累積のタイミングや動的な計算グラフ、可変形状が障壁になりやすいためです。ネットワーク全体のグラフ化が難しいときは、torch.cuda.make_graphed_callablesでキャプチャ可能なサブモジュールだけをグラフ化し、動的な部分はEager実行のまま残す部分適用が現実的です。

グラフ化した後にカーネルの入力を変えるにはどうすればよいですか？

構造が同じで入力アドレスや引数だけが変わるなら、cudaGraphExecUpdateで既存の実行インスタンスを差分更新します。インスタンス化をやり直さずに済むため更新コストが小さくなります。ノードやエッジ自体が変わって更新できない場合のみ、戻り値を確認してインスタンス化し直してください。

関連記事

• CUDAとは？定義と基本的な仕組みについての解説
• GPUだけでは解決できない低レイテンシ推論の構造的ボトルネック
• PyTorch Lightningとは何か？概要と誕生の背景を解説
• DGX Sparkとは？スペック・できること・価格をわかりやすく解説