NVIDIA、CUDAにおける「カーネルフュージョン」技術を解説　メモリ帯域のボトルネックと起動オーバーヘッドを削減

要点

• NVIDIAが開発者向けブログで、GPUのメモリ帯域幅を最適化し、起動オーバーヘッドを削減する「カーネルフュージョン」について解説した。

• GPUの高速な演算性能に対してメモリ帯域がボトルネックとなる問題を、複数操作の1カーネル化によって解決する。

• 処理全体の実行順序をまとめる「CUDAグラフ」とは異なるレイヤーの技術であり、両者は相互補完的に併用できる。

• CUDA 13.2で登場したモダンなC++ランタイム「CCCL」を用いて、手動フュージョンの具体的なコード例とパフォーマンス改善効果を示した。

• NVIDIAは2026年7月10日（現地時間）、公式の開発者向け技術ブログ「NVIDIA Technical Blog」において、CUDAにおける「カーネルフュージョン（Kernel Fusion）」を用いたGPUコードの最適化手法に関する解説記事を公開した。本記事では、演算器の処理能力に比べてデバイスメモリの転送速度が遅いために生じるメモリ帯域幅のボトルネックを解消するため、複数のGPU操作を単一のカーネルに統合する具体的なアプローチやその効果が紹介されている。

• GPU向けのプログラムを最適化する手法は数多く存在するが、中でも頻繁に直面する課題がメモリ帯域幅の限界である。現在のGPUは演算処理の速度が極めて高速であるため、広帯域なデバイスメモリ（GPUボード上に搭載された大容量のメインメモリ）を使用していたとしても、メモリ転送が追いつかずにGPUの演算能力（カーネル）を十分に使い切れない事態が発生しやすい。

• このような状況において有効なのが、複数のGPU操作を1つのデバイスカーネル（GPU上で動作するプログラムの最小単位）に結合する「カーネルフュージョン」と呼ばれる手法である。この手法では、処理の途中で発生する中間結果をグローバルメモリ（すべてのスレッドからアクセス可能なGPUのメインメモリ領域）に都度書き出さず、プロセッサに近い高速なレジスタ（演算器直近の記憶領域）に保持したまま次の処理へと受け渡す。これにより、メモリの往復トラフィックを削減すると同時に、カーネルを個別に起動する際の制御オーバーヘッドも低減させることができるという。

• ブログでは、カーネルの本体を結合する今回のフュージョンと、「CUDAグラフ（CUDA Graphs）」による効率化の違いについても解説されている。CUDAグラフは、カーネルの起動やメモリコピー、同期といった一連の処理シーケンスを1つの再利用可能なオブジェクトとしてまとめ、CPU（ホスト）側から1回の呼び出しでまとめて処理を実行できるようにする技術である。これはホスト側のオーバーヘッドをマイクロ秒単位で削減するものの、個々のカーネル本体を結合するわけではないため、グラフ内部のカーネル同士は独立して実行され、中間結果はグローバルメモリを経由する。したがって、カーネルフュージョンとCUDAグラフは異なるレイヤーで作用する補完的な技術であり、これらを組み合わせて使用することが可能であると説明している。

• 具体的な応用例として、配列の各要素の絶対値を求めた上で、そのすべての合計を計算する「sum(abs(x))」という処理が紹介されている。

• 従来の素直な（ナイーブな）実装では、2つの独立したカーネルが使用される。まず、入力配列と同じサイズの一時的な中間バッファに対して絶対値の計算結果を書き出す「abs_kernel」を実行し、その後にその中間バッファを読み込んで単一の数値に縮小（リダクション）する「sum_kernel」を実行する。このアプローチは正しく動作するものの、中間結果をグローバルメモリに一時保存するための大きなデータ転送が発生する。NVIDIAのプロファイリングツール「Nsight Systems」による解析結果では、処理時間全体の65%が「abs_kernel」に、35%が「sum_kernel」に費やされており、データ転送による無駄が全体のパフォーマンスに大きな影響を与えていることが可視化されている。

• この課題を解決するため、ブログでは2つのカーネルを1つの「sum_abs_kernel」に手動で書き換えるアプローチを示している。

• このコード例には、CUDA 13.2で導入された最新のC++ランタイム「NVIDIA CCCL（CUDA C++ Core Libraries）」が採用されており、メモリ範囲を安全に扱う「cuda::std::span」やカーネルの起動を行う「cuda::launch」といったモダンなインターフェースが使用されている。

• 手動でフュージョンされた「sum_abs_kernel」では、絶対値の計算がインラインで実行され、中間結果はグローバルメモリへ書き出されることなくレジスタ内で処理される。さらに、スレッドがデータのサイズに応じてループ処理を行う「グリッドストライド・ループ（Grid-stride loop）」を適用することで、各スレッドが複数の要素を処理しながら部分和を単一のレジスタ上に累積していく。これにより、部分和がグローバルメモリに書き戻されるのを防いでいる。

• また、ブロック単位での合計処理には、高速な共有メモリ（同じスレッドブロック内で共有される高速メモリ領域）を利用する「cub::BlockReduce」が用いられる。最終的に、各ブロックから1つのスレッドだけが「cuda::atomic_ref」を介して、グローバルな結果バッファに対してアトミック（競合が発生しない制御下）に部分和を加算して最終的な合計値を算出する。

• このような手動カーネルフュージョンを適用することで、中間バッファそのものが不要となり、Nsight Systemsのプロファイル上でも2つのカーネル実行が1つに集約され、実行時間が大幅に短縮されることが示されている。