CUDA execution model

Warps, not threads

The hardware issues instructions for 32 threads at a time. A 256-thread block is eight warps, and warpSize is 32 on every current part. Two threads share a warp when their threadIdx.x values fall in the same aligned range of 32, which is why block sizes that are multiples of 32 leave no lane idle.

Divergence serialises

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i % 2 == 0) c[i] = f(i);
else            c[i] = g(i);

Half the lanes take each path, so the warp issues f with half its lanes masked, then g with the other half. The cost is the sum of the two paths, not the longer one. A short branch may be if-converted into predicated instructions, which removes the branch but still executes both bodies.

Blocks schedule independently

There is no order between blocks and no barrier between them in the programming model. Block 7 may run first, last, or at the same time as block 3 on another SM. A spin-wait for a flag written by another block deadlocks whenever the waiting blocks fill the device, which is the normal case for a grid larger than one wave.

Latency hiding

Each scheduler holds a pool of resident warps and issues from one that is ready. A warp stalled on a global load is skipped and the cycles go to another warp, which is why occupancy is worth measuring at all: it is the supply of warps available to cover latency. Volta and later add independent thread scheduling, so divergent lanes can make progress separately, but the issue slots are still shared. That is why __syncwarp and the _sync shuffle forms exist; confirm their masks in the guide for your toolkit.

Things that break

A barrier in divergent code is the classic hang, and it is input-dependent: the same kernel returns for a size that keeps every warp convergent and hangs for one that does not. Assuming blocks run in index order produces a kernel that works while the grid fits in a single wave and fails when it does not. Both failures look like a hardware fault and are neither.