CUDA C++ 執行模型 (CUDA C++ Execution Model)
重點總覽
CUDA C++ 是 C++ 的延伸，本節以「修改 ISO C++ 標準 [intro.progress] 條款」的方式，定義 host thread、device thread 與 CUDA API 三層的前向進度 (forward progress) 保證。核心目標：對所有 device thread 提供 parallel forward progress [intro.progress.9]，讓既有 C++ 程式更容易用 CUDA C++ 平行化。關鍵在於釐清「誰保證最終會推進、在什麼前提下、推進的傳染範圍多大」。
| 項目 | 重點 |
|---|---|
| concurrent FP [intro.progress.7] | 只要 thread 還沒終止，最終一定會推進（不論他人是否推進） |
| parallel FP [intro.progress.9] | 尚未執行任何步驟前，不保證會開始；一旦執行一步，就升級為 concurrent FP |
| Host thread | 前向進度為 implementation-defined；通用平台應提供 concurrent FP |
| 連動關係 | host 提供 concurrent FP ⇒ CUDA 對 device thread 提供 parallel FP |
| Device thread 傳染範圍 | 一旦某 thread 推進：Cooperative Grid → 整個 grid；否則 → 整個 thread-block cluster |
| CUDA API 保證 | 每次 API 呼叫最終必「return」或「確保至少一個 device thread 推進」 |
| Dependencies | device thread 在所有依賴完成前不得啟動（stream/event/kernel 依賴） |
前向進度的兩種等級
本節大量引用 ISO C++ 草案的進度條款，先記住兩個等級：
- [intro.progress.7] concurrent forward progress：只要 thread 尚未終止，實作保證它最終一定會推進。與其他 thread 是否在推進無關；「最終」意指一段未指定但有限的時間內。
- [intro.progress.9] parallel forward progress：thread 尚未執行任何步驟前，實作不要求它最終會開始推進；但一旦執行了一步，它便提供 concurrent FP 保證。
Note
parallel FP 不規定 thread「何時開始」，這由建立該 thread 的實體決定（例如必須有 API 或排程把它排上去）。這正是後面 device thread「不保證自動開始、要靠 API/依賴觸發」的根源。
換言之，一個本身提供 concurrent FP 的執行體（如 SM）若以任意順序逐一執行一組工作（如各 warp/grid），對這組工作而言即構成 parallel FP——這正對應 GPU 的排程模型。
Host threads（5.8.1）
- host 實作為了執行
main、std::thread、std::jthread而建立的 thread，其前向進度屬於 host 實作的 implementation-defined behavior。 - 通用 (general-purpose) host 實作應提供 concurrent forward progress。
- 連動保證：若 host 提供 concurrent FP [intro.progress.7]，則 CUDA C++ 對 device thread 提供 parallel FP [intro.progress.9]。
Important
整條保證鏈是有條件的：device thread 的 parallel FP 建立在 host 提供 concurrent FP 之上。若 host 平台本身不保證 concurrent FP，下游的 device 保證也不成立。
Device threads（5.8.2）
推進的傳染範圍
一旦某個 device thread 開始推進 (makes progress)：
- 若它屬於 Cooperative Grid：其 grid 內所有 device thread 最終都會推進。
- 否則：其 thread-block cluster 內所有 device thread 最終都會推進。
- 由上可推得：其所在 thread block 內所有 thread 最終都會推進。
Warning
其他 thread-block cluster 的 thread不保證最終會推進。換言之，推進的傳染範圍止於 cluster（非 Cooperative Grid 時），不會自動擴及整個 grid。
某 device thread 推進
│
├─ Cooperative Grid? ──Yes──► 整個 grid 都最終推進
│ ──No───► 整個 thread-block cluster 都最終推進
│ └─（蘊含）整個 thread block 都最終推進
└─ 其他 cluster ─────────────────► 不保證
修改後的 [intro.progress.1]：host vs device thread
標準假設「任何 thread 最終會做下列其一」。CUDA 對 device thread 刪減了若干項（粗體為差異）：
| 編號 | Host thread 最終會… | Device thread 最終會… |
|---|---|---|
| 1 | terminate | terminate |
| 2 | 呼叫 std::this_thread::yield |
（無此項） |
| 3 | 呼叫 library I/O 函式 | 呼叫 library I/O 函式 |
| 4 | 透過 volatile glvalue 存取 | volatile glvalue 存取，但排除 automatic storage duration 物件 |
| 5 | 執行 synchronization 或 atomic 操作 | synchronization 或 atomic read 操作，但排除 automatic storage duration 物件 |
| 6 | 繼續執行 trivial infinite loop | （無此項） |
- 後文以
host.threads.<id>、device.threads.<id>指稱上表各子項。 - device 少了 yield (host.2) 與 trivial infinite loop (host.6)；且 volatile/atomic 對 **automatic storage（區域變數）**不算進度。
- atomic 只認 read（device.4），不認單純 atomic write/fence。
Note
文件指出：對 automatic storage 只做 volatile/atomic 而造成的 undefined behavior，部分屬已知缺陷、未來可能修。但「只做 atomic write 或 fence 不給前向進度」是刻意設計，為效能取捨，因為對它們提供前向進度收益極小卻會整體拖慢效能。
Device 範例（Execution.Model.Device.0–4）
| 範例 | 程式重點 | 結果 | 依據 |
|---|---|---|---|
| Device.0 | 對 cuda::atomic_ref<int, thread_scope_device>（非 automatic storage）spin/store |
grid 最終終止 | device.threads.4 成立 |
| Device.1 | while(true) cuda::std::this_thread::yield(); |
可能無人推進 | device 不支援 host.threads.2 |
| Device.2 | volatile bool True=true; while(True);（automatic storage） |
可能無人推進 | host.threads.4 對 automatic storage 不適用（device.3 例外） |
| Device.3 | cuda::atomic<bool, thread_scope_thread> True=true; while(True.load()); |
可能無人推進 | host.threads.5 對 automatic storage 不適用（device.4 例外） |
| Device.4 | while(true) { /* empty */ } |
可能無人推進 | device 不支援 host.thread.6（trivial infinite loop） |
// Device.0：對 device-scope atomic（非區域變數）loop，最終會終止
__global__ void ex0atomic_ref<int, cuda::thread_scope_device> atom {
if (threadIdx.x == 0) {
while memory_order_relaxed) == 0; // 等待
} else if (threadIdx.x == 1) {
atom.store(1, cuda::memory_order_relaxed); // 設旗標
}
}
CUDA APIs（5.8.3）
核心保證
- 任一 CUDA API 呼叫最終必做其一：return，或確保至少一個 device thread 推進。
- query 函式（如
cudaStreamQuery、cudaEventQuery）不得在沒有任何 device thread 推進的情況下，持續回傳cudaErrorNotReady。 - 推進的那個 device thread 不必與該 API 呼叫「相關」：對某 stream/process 操作的 API，可能讓另一 stream/process 的 device thread 推進。
Tip
粗測前向進度是否符規（非充分）：設環境變數
CUDA_DEVICE_MAX_CONNECTIONS=1 CUDA_LAUNCH_BLOCKING=1 後跑程式，檢查是否仍會終止；若不會，即有 bug。此法能抓到許多（但非全部）前向進度 bug。
API 範例（Execution.Model.API.1–4）
| 範例 | 模式 | 結果 | 關鍵理由 |
|---|---|---|---|
| API.1 | hello_world<<<1,2>>>() 內 __syncthreads() + cudaDeviceSynchronize() |
終止、回 cudaSuccess（前提：無其他 process/stream 的 device thread 一直被選去推進，否則本 grid 可能被餓死） | sync API 確保至少一 thread 啟動 → 同 block 全啟動 → 過 barrier → grid 退出 → sync 解除 |
| API.2 | host 端 while(flag==0) spin，之後才呼叫 cudaDeviceSynchronize() |
可能無人推進 | sync API 只在 device thread 已推進並設旗標後才被呼叫，形成死結 |
| API.3 | 同 API.2，spin 前插入單次 cudaStreamQuery(0) |
可能無人推進 | 單一 query 不保證 device thread 啟動，需重複呼叫 |
| API.4 | 在 spin-loop 內反覆呼叫 cudaStreamQuery(0) |
終止 | 重複 query 保證 device thread 最終推進並設旗標 |
// API.4：在 spin-loop 內反覆 query，保證 producer 最終推進
cuda::atomic<int, cuda::thread_scope_system> flag = 0;
__global__ void producer() { flag.store(1); }
int main() {
cudaHostRegister(&flag, sizeof(flag));
producer<<<1,1>>>();
while (flag.load() == 0) {
(void)cudaStreamQuery(0); // 反覆呼叫才有前向進度保證
}
return cudaDeviceSynchronize();
}
Warning
API.2 vs API.4 是最易混淆點：在 spin-loop 之外、僅呼叫一次 sync/query 不足以解死結；唯有迴圈內持續呼叫 query（或一定會推進整批工作的 sync API）才能保證 device 端最終推進。
Dependencies（5.8.3.1）
- device thread 在其所有依賴完成前不得啟動。
- 依賴常經由 CUDA Stream Commands 建立，可包含對 CUDA Events、CUDA Kernels 完成的依賴等。
| 範例 | 配置 | 結果 | 理由 |
|---|---|---|---|
| Stream.0 | first 與 second 各自不同 stream（無依賴） |
可能無人推進 | CUDA 只保證「某一」thread 推進、不保證是哪一個；可能總挑 second，使其 spin-loop 永不解開（second 餓死 first） |
| Stream.1 | first、second 同一 stream（有順序依賴） |
終止 | stream 依賴強制兩 grid 依序執行，first 先設旗標，second 才解開 |
Important
同一 stream 上的命令依序執行，建立了前向進度所需的依賴；跨不同 stream 且無事件/依賴連結時，排程可任意選擇推進對象，可能造成 starvation。要可靠地讓「等待者」被「設旗標者」解開，需用 stream ordering 或顯式依賴串接。
考試/測驗重點
| 題型 | 關鍵答案 |
|---|---|
| CUDA 對 device thread 提供哪種前向進度？ | parallel forward progress [intro.progress.9] |
| concurrent 與 parallel FP 差別？ | concurrent：未終止前一定推進；parallel：執行第一步前不保證開始，執行一步後升級為 concurrent |
| device 的 parallel FP 前提？ | host 實作提供 concurrent FP [intro.progress.7] |
| host thread 前向進度由誰定？ | implementation-defined；通用平台應提供 concurrent FP |
| 某 device thread 推進後，誰跟著保證推進？ | Cooperative Grid → 整個 grid；否則 → 整個 thread-block cluster（蘊含整個 thread block） |
| 其他 cluster 的 thread 保證推進嗎？ | 不保證 |
| device thread 比 host thread 少了哪兩項「最終會做」的選項？ | std::this_thread::yield（host.2）與 trivial infinite loop（host.6） |
| 對 automatic storage（區域變數）做 volatile/atomic 算前向進度嗎？ | 不算，被排除 |
| device 認哪種 atomic 為進度？ | atomic read；單純 atomic write/fence 不算（刻意設計，為效能） |
while(true) cuda::std::this_thread::yield(); 結果？ |
可能無人推進（device 無 host.2） |
volatile bool True=true; while(True); 結果？ |
可能無人推進（automatic storage 例外 device.3） |
while(true) {} 結果？ |
可能無人推進（device 無 host.6） |
| CUDA API 呼叫的最低保證？ | 最終 return，或確保至少一個 device thread 推進 |
| 推進的 device thread 必須與該 API 相關嗎？ | 不必相關（可在其他 stream/process） |
| query 函式可一直回 cudaErrorNotReady 嗎？ | 不可在無 device thread 推進的情況下持續回傳 |
| 測前向進度的環境變數？ | CUDA_DEVICE_MAX_CONNECTIONS=1 CUDA_LAUNCH_BLOCKING=1（非充分，抓不到所有 FP bug） |
| 單次 vs 反覆 query 的差別（API.3 vs API.4）？ | 單次不保證 device 啟動；spin-loop 內反覆 query 才保證最終推進 |
| 兩 kernel 跨不同 stream、等待者 spin（Stream.0）？ | 無依賴 → 可能 starvation、不終止 |
| 同一 stream 的兩 kernel（Stream.1）？ | 有順序依賴 → 依序執行 → 終止 |
| device thread 何時可啟動？ | 在其所有依賴（stream/event/kernel）完成之後 |