非同步執行：Streams 與 Events (Async Streams and Events)

重點總覽

什麼是非同步並行執行 (Asynchronous Concurrent Execution)

CUDA 允許多個任務並行 (concurrent)、也就是重疊 (overlapping) 執行，具體可重疊的工作包含：

• host 上的運算
• device 上的運算
• host → device 的 memory transfer
• device → host 的 memory transfer
• 同一個 device 內部的 memory transfer
• 跨多個 device 之間的 memory transfer

並行性是透過非同步介面 (asynchronous interface) 來表達：dispatch 函式呼叫或 kernel launch 會立即返回。非同步呼叫通常在被派發的操作「完成前」就返回，甚至可能在操作「開始前」就返回。返回後 application 可同時去做其他工作；當需要最終結果時，才必須做某種形式的同步 (synchronization) 確認該操作已完成。

三種同步方式

• 先前在「Synchronizing CPU and GPU」介紹的 cudaDeviceSynchronize() 是 blocking 呼叫，等待所有先前發出的工作完成。
• 之所以需要它，正是因為 kernel launch 是非同步的、會立即返回。
• 非同步執行的核心 API 元件就是 CUDA Streams 與 CUDA Events。

非同步重疊（理想時間軸）
Host  : |== CPU work ====================================|
H2D   :   |== copy in ==|
Kernel:                 |==== compute ====|
D2H   :                                   |== copy out ==|
                        ^ 三類工作交錯重疊，總時間 < 各段相加

CUDA Streams

CUDA stream 是一種抽象，讓程式設計者表達一連串操作的序列。它運作得像一條 work-queue：程式把 memory copy、kernel launch 等操作加入佇列，依序執行。

• stream 內的操作依照 enqueue 的順序循序 (sequential) 執行，前面完成並 dequeue 後，下一個才到佇列前端被考慮執行。
• application 可同時使用多個 stream；此時 runtime 會依 GPU 資源狀態，從有工作可做的 stream 中挑選任務執行，達成跨 stream 並行。
• stream 可被指定 priority，作為排程的 hint，但不保證特定執行順序。
• stream 內的 API 呼叫與 kernel launch，相對於 host thread 都是非同步的。
• 可用「等待 stream 清空」或「device 層級同步」兩種方式與 stream 同步。

單一 stream：佇列，先進先出、循序
  stream ─► [ copyH2D ] ─► [ kernel ] ─► [ copyD2H ]  (依序執行)

多個 stream：runtime 視資源並行排程
  stream1 ─► [ kernelA ] ─► [ copyA ]
  stream2 ─► [ kernelB ] ─► [ copyB ]
            └─ A 與 B 可在硬體允許下重疊 ─┘

建立與銷毀 Stream

用 cudaStreamCreate() 建立，初始化一個之後用來識別 stream 的 handle；用 cudaStreamDestroy() 銷毀。

cudaStream_t stream;            // Stream handle
cudaStreamCreate(&stream);      // Create a new stream

// stream based operations ...

cudaStreamDestroy(stream);      // Destroy the stream

• handle 型別為 cudaStream_t。
• 若呼叫 cudaStreamDestroy() 時 device 仍在該 stream 內工作，stream 會先完成所有工作才被銷毀。

在 Stream 啟動 Kernel

照常用 triple-chevron 語法，把 stream 當作第四個額外參數傳入即可。

kernel<<<grid, block, shared_mem_size, stream>>>(...);

• 四個參數依序為 grid、block、shared memory 大小、stream。
• kernel launch 為非同步，呼叫立即返回；launch 成功後 kernel 在 stream 中執行，CPU 與其他 stream 可同時做別的事。

在 Stream 啟動 Memory Transfer

用 cudaMemcpyAsync() 把傳輸排入 stream。它與 cudaMemcpy() 類似，但多一個 stream 參數。

// Copy `size` bytes from `src` to `dst` in stream `stream`
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);

• 此呼叫立即返回（對比 cudaMemcpy() 會阻塞到傳輸完成）；要安全讀取結果須先同步確認完成。
• 其他傳輸函式如 cudaMemcpy2D() 也有非同步變體。

Stream 同步

最簡單的同步就是等待 stream 清空，有 blocking 與 non-blocking 兩種：

// blocking：阻塞直到 stream 內所有工作完成
cudaStreamSynchronize(stream);
// 此後可安全存取 stream 操作的結果

// non-blocking：只「偷看」stream 是否清空
// 清空回傳 cudaSuccess；尚未清空回傳 cudaErrorNotReady
cudaError_t status = cudaStreamQuery(stream);
switch (status) {
  case cudaSuccess:
    std::cout << "The stream is empty" << std::endl;
    break;
  case cudaErrorNotReady:
    std::cout << "The stream is not empty" << std::endl;
    break;
  default:
    // An error occurred - we should handle this
    break;
}

CUDA Events

CUDA event 是插入 stream 的標記 (marker)，像追蹤粒子 (tracer particle) 一樣追蹤 stream 中任務的進度。

• 若只靠 stream 同步，只能知道「stream 是否清空」。例如在 stream 排入兩個 kernel，若某操作只相依於第一個 kernel 的輸出，沒有 event 時只能等整個 stream 清空（此時兩個 kernel 都跑完了）才能安全開始。
• 改用 event：在第一個 kernel 之後、第二個 kernel 之前插入 event，等該 event 到達 stream 前端，就能在 kernel1 完成但 kernel2 尚未開始時安全啟動相依操作。
• 以此方式串接，可在操作與 stream 之間建立相依圖 (graph of dependencies)，此類比直接對應到後續的 CUDA Graphs。
• event 還會保存時間資訊，可用來為 kernel launch 與 memory transfer 計時。

event 表達細粒度相依：
  stream ─► [ kernel1 ] ─► (event) ─► [ kernel2 ]
                            │
                            └─ event 到前端 ⇒ kernel1 已完成
                               即可啟動相依工作（kernel2 仍可能在跑）

建立與銷毀 Event

cudaEvent_t event;
cudaEventCreate(&event);   // Create the event
// ... do some work involving the event ...
cudaEventDestroy(event);   // 不再需要時銷毀

• application 有責任在 event 不再需要時銷毀它。

把 Event 插入 Stream

用 cudaEventRecord() 把 event 排入 stream。

cudaEvent_t  event;
cudaStream_t stream;
cudaEventCreate(&event);
cudaEventRecord(event, stream);   // Insert the event into the stream

計時 (Timing)

event 到達 stream 前端時會記錄一個 timestamp。用兩個 event 夾住一段操作（如 kernel），即可精準量測其執行時間。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);            // 記錄起點
kernel<<<grid, block, 0, stream>>>(...);   // 啟動 kernel
cudaEventRecord(stop, stream);             // 記錄終點

cudaStreamSynchronize(stream);             // 等 stream 完成，兩 event 皆已觸發

float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);   // 取得毫秒數
std::cout << "Kernel execution time: " << elapsedTime << " ms" << std::endl;

• cudaEventElapsedTime(&ms, start, stop) 回傳兩 event 之間的時間，單位為毫秒 (ms)。
• 取時間前必須先確保兩個 event 都已觸發（此處用 cudaStreamSynchronize）。

查詢 Event 狀態

如同查 stream，event 也能用 blocking 或 non-blocking 方式查詢。

blocking — cudaEventSynchronize()：阻塞到該 event 完成。可在第一個 kernel 後、第二個 kernel 前記錄 event，等該 event 即可在 kernel2 完成前就啟動相依工作。

cudaStreamCreate(&stream);
cudaEventCreate(&event);

kernel<<<grid, block, 0, stream>>>(...);
cudaEventRecord(event, stream);            // kernel1 後記錄 event
kernel2<<<grid, block, 0, stream>>>(...);

cudaEventSynchronize(event);               // 等到 event ⇒ kernel1 必已完成
dependentCPUtask();                        // 可能早於 kernel2 結束就開始

cudaStreamSynchronize(stream);             // 等整個 stream ⇒ kernel2 也完成

non-blocking — cudaEventQuery()：立即返回，cudaSuccess 代表 event（及其之前的工作）已完成。下例在 CPU 工作迴圈中偶爾「偷看」kernel1 是否完成，一旦完成就在另一個 stream 啟動 D2H 拷貝，達成 CPU 工作、GPU kernel、D2H 拷貝三者重疊。

cudaMallocHost(&h_data, size);             // pinned host 記憶體，才能真正非同步
cudaStreamCreate(&stream1);                // 處理用 stream
cudaStreamCreate(&stream2);                // 拷貝用 stream
cudaEventCreate(&event);

kernel1<<<grid, block, 0, stream1>>>(d_data, size);
cudaEventRecord(event, stream1);           // kernel1 後插入 event
kernel2<<<grid, block, 0, stream1>>>();

while (not allCPUWorkDone() || not copyStarted) {
  doNextChunkOfCPUWork();
  if (not copyStarted) {
    if (cudaEventQuery(event) == cudaSuccess) {   // kernel1 完成？
      cudaMemcpyAsync(h_data, d_data, size,
                      cudaMemcpyDeviceToHost, stream2);
      copyStarted = true;
    }
  }
}
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

考試/測驗重點

Related Notes

• 02-Programming-GPUs/03-CUDA-Cpp-Sync-and-Workflow
• 02-Programming-GPUs/02-CUDA-Cpp-Memory-Management
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 01-Introduction-to-CUDA/02-Execution-Model-and-SIMT
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps