Programmatic Dependent Launch 深入 (PDL Deep Dive)

重點總覽

Programmatic Dependent Launch（PDL）讓「相依的 secondary kernel」可以在它所依賴的 primary kernel 尚未執行完成前，就先在同一條 CUDA stream 中提早啟動，藉此把 secondary kernel 與 primary kernel 不相依的部分重疊執行，並隱藏 secondary kernel 的 launch latency。需要 compute capability 9.0 以上才能提供真正的重疊執行。

Background：為何相依 grid 能提早啟動

CUDA 應用透過在 GPU 上 launch 並執行多個 kernel 來使用 GPU。傳統時間軸上，secondary_kernel 在 primary_kernel 執行完成後才 launch，這種序列化執行通常是必要的，因為 secondary 依賴 primary 產生的結果資料。

• 若 secondary 與 primary 完全沒有相依，本來就可用 CUDA Streams 並行 launch。
• 即使 secondary 相依於 primary，仍有並行潛力：幾乎所有 kernel 都有 preamble section（前導段），會執行「清零 buffer、載入常數」等不依賴 primary 結果的工作。
• 這段 preamble 可以與 primary 並行執行而不影響結果；同時並行 launch 也能把 secondary 的 launch latency 藏在 primary 的執行之中。

傳統序列化：
  primary_kernel  [=========]
  secondary_kernel            [preamble][dependent work]
                              ^ 等 primary 完成才 launch

PDL 重疊執行：
  primary_kernel  [=========]
  secondary_kernel      [preamble][sync][dependent work]
                        ^ 提早 launch，preamble 與 primary 重疊

PDL 對 CUDA kernel launch API 做了改動以達成上述重疊，這些 API 需要至少 compute capability 9.0 才能提供 overlapping execution。

API Description：trigger 與 grid sync

在 PDL 中，primary 與 secondary kernel 被 launch 到同一條 CUDA stream。分工如下：

• Primary kernel：當它準備好讓 secondary 啟動時，所有 thread block 都應呼叫 cudaTriggerProgrammaticLaunchCompletion()。
• Secondary kernel：必須用 extensible launch API（cudaLaunchKernelEx）並帶 launch attribute 啟動；在存取 primary 結果前須呼叫 cudaGridDependencySynchronize()。

__global__ void primary_kernel() {
    // 在啟動 secondary 之前應完成的初始工作
    // 觸發 secondary kernel
    cudaTriggerProgrammaticLaunchCompletion();
    // 可與 secondary kernel 同時進行的工作
}

__global__ void secondary_kernel() {
    // 不相依的工作（preamble，可與 primary 重疊）
    // 會阻塞，直到所有相依的 primary kernel 完成並把結果 flush 到 global memory
    cudaGridDependencySynchronize();
    // 相依的工作（需要 primary 結果）
}

cudaTriggerProgrammaticLaunchCompletion 放在 primary「初始工作之後」，之後的程式碼便可與 secondary 同時進行；cudaGridDependencySynchronize 是 secondary 取得 primary 結果前的屏障。

secondary kernel 用 extensible launch API 並設定 cudaLaunchAttributeProgrammaticStreamSerialization attribute：

cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;
configSecondary.attrs = attribute;
configSecondary.numAttrs = 1;
primary_kernel<<<grid_dim, block_dim, 0, stream>>>();
cudaLaunchKernelEx(&configSecondary, secondary_kernel);

設定 programmaticStreamSerializationAllowed = 1 後，CUDA driver 便可以安全地提早 launch secondary，不必等待 primary 完成與其 memory flush。

Trigger 觸發時機與安全性

CUDA driver 在「所有 primary thread block 都已 launch 並執行了 cudaTriggerProgrammaticLaunchCompletion」時，可以 launch secondary kernel。

• 若 primary kernel 沒有執行 trigger，trigger 會在 primary 所有 thread block 退出後隱式發生（implicitly occurs）。
• 不論哪種情況，secondary 的 thread block 都可能在 primary 寫入的資料變為可見之前就 launch。
• 因此當 secondary 設定了 PDL 時，必須用 cudaGridDependencySynchronize 或其他手段，確認 primary 的結果資料已就緒。

trigger 觸發路徑：
  primary 顯式呼叫 cudaTriggerProgrammaticLaunchCompletion()
        └─► driver 可 launch secondary（所有 primary block 已 launch+trigger）
  primary 未呼叫 trigger
        └─► 隱式：所有 primary thread block 退出後才觸發

Use in CUDA Graphs

PDL 可在 CUDA Graphs 中使用，途徑有二：經由 stream capture，或直接用 edge data。

• 用 edge data 時，在連接兩個 kernel node 的 edge 上設定 cudaGraphDependencyType 為 cudaGraphDependencyTypeProgrammatic。
• 此 edge type 會讓 upstream kernel 對 downstream kernel 中的 cudaGridDependencySynchronize() 可見。
• 此 type 必須搭配 outgoing port：cudaGraphKernelNodePortLaunchCompletion 或 cudaGraphKernelNodePortProgrammatic 二者之一。

// edge data 設定（連接 upstream → downstream kernel node）
cudaGraphEdgeData edgeData;
edgeData.type      = cudaGraphDependencyTypeProgrammatic;
edgeData.from_port = cudaGraphKernelNodePortProgrammatic;        // 或 ...PortLaunchCompletion

把 edge type 設為 Programmatic，等於在 graph 中宣告「downstream 對 upstream 有 PDL 相依」，使 downstream 的 cudaGridDependencySynchronize() 能感知 upstream。

Stream 屬性 ↔ Graph edge 對應

stream capture 下，三種 stream 端設定會轉成對應的 graph edge：

三者的 edgeData.type 皆為 cudaGraphDependencyTypeProgrammatic。重點是 triggerAtBlockStart 的值（0 或 1）決定了 outgoing port 是 Programmatic 還是 LaunchCompletion。

            triggerAtBlockStart = 0  ─►  from_port = cudaGraphKernelNodePortProgrammatic
edgeData
type =      triggerAtBlockStart = 1  ─►  from_port = cudaGraphKernelNodePortLaunchCompletion
Programmatic
            StreamSerialization      ─►  from_port = cudaGraphKernelNodePortProgrammatic

考試/測驗重點

Related Notes

• 04-CUDA-Features/03-CUDA-Graphs-Structure-and-Capture
• 04-CUDA-Features/04-CUDA-Graphs-Updating-and-Conditional
• 04-CUDA-Features/07-Cooperative-Groups-Deep-Dive
• 04-CUDA-Features/16-Work-Stealing-Cluster-Launch-Control
• 03-Advanced-CUDA/02-Advanced-Streams-and-Dependent-Launch
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps