5.8 CUDA C++ Execution Model（CUDA C++ 执行模型）

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}

     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 Example.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}

     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}