進階 Streams 與相依啟動 (Advanced Streams and Dependent Launch)

重點總覽

More on Streams and Events（總覽）

CUDA stream 的基本語意：同一個 stream 上提交的操作預設被序列化（serialized），前一個未完成下一個不能開始。唯一的例外是新增的 Programmatic Dependent Launch（PDL） 功能。要啟用並行，可用多個 stream，也可用 CUDA Graphs，兩種方式還能結合。

不同 stream 的工作要能並行執行，需滿足特定條件：沒有 event 相依、沒有 implicit synchronization、硬體資源足夠等。

• 不同 stream 的獨立操作之間，若 host 對 NULL stream 提交任何 CUDA 操作，就無法並行——除非這些 stream 是 non-blocking stream。
• non-blocking stream：用 cudaStreamCreateWithFlags() 搭配 cudaStreamNonBlocking flag 建立。為提升並行潛力，建議一律建立 non-blocking stream。

用 Event 表達跨 stream 相依

• 表達 stream 之間的相依，建議使用 non-timing event：呼叫 cudaStreamWaitEvent()，讓某 stream 之後提交的操作等待先前（如另一 stream 上）記錄的 event 完成。
• CUDA event 預設帶 timing 資訊（可用於 cudaEventElapsedTime()）。但若 event 只用來表達相依、不需計時，建議用 cudaEventCreateWithFlags() 搭配 cudaEventDisableTiming flag 建立，以提升效能。

Explicit Synchronization（顯式同步）

四種不同粒度的顯式同步方法：

Table 4：與 host 的顯式同步選項摘要

Implicit Synchronization（隱式同步）

來自不同 stream 的兩個指令，若 host thread 在它們之間發出下列任一操作，就無法並行：

• page-locked（pinned）host 記憶體配置
• device 記憶體配置
• device 記憶體 set（memset）
• 在同一 device 記憶體兩個位址之間的 memory copy
• 對 NULL stream 的任何 CUDA 指令
• L1 / shared memory 配置之間的切換

需要做相依檢查（dependency check）的操作包含：同一 stream 內、位於被檢查 launch 之前的任何指令，以及對該 stream 的任何 cudaStreamQuery() 呼叫。

host thread 時間軸（避免隱式序列化）：

    stream A: kA1 ──────────┐
    stream B: kB1 ──────────┤  ← A、B 可並行
                            │
    [中間插入 cudaMalloc / NULL-stream op / L1↔shared 切換]
                            ▼
    → 此後 A、B 之間被隱式序列化（無法重疊）

Stream Priorities（Stream 優先級）

• 建立時用 cudaStreamCreateWithPriority() 指定相對優先級。
• 可用優先級範圍由 cudaDeviceGetStreamPriorityRange() 取得，順序為 [greatest priority, least priority]。
• runtime 時，GPU scheduler 用 stream priority 決定任務挑選順序，但這是 hint 而非保證。
• 挑選要 launch 的工作時，高優先級 stream 的待處理任務優先於低優先級。
• 高優先級任務不會搶佔（preempt）已在執行的低優先級任務。
• GPU 不會在執行中重新評估 work queue；提高某 stream 優先級不會中斷正在進行的工作。

// get the range of stream priorities for this device
int leastPriority, greatestPriority;
cudaDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, greatestPriority);
cudaStreamCreateWithPriority(&st_low,  cudaStreamNonBlocking, leastPriority);

重點：先取得 [greatestPriority, leastPriority] 範圍，再分別建立最高 / 最低優先級的 non-blocking stream。

Programmatic Dependent Kernel Launch（PDL，相依 kernel 啟動）

CUDA stream 語意讓 kernel 依序執行：第二個 kernel 開始時，第一個 kernel 寫出的相依資料保證已就緒。但常見的情況是——第一個 kernel 已把後續所需資料寫入 global memory，卻還有其他工作要做；而第二個 kernel 在需要那份資料之前也有一段獨立工作。此時可讓兩個 kernel 部分重疊（假設硬體資源足夠），連第二個 kernel 的 launch overhead 都能一起被覆蓋掉。

可達成的重疊程度，除了硬體資源外，取決於 kernel 的具體結構：

• 第一個 kernel 在其執行的哪個時點完成第二個 kernel 所依賴的工作？
• 第二個 kernel 在其執行的哪個時點才開始用到第一個 kernel 的資料？

由於高度依賴具體 kernel，難以完全自動化，因此 CUDA 提供 PDL 讓開發者自行指定兩個 kernel 之間的同步點。

PDL 的三個組成

__global__ void primary_kernel() {
    // Initial work that should finish before starting secondary kernel

    // Trigger the secondary kernel
    cudaTriggerProgrammaticLaunchCompletion();

    // Work that can coincide with the secondary kernel
}

__global__ void secondary_kernel() {
    // Initialization, Independent work, etc.

    // Will block until all primary kernels the secondary kernel is dependent on
    // have completed and flushed results to global memory
    cudaGridDependencySynchronize();

    // Dependent work
}

// Launch the secondary kernel with the special attribute
cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;

cudaLaunchConfig_t config = {0};
config.gridDim         = grid_dim;
config.blockDim        = block_dim;
config.dynamicSmemBytes = 0;
config.stream          = stream;
config.attrs           = attribute;   // Add special attribute for PDL
config.numAttrs        = 1;

// Launch primary kernel (normal launch)
primary_kernel<<<grid_dim, block_dim, 0, stream>>>();
// Launch secondary (dependent) kernel via launch config with the attribute
cudaLaunchKernelEx(&config, secondary_kernel);

重點：primary 用一般 <<<>>> 啟動；secondary 必須透過帶 PDL attribute 的 cudaLaunchConfig_t 用 cudaLaunchKernelEx() 啟動。cudaTriggerProgrammaticLaunchCompletion() 後到 grid 結束之間的工作，可與 secondary 的獨立段重疊。

PDL 時序（arrive / wait）

時間 ───────────────────────────────────────────────►

primary :  [初始工作]──Trigger()──[可與 secondary 重疊的剩餘工作]──┐(flush 到 global mem)
                          │                                        │
                          ▼ 觸發 secondary 可開始                   │
secondary:                [獨立工作]──GridDependencySync()··········▼ 解除阻塞──[相依工作]
                                          ▲ 在此阻塞，直到 primary 完成並 flush 結果
重疊區段：primary 的剩餘工作  ‖  secondary 的獨立工作（含 launch overhead 被覆蓋）

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/03-Batched-Transfers-and-Env-Vars
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps