Stream 回呼、排序與 CUDA Graphs (Callbacks, Ordering, Graphs)

重點總覽

Callback Functions from Streams（從 stream 啟動 host 函式）

CUDA 允許從 stream 內部在 host 端啟動函式。目前有兩個 API：cudaLaunchHostFunc() 與 cudaAddCallback()（即 cudaStreamAddCallback()）。

cudaLaunchHostFunc() 簽名與用法：

cudaError_t cudaLaunchHostFunc(cudaStream_t stream,
                               void (*func)(void *), void *data);

// host 函式本體：簡單的 C 函式
void hostFunction(void *data);  // data 指向使用者自訂的資料結構

• stream：要把 callback 排入的 stream。
• func：要啟動的 callback 函式。
• data：傳給 callback 的資料指標。

執行保證（與 unified memory 搭配時）：

• callback 執行期間，該 stream 被視為 idle，因此函式可使用 attach 到此 stream 的記憶體。
• 函式開始執行的效果，等同於在它之前同步一個記錄於同 stream 的 event；因此會同步在它之前「joined」的各 stream。
• 對任何 stream 加入 device 工作，不會讓 stream 變為 active，直到所有先前的 host 函式與 stream callback 都執行完。
• 函式完成本身不會讓 stream 變 active（除非後面有 device 工作）；可藉由在 stream 結尾的 host 函式發信號來做 stream 同步。

Using cudaStreamAddCallback()（保留供既有程式參考）

cudaError_t cudaStreamAddCallback(cudaStream_t stream,
        cudaStreamCallback_t callback, void* userData, unsigned int flags);

// callback 簽名（與 cudaLaunchHostFunc 不同：多了 stream 與 status）
void callbackFunction(cudaStream_t stream, cudaError_t status, void *userData);

• flags 目前必須為 0（保留未來相容）。
• callback 多接收 status：觸發 callback 的 stream 操作之目前錯誤狀態（可能由先前操作設定）。
• 同樣地，host 函式完成前 stream 不會推進，且不可呼叫任何 CUDA 函式。

Asynchronous Error Handling（非同步錯誤處理）

在 stream 中，錯誤可能源自任何操作（kernel launch、記憶體傳輸）。這些錯誤可能直到 stream 被同步時才回報給使用者（例如等 event 或 cudaStreamSynchronize()）。

cudaStreamSynchronize(stream);
// 看最後錯誤但不清除
cudaError_t err = cudaPeekAtLastError();
if (err != cudaSuccess) {
    printf("Error name: %s\n", cudaGetErrorName(err));
    printf("Error desc: %s\n", cudaGetErrorString(err));
}
// 看最後錯誤並清除
cudaError_t err2 = cudaGetLastError();
// 再 peek/get 一次將回傳 cudaSuccess（已被清除）
cudaError_t err3 = cudaGetLastError();  // == cudaSuccess

兩者皆回傳 cudaError_t。cudaPeekAtLastError() 不清除、cudaGetLastError() 會清除，這是測驗常見對比。

CUDA Stream Ordering（stream 排序語意）

CUDA stream 是 in-order streams：stream 內操作的執行順序 = enqueue 順序，操作不能 leap-frog（越過）其他操作。

• 記憶體操作（如 copy）由 runtime 追蹤，一定會在下一個操作前完成，以確保 dependent kernel 能安全存取被傳輸的資料。
• 少數例外：為效能而放寬語意，例如 programmatic dependent kernel launch（用特殊 attribute 與 launch 機制讓兩 kernel 重疊），或 batched memory copy（runtime 可並行執行非重疊的批次 copy）。

單一 stream（in-order）：
 enqueue:  [memcpy H2D] -> [kernel A] -> [kernel B] -> [memcpy D2H]
 execute:  [memcpy H2D] -> [kernel A] -> [kernel B] -> [memcpy D2H]   同序、不越位

Blocking and non-blocking streams 與 default stream

CUDA 有兩種 stream：blocking 與 non-blocking。名稱有點誤導——blocking/non-blocking 只指它們如何與 default stream 同步。

• cudaStreamCreate() 建立的是 blocking stream（預設）。
• 要建 non-blocking stream，需用 cudaStreamCreateWithFlags() 搭配 cudaStreamNonBlocking。
• 兩者皆用 cudaStreamDestroy() 銷毀。

cudaStream_t stream;
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

Legacy Default Stream（傳統預設 stream / NULL stream）

當 kernel launch 或 blocking cudaMemcpy() 未指定 stream 時，使用 legacy default stream（又稱 NULL stream 或 stream ID 0）。它在所有 host thread 間共享，且為 blocking stream。

• 當操作排入此 default stream 時，它會與所有其他 blocking stream 同步——亦即等所有 blocking stream 完成才能執行。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);          // blocking
cudaStreamCreate(&stream2);          // blocking
kernel1<<<grid, block, 0, stream1>>>(...);
kernel2<<<grid, block>>>(...);       // 進入 default stream
kernel3<<<grid, block, 0, stream2>>>(...);
cudaDeviceSynchronize();

上例中 kernel2（default stream）會等 kernel1 完成，kernel3 會等 kernel2 完成——即使三者原本可並行。改用 non-blocking stream 可避免此同步：

cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);
// default stream 不再與 stream1/stream2 同步，三 kernel 原則上可並行
// 因此不能假設任何執行順序，需顯式同步（如 cudaDeviceSynchronize）

Blocking（default stream 介入）：
  stream1: kernel1 ───┐
  default:            kernel2 ──┐         （等 kernel1）
  stream2:                      kernel3   （等 kernel2）

Non-blocking（default 不同步）：
  stream1: kernel1
  default: kernel2     三者原則上可並行 → 需顯式同步
  stream2: kernel3

Per-thread Default Stream（每 thread 預設 stream）

自 CUDA 7 起，可讓每個 host thread 擁有獨立的 default stream，取代共享的 legacy default stream。啟用方式（擇一）：

• nvcc 編譯選項 --default-stream per-thread，或
• 定義 preprocessor macro CUDA_API_PER_THREAD_DEFAULT_STREAM。

啟用後，每 host thread 的 default stream 不會像 legacy 那樣與其他 stream 同步；於是上面的 legacy 範例會呈現與 non-blocking 範例相同的同步行為。

Explicit Synchronization（顯式同步）

cudaStreamWaitEvent() 是建立跨 stream 依賴的主要工具（見 events 筆記）。

Implicit Synchronization（隱式同步）

• 來自不同 stream 的兩個操作，若之間插入了任何 NULL stream 的 CUDA 操作，便不能並行——除非這些 stream 是 non-blocking（以 cudaStreamNonBlocking 建立）。

提升並行潛力的準則：

• 所有獨立操作應在依賴操作之前發出（issue）。
• 任何同步都應盡可能延後。

Miscellaneous and Advanced（雜項與進階）

Stream Prioritization（stream 優先級）

用 cudaStreamCreateWithPriority() 建立帶優先級的 stream，接收 stream handle 與 priority level 兩參數。數字越小 = 優先級越高，預設優先級為 0。可用 cudaDeviceGetStreamPriorityRange() 查詢裝置/context 的優先級範圍。

int minPriority, maxPriority;
cudaDeviceGetStreamPriorityRange(&minPriority, &maxPriority);
cudaStream_t stream1, stream2;
// cudaStreamDefault → blocking（相對 legacy default stream）；也可用 cudaStreamNonBlocking
cudaStreamCreateWithPriority(&stream1, cudaStreamDefault, minPriority); // 最低優先級
cudaStreamCreateWithPriority(&stream2, cudaStreamDefault, maxPriority); // 最高優先級

CUDA Graphs with Stream Capture（入門）

多個 stream 加上 cudaStreamWaitEvent 的跨 stream 依賴，可表達完整的 DAG（有向無環圖）操作。若某段操作 DAG 需反覆執行多次，可用 CUDA Graphs 降低重複從 host 發出同一串 API 呼叫的延遲與 CPU 開銷：圖只需指定一次，之後可多次執行。

CUDA Graphs 三步驟：

1. Capture：第一次執行時把 graph 捕捉下來（或用 graph API 手動組合）。
2. Instantiate：捕捉後做一次，建立執行所需的 runtime 結構，使後續 launch 盡量快。
3. Execute (launch)：重複執行已 instantiate 的 graph 多次；因結構已就緒，CPU 開銷最小化。

bool graphCreated = false;
cudaGraph_t graph;
cudaGraphExec_t instance;
for (int istep = 0; istep < NSTEP; istep++) {
    if (!graphCreated) {
        cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);  // 開始 capture
        for (int ikrnl = 0; ikrnl < NKERNEL; ikrnl++)
            shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
        cudaStreamEndCapture(stream, &graph);                          // 結束 capture
        cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);         // instantiate（一次）
        graphCreated = true;
    }
    cudaGraphLaunch(instance, stream);   // 重複 launch（多次）
}
cudaStreamSynchronize(stream);

傳統重複 launch：  host: [launchK][launchK]...[launchK]  每次都付 CPU 開銷
CUDA Graphs：      capture+instantiate (一次) → launch graph × N  CPU 開銷最小

Summary of Asynchronous Execution（非同步執行總結）

• 非同步 API 讓我們表達任務並行/重疊；實際達成的並行度取決於硬體資源與 compute capability。
• CUDA 非同步執行的關鍵抽象：streams、events、callback 函式。
• 同步可在 event / stream / device 三個層級進行。
• default stream 是 blocking stream，會與其他 blocking stream 同步，但不與 non-blocking stream 同步。
• 可用 per-thread default stream（--default-stream per-thread 或 CUDA_API_PER_THREAD_DEFAULT_STREAM）避免 default stream 行為。
• stream 可有不同優先級（hint，記憶體傳輸可能不遵守）。
• CUDA 提供降低/重疊開銷的機制：CUDA Graphs、Batched Memory Transfers、Programmatic Dependent Kernel Launch。

考試/測驗重點

Related Notes

• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 02-Programming-GPUs/03-CUDA-Cpp-Sync-and-Workflow
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 02-Programming-GPUs/17-NVCC-Compiler
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps