3 — The CUDA execution model

Requirements: chapter 2. To run the examples: a GPU, the CUDA Toolkit and Nsight Compute.

CUDA lets you write code for one thread. The hardware does not run one thread. Learning CUDA properly means learning where that fiction stops.

The three levels

1. Thread. The basic unit, with its own registers and program counter.
2. Thread block. Up to 1024 threads on a single SM. Threads in a block can synchronize with __syncthreads() and share data through shared memory.
3. Grid. The array of blocks launched by one kernel call. The hardware assigns blocks to SMs with available capacity (CUDA guide).

There’s also an optional fourth level from compute capability 9.0: thread block clusters. The portable maximum is 8 blocks per cluster, and H100 allows a non-portable 16 (Hopper Tuning Guide).

Here’s how you launch a grid of 4 blocks with 256 threads each:

myKernel<<<4, 256>>>(args);

The warp

An SM executes in warps of 32 threads. The first warp in a block holds threadIdx.x 0 to 31, the next holds 32 to 63, and so on. This is SIMT, single instruction, multiple threads: active threads in a warp execute the same instruction at a given clock (Volta whitepaper).

One practical rule follows: pick block sizes as multiples of 32. Otherwise the last warp runs partly empty.

What happens at a branch

What happens here?

if (threadIdx.x % 2 == 0) {
    do_even_work();
} else {
    do_odd_work();
}

Half the warp wants one path and half wants the other, but they share a single issue slot. This is warp divergence.

Before Volta, a warp shared one program counter and an active mask. The hardware masked off the odd threads, ran the even path, flipped the mask, ran the odd path, and reconverged.

Since Volta, each thread has its own program counter and call stack, and a hardware schedule optimizer groups threads at the same PC. The compiler ptxas still marks reconvergence, using BSSY and BSYNC convergence-barrier instructions (Volta whitepaper, Binary Utilities). This stops diverged threads deadlocking on each other.

The cost hasn’t gone away. When a warp splits, each group uses only part of the warp’s lanes. One preprint measured about a 2× cost for two paths, and observed paths running to completion in turn (arXiv 2607.23402). That’s a report on specific code, not a rule, and NVIDIA doesn’t publish its warp scheduling policy. So measure yours.

Synchronization

Two rules catch new CUDA programmers:

• Since Volta, __syncthreads() is enforced per thread. Code that assumed a warp moves in lockstep is wrong.
• If you need warp-level synchronization, use the *_sync intrinsics or __syncwarp() (CUDA guide).

Hands-on: measure warp divergence

Needs: a GPU, the CUDA Toolkit and Nsight Compute.

The same kernel takes both paths. Only the condition changes: perWarp=1 splits inside each warp, and perWarp=0 splits by block, so each warp agrees with itself.

1. Create div.cu:

   #include <cstdio>
   #include <chrono>
   #include <cuda_runtime.h>
   
   __global__ void diverge(unsigned *d, int iters, int perWarp) {
       int i = blockIdx.x * blockDim.x + threadIdx.x;
       int odd = perWarp ? (threadIdx.x & 1) : (blockIdx.x & 1);
       unsigned x = d[i];
       if (odd) {
           for (int k = 0; k < iters; ++k) x = x * 3u + 1u;
       } else {
           for (int k = 0; k < iters; ++k) x = (x ^ (x >> 3)) + 7u;
       }
       d[i] = x;
   }
   
   int main() {
       const int blocks = 1024, threads = 256, iters = 2000;
       unsigned *d;
       cudaMalloc(&d, sizeof(unsigned) * blocks * threads);
       cudaMemset(d, 0, sizeof(unsigned) * blocks * threads);
       for (int mode = 0; mode < 2; ++mode) {
           auto t0 = std::chrono::steady_clock::now();
           diverge<<<blocks, threads>>>(d, iters, mode);
           cudaDeviceSynchronize();
           std::chrono::duration<double, std::milli> ms =
               std::chrono::steady_clock::now() - t0;
           printf("perWarp=%d: %.3f ms\n", mode, ms.count());
       }
       cudaFree(d);
   }

2. Build it, run it, and inspect the code:

   nvcc -O3 div.cu -o div
   ./div
   cuobjdump -sass div | grep -E 'BSSY|BSYNC|WARPSYNC|BRA'

3. Profile it:

   ncu -o div --section SourceCounters --section InstructionStats ./div
   ncu --import div.ncu-rep --page source --print-source sass

4. Look for: BSSY and BSYNC around the branch, and low divergent branch targets at perWarp=0 but high at perWarp=1. The metrics are smsp__branch_targets_threads_divergent and smsp__branch_targets_threads_uniform (Profiling Guide). Query valid suffixes on your chip first.

5. Expected: time should rise with divergence.

If it doesn’t match: if there’s no BSSY, the compiler probably made the code branch-free. Change the loop bodies and try again. Also note that the second run starts from the first run’s data, so read only the timing, not the values.