5.8. CUDA C++ 执行模型
CUDA C++ 旨在为所有设备执行线程提供并行前向进展 [intro.progress.9]，以便使用 CUDA C++ 对现有 C++ 应用程序进行并行化。
[intro.progress]
[intro.progress.7]：对于提供并发前向进展保证的执行线程，实现确保只要该线程尚未终止，它最终就会取得进展。
[注意 5：无论其他执行线程（如果有）是否曾经或正在取得进展，这一点都适用。最终满足此要求意味着，这将在未指定但有限的时间内发生。注意结束]
[intro.progress.9]：对于提供并行前向进展保证的执行线程，如果它尚未执行任何执行步骤，则不要求实现确保该线程最终会取得进展；一旦此线程执行了一个步骤，它就提供并发前向进展保证。
[注意 6：这并未规定何时启动此执行线程的要求，通常由创建此执行线程的实体来规定。例如，一个提供并发前向进展保证、并以任意顺序逐一执行某个任务集中的任务的执行线程，满足这些任务的并行前向进展要求。注意结束]
CUDA C++ 编程语言是 C++ 编程语言的扩展。本节记录了对当前ISO 国际标准 ISO/IEC 14882《编程语言 C++》草案中[intro.progress]一节的修改和扩展。修改过的各节会明确指出，差异以粗体显示。所有其他各节均为新增内容。
5.8.1. 主机线程
由主机实现创建、用于执行main、std::thread和std::jthread的执行线程所提供的前向进展，是由主机实现定义的行为[intro.progress]。通用主机实现应提供并发前向进展。
如果主机实现提供并发前向进展 [intro.progress.7]，则 CUDA C++ 为设备线程提供并行前向进展 [intro.progress.9]。
5.8.2. 设备线程
一旦某个设备线程取得进展：
如果它属于一个协作网格，则其网格中的所有设备线程最终都必须取得进展。
否则，其线程块簇中的所有设备线程最终都必须取得进展。
[注意：不保证其他线程块簇中的线程最终会取得进展。注意结束。]
[注意：这意味着其线程块内的所有设备线程最终都必须取得进展。注意结束。]
将[intro.progress.1]修改如下（修改内容以粗体显示）：
实现可以假定任何主机线程最终都会执行以下操作之一：
终止，
调用库 I/O 函数，
通过 volatile 泛左值执行访问，
执行同步操作或原子操作，或者
继续执行平凡无限循环（[stmt.iter.general]）。
实现可以假定任何设备线程最终都会执行以下操作之一：
终止，
调用库 I/O 函数，
通过 volatile 泛左值执行访问，但所指示的对象具有自动存储期的情况除外，或者
执行同步操作或原子读取操作，但所指示的对象具有自动存储期的情况除外。
[注意：设备线程相对于主机线程的某些当前限制是我们已知的实现缺陷，我们可能会随着时间的推移修复这些缺陷。例如，设备线程最终仅对具有自动存储期的对象执行 volatile 操作或原子操作时会产生未定义行为。然而，设备线程相对于主机线程的其他限制是有意作出的选择。这些选择使某些性能优化成为可能，而如果设备线程严格遵循 C++ 标准，这些优化就无法实现。例如，为最终仅执行原子写入或栅栏的程序提供前向进展，会降低整体性能，却几乎没有实际益处。注意结束。]
对[intro.progress.1]的修改所导致的主机线程与设备线程之间前向进展保证差异的示例。
以下示例分别使用“host.threads.<id>”和“device.threads.<id>”来引用上述主机线程和设备线程的实现假设中逐项列出的子条款。
1// Example: Execution.Model.Device.0
2// Outcome: grid eventually terminates per device.threads.4 because the atomic object does not have automatic storage duration.
3__global__ void ex0(cuda::atomic_ref<int, cuda::thread_scope_device> atom) {
4 if (threadIdx.x == 0) {
5 while(atom.load(cuda::memory_order_relaxed) == 0);
6 } else if (threadIdx.x == 1) {
7 atom.store(1, cuda::memory_order_relaxed);
8 }
9}
1// Example: Execution.Model.Device.1
2// Allowed outcome: No thread makes progress because device threads don't support host.threads.2.
3__global__ void ex1() {
4 while(true) cuda::std::this_thread::yield();
5}
1// Example: Execution.Model.Device.2
2// Allowed outcome: No thread makes progress because device threads don't support host.threads.4
3// for objects with automatic storage duration (see exception in device.threads.3).
4__global__ void ex2() {
5 volatile bool True = true;
6 while(True);
7}
1// Example: Execution.Model.Device.3
2// Allowed outcome: No thread makes progress because device threads don't support host.threads.5
3// for objects with automatic storage duration (see exception in device.threads.4).
4__global__ void ex3() {
5 cuda::atomic<bool, cuda::thread_scope_thread> True = true;
6 while(True.load());
7}
1// Example: Execution.Model.Device.4
2// Allowed outcome: No thread makes progress because device threads don't support host.thread.6.
3__global void ex4() {
4 while(true) { /* empty */ }
5}
5.8.3. CUDA API
CUDA API 调用最终必须返回，或者确保至少一个设备线程取得进展。
CUDA 查询函数（例如cudaStreamQuery、cudaEventQuery等）不得在没有设备线程取得进展的情况下持续返回cudaErrorNotReady。
[注意：该设备线程不必与 API 调用“相关”，例如，对一个流或进程执行操作的 API 可以确保另一个流或进程上的设备线程取得进展。注意结束。]
[注意：测试程序是否符合 CUDA API 前向进展要求的一种简单但不充分的方法是，在设置以下环境变量后运行程序：
CUDA_DEVICE_MAX_CONNECTIONS=1 CUDA_LAUNCH_BLOCKING=1，然后检查程序是否仍会终止。如果不会，则程序存在错误。此方法并不充分，因为它无法发现所有前向进展错误，但确实可以发现许多此类错误。注意结束。]
CUDA API 前向进展保证的示例。
1// Example: Execution.Model.API.1
2// Outcome: if no other device threads (e.g., from other processes) are making progress,
3// this program terminates and returns cudaSuccess.
4// Rationale: CUDA guarantees that if the device is empty:
5// - `cudaDeviceSynchronize` eventually ensures that at least one device-thread makes progress, which implies that eventually `hello_world` grid and one of its device-threads start.
6// - All thread-block threads eventually start (due to "if a device thread makes progress, all other threads in its thread-block cluster eventually make progress").
7// - Once all threads in thread-block arrive at `__syncthreads` barrier, all waiting threads are unblocked.
8// - Therefore all device threads eventually exit the `hello_world`` grid.
9// - And `cudaDeviceSynchronize`` eventually unblocks.
10__global__ void hello_world() { __syncthreads(); }
11int main() {
12 hello_world<<<1,2>>>();
13 return (int)cudaDeviceSynchronize();
14}
1// Example: Execution.Model.API.2
2// Allowed outcome: eventually, no thread makes progress.
3// Rationale: the `cudaDeviceSynchronize` API below is only called if a device thread eventually makes progress and sets the flag.
4// However, CUDA only guarantees that `producer` device thread eventually starts if the synchronization API is called.
5// Therefore, the host thread may never be unblocked from the flag spin-loop.
6cuda::atomic<int, cuda::thread_scope_system> flag = 0;
7__global__ void producer() { flag.store(1); }
8int main() {
9 cudaHostRegister(&flag, sizeof(flag));
10 producer<<<1,1>>>();
11 while (flag.load() == 0);
12 return cudaDeviceSynchronize();
13}
1// Example: Execution.Model.API.3
2// Allowed outcome: eventually, no thread makes progress.
3// Rationale: same as Example.Model.API.2, with the addition that a single CUDA query API call does not guarantee
4// the device thread eventually starts, only repeated CUDA query API calls do (see Execution.Model.API.4).
5cuda::atomic<int, cuda::thread_scope_system> flag = 0;
6__global__ void producer() { flag.store(1); }
7int main() {
8 cudaHostRegister(&flag, sizeof(flag));
9 producer<<<1,1>>>();
10 (void)cudaStreamQuery(0);
11 while (flag.load() == 0);
12 return cudaDeviceSynchronize();
13}
1// Example: Execution.Model.API.4
2// Outcome: terminates.
3// Rationale: same as Execution.Model.API.3, but this example repeatedly calls
4// a CUDA query API in within the flag spin-loop, which guarantees that the device thread
5// eventually makes progress.
6cuda::atomic<int, cuda::thread_scope_system> flag = 0;
7__global__ void producer() { flag.store(1); }
8int main() {
9 cudaHostRegister(&flag, sizeof(flag));
10 producer<<<1,1>>>();
11 while (flag.load() == 0) {
12 (void)cudaStreamQuery(0);
13 }
14 return cudaDeviceSynchronize();
15}
5.8.3.1. 依赖项
设备线程在其所有依赖项完成之前不得启动。
由依赖项产生的 CUDA API 前向进展保证的示例
1// Example: Execution.Model.Stream.0
2// Allowed outcome: eventually, no thread makes progress.
3// Rationale: while CUDA guarantees that one device thread makes progress, since there
4// is no dependency between `first` and `second`, it does not guarantee which thread,
5// and therefore it could always pick the device thread from `second`, which then never
6// unblocks from the spin-loop.
7// That is, `second` may starve `first`.
8cuda::atomic<int, cuda::thread_scope_system> flag = 0;
9__global__ void first() { flag.store(1, cuda::memory_order_relaxed); }
10__global__ void second() { while(flag.load(cuda::memory_order_relaxed) == 0) {} }
11int main() {
12 cudaHostRegister(&flag, sizeof(flag));
13 cudaStream_t s0, s1;
14 cudaStreamCreate(&s0);
15 cudaStreamCreate(&s1);
16 first<<<1,1,0,s0>>>();
17 second<<<1,1,0,s1>>>();
18 return cudaDeviceSynchronize();
19}
1// Example: Execution.Model.Stream.1
2// Outcome: terminates.
3// Rationale: same as Execution.Model.Stream.0, but this example has a stream dependency
4// between first and second, which requires CUDA to run the grids in order.
5cuda::atomic<int, cuda::thread_scope_system> flag = 0;
6__global__ void first() { flag.store(1, cuda::memory_order_relaxed); }
7__global__ void second() { while(flag.load(cuda::memory_order_relaxed) == 0) {} }
8int main() {
9 cudaHostRegister(&flag, sizeof(flag));
10 cudaStream_t s0;
11 cudaStreamCreate(&s0);
12 first<<<1,1,0,s0>>>();
13 second<<<1,1,0,s0>>>();
14 return cudaDeviceSynchronize();
15}