# Programmatic Dependent Launch ## 定義 Programmatic Dependent Launch(PDL)は、同一CUDAストリーム上で依存関係にある2つのカーネル(先行カーネルAと後続カーネルB)を、ホスト(CPU)を介さずに部分的にオーバーラップさせて実行させるCUDAの機構である。先行カーネルが`cudaTriggerProgrammaticLaunchCompletion()`を呼び出した時点で、後続カーネルに必要なデータ(L2キャッシュ・共有メモリ・グローバルメモリへのフラッシュ)が揃ったことをデバイス側で通知し、後続カーネルは`cudaGridDependencySynchronize()`を抜けてすぐに処理を開始できる。ホスト側は`cudaLaunchConfig_t`に`cudaLaunchAttributeProgrammaticStreamSerialization`属性を設定して`cudaLaunchKernelExC()`で後続カーネルを起動することでPDLを有効化する。先行カーネルは全ての必要なデータを生成・同期し終えるまでトリガーを呼んではならない。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Programmatic Dependent Launch) ## カーネル間オーバーラップとしての位置づけ PDLは[[CUDA Stream]]によるカーネル間オーバーラップ(異なるストリーム間での並行実行)とは異なるレイヤーの技術で、**同一ストリーム上の依存関係を持つ2カーネル間**のオーバーラップを対象とする。先行カーネルのエピローグ(ティアダウン処理)と後続カーネルのプロローグを重ねることで、カーネル起動・終了境界のオーバーヘッドそのものを隠す。書籍第11章末では、これをワープ特化(カーネル内パイプライニング)・スレッドブロッククラスタ(ブロック間協調・TMAマルチキャスト)と組み合わせた3層構成の高性能GEMM実装例が示される: (1) カーネル内パイプライニング(ワープ特化によるロード/計算のオーバーラップ)、(2) カーネル間オーバーラップ(PDLによる先行カーネルのプロローグ完了後すぐの後続カーネル起動)、(3) ブロック間協調(スレッドブロッククラスタによるタイルの共有・動的負荷分散)。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]] §Combining PDL and Thread Block Clusters with Warp Specialization) ## 横断的知見 - **PDLは書籍内で初出の機構であり、他ソースからの裏付けはまだない**: 現時点で本概念に言及するソースは[[@2025__OReilly__AI Systems Performance Engineering - Chapter 11 Inter-Kernel Pipelining, Synchronization, and CUDA Stream-Ordered Memory Allocations]]のみ。第12章(CUDA Graphs、デバイス起動型カーネルオーケストレーション)でも関連する可能性が高く、追加ソースが入り次第、突き合わせを追記する。 ## 未解決の問い - PDLはCUDA Graphキャプチャ(`cudaStreamBeginCapture`)と組み合わせて使えるか。カーネル境界のオーバーヘッド隠蔽という目的はCUDA Graphの狙いと重なるため、両者は補完的か代替的か。 - 書籍は「PDL+スレッドブロッククラスタ+ワープ特化」の組み合わせが実務上正当化される場面は限定的だと述べるが、具体的にどのような推論エンジン(超低レイテンシ)がこの組み合わせを採用しているかの実例は示されていない。第12章以降で具体例が出るか確認する。 ## 関連 - 概念: [[CUDA Stream]] / [[CUDA]] / [[CUDAGraph]] - ソース: [[@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.