SIMT Execution: How Thousands of GPU Threads Compute in Parallel, with Worked Examples

SIMT versus SIMD: the same hardware idea, a different contract

CPU vector units are SIMD: the programmer or the auto-vectorizer packs data into 8- or 16-wide registers and expresses conditionals as masked blends, so the vector width is visible in the code.

SIMT moves that bookkeeping into hardware. Each thread has its own registers, its own index and, logically, its own program counter. The hardware groups 32 consecutive threads of a block into a warp and issues each instruction once for all of them. When all 32 threads want the same instruction, which is the common case, the warp executes at full width. When they disagree, the hardware keeps a per-lane active mask and executes each distinct path with only the relevant lanes enabled. Correctness never depends on the warp size; performance does.

That contract has a useful consequence: a kernel written for one thread is always correct, and making it fast is a matter of making neighbouring threads behave alike. Neighbouring means adjacent threadIdx.x values, because warps are formed from consecutive thread indices, x fastest, then y, then z. See CUDA grids, blocks and threads for how those indices are laid out.

How a warp executes a branch

When a warp reaches a conditional branch, each lane evaluates the condition. If every active lane agrees, the warp simply jumps and nothing is lost. If they disagree, the warp has diverged. The hardware executes one side with the lanes that took it and the others masked off, then the other side with the complementary mask, and the lanes reconverge afterwards.

The consequence is a simple cost model. For a two-way branch inside a warp, time is roughly time(path A) plus time(path B) whenever both sides are taken by at least one lane, rather than the maximum of the two. For an n-way switch on a lane-dependent value, a warp can pay for up to n paths. Loops are the same idea stretched over time: a warp keeps iterating until its slowest lane finishes, and lanes that finished early sit masked for the remaining iterations.

Divergence is a per-warp property: different warps taking different paths cost nothing extra. A branch on a warp, block or kernel-argument value is uniform; one on a lane index or per-thread data may not be.

// Divergence cost, measured the simple way: time the same work with a
// lane-dependent branch and with a warp-uniform branch.
__device__ float heavy(float x, int n) {
    for (int i = 0; i < n; ++i) x = x * 1.000001f + 0.5f;
    return x;
}

__global__ void branch_on_lane(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = i;
    if (threadIdx.x % 2 == 0) v = heavy(v, n);        // half the lanes
    else                      v = heavy(v + 1.f, n);  // the other half
    out[i] = v;                                       // warp runs both paths
}

__global__ void branch_on_warp(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = i;
    if ((threadIdx.x / 32) % 2 == 0) v = heavy(v, n); // whole warp agrees
    else                             v = heavy(v + 1.f, n);
    out[i] = v;                                       // one path per warp
}
// Same instruction count per thread; branch_on_lane takes about twice as long
// when heavy() dominates, because each warp executes both loop bodies.

Worked example: SIMT efficiency in numbers

A useful metric is SIMT efficiency: the fraction of lane slots that do useful work. It equals the average number of active lanes per issued instruction, divided by 32.

Example 1: even-odd branch. In the branch_on_lane kernel above, both paths cost the same, say T cycles of issue time. Each warp runs both, taking 2T, with 16 lanes active in each pass. SIMT efficiency is 16/32, or 50 percent. The branch_on_warp version does exactly the same arithmetic per thread and takes T per warp at 100 percent efficiency, so the kernel runs about twice as fast when the loop dominates.

Example 2: variable-length loops. Suppose each thread processes one token sequence and loops over its length. Within one warp the lengths are 10 for 31 lanes and 200 for one lane. The warp runs 200 iterations. Useful lane-iterations are 31 x 10 + 200 = 510, available lane-iterations are 32 x 200 = 6,400, so efficiency is about 8 percent. Sorting sequences by length before assigning them to threads, the GPU analogue of bucketing batches by length, raises this close to 100 percent because each warp then holds similar lengths.

Predication: when the compiler avoids the branch

For short conditional bodies the compiler often does not emit a branch at all. It computes both results and selects one per lane, or it attaches a predicate to each instruction so that lanes with a false predicate do not write their result. A ternary such as y = (x > 0) ? x : 0.1f * x typically becomes a compare and a select, with no control flow and therefore no divergence to speak of. Both sides are evaluated, which is cheap when they are a handful of instructions.

Predication stops paying when the bodies are long, because then the warp executes every instruction of both sides anyway, or when one side contains memory operations that should not run. Write short conditionals as selects, keep expensive work outside branches, and read the SASS with cuobjdump -sass when a hot branch matters.

Independent thread scheduling since Volta

Before Volta, a warp shared one program counter and one call stack, and reconvergence points were fixed by the compiler. This made some intra-warp synchronization patterns deadlock: if lane 0 spins waiting for a flag that lane 1 will set after the branch, and the hardware chose to run lane 0's path to completion first, lane 1 never ran.

From Volta (compute capability 7.0) onward, each thread keeps its own program counter and call stack, and the scheduler can interleave execution of divergent paths within a warp. Independent scheduling does not make divergence free: at any instant the warp still issues one instruction for one group of lanes, so divergent paths are still serialized and the cost model above still holds.

Independent scheduling changed one programming rule. Code written for older GPUs often assumed that the 32 threads of a warp run in lockstep and therefore exchanged data through shared memory without a barrier, so-called warp-synchronous programming. That assumption is no longer guaranteed. The fix is explicit: use the _sync variants of warp intrinsics, which take a mask of participating lanes, and call __syncwarp() where lanes of a warp communicate through shared memory. The older unsuffixed shuffle and vote intrinsics are deprecated for this reason.

Warp primitives: letting 32 lanes cooperate

Because a warp's lanes share an instruction stream, the hardware can move data between their registers directly. __shfl_sync(mask, v, src) returns lane src's value of v; __shfl_down_sync, __shfl_up_sync and __shfl_xor_sync read from a lane at a relative offset or XOR distance. Votes compress one boolean per lane into a warp-wide result: __ballot_sync(mask, pred) returns a 32-bit word with bit k set if lane k's predicate is true, and __any_sync and __all_sync reduce it further. __activemask() reports which lanes are currently executing.

The mask argument names the lanes that participate, and every lane named must reach the same intrinsic. Passing 0xffffffff from inside a divergent branch where half the lanes are absent is undefined behaviour; the usual pattern is to keep all lanes on the call path and let inactive ones contribute a neutral value, such as zero for a sum.

The canonical use is a reduction. Five shuffle steps with offsets 16, 8, 4, 2 and 1 sum 32 values into lane 0, entirely in registers. A block-level sum does this once per warp, parks the 32 partials in shared memory, and reduces them with one more warp. The kernel below is the building block behind softmax denominators, layer-norm statistics and loss sums in training code.

// Warp-level sum with shuffles: no shared memory, no __syncthreads.
__device__ float warp_sum(float v) {
    unsigned mask = 0xffffffffu;                 // all 32 lanes participate
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(mask, v, offset);  // read v from lane (id + offset)
    return v;                                    // lane 0 holds the warp total
}

// Block sum: each warp reduces, lane 0 of each warp writes a partial,
// then the first warp reduces the partials.
__global__ void block_sum(const float* x, float* out, int n) {
    __shared__ float partial[32];                // up to 1024 threads = 32 warps
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = (i < n) ? x[i] : 0.f;              // pad instead of branching away
    v = warp_sum(v);
    int lane = threadIdx.x % 32, warp = threadIdx.x / 32;
    if (lane == 0) partial[warp] = v;
    __syncthreads();
    int nwarps = (blockDim.x + 31) / 32;
    if (warp == 0) {
        v = (lane < nwarps) ? partial[lane] : 0.f;
        v = warp_sum(v);
        if (lane == 0) atomicAdd(out, v);
    }
}

Ballots enable warp-aggregated operations. In stream compaction, keeping the elements that pass a filter, the warp counts its keepers with __popc on the ballot word, one lane performs a single atomic to reserve space for all of them, and each lane derives its slot from the bits below its own. Contention on the counter drops by up to 32x.

Latency hiding: why a GPU wants thousands of threads

A global-memory load takes hundreds of cycles. A SIMT processor does not stall the whole core while it waits; it switches to another warp whose operands are ready. Each SM keeps many warps resident at once, up to 64 on an H100, with their registers already loaded, so switching costs nothing. Thousands of threads are not there to run simultaneously in the arithmetic sense; they are there so the schedulers always find some warp ready to issue.

This is why a GPU kernel launched with too few threads is slow even when it does little arithmetic, and why register usage per thread matters: registers are divided among resident threads, so heavier threads mean fewer warps to hide latency. GPU occupancy and register pressure cover that trade.

Failure modes

• Data-dependent branches on per-thread values. Filtering, early exits and per-token control flow halve or worse the warp's throughput. Sort, bucket or partition data so each warp sees similar values, or restructure the work so a warp handles one item cooperatively.
• Warp-synchronous code without _sync. Legacy reductions that drop __syncthreads for the last 32 elements and rely on lockstep execution can give wrong answers on Volta and later. Use shuffles or add __syncwarp().
• Wrong masks. Calling a full-mask shuffle where some lanes have exited, for example after an early return for out-of-range indices, is undefined. Keep all lanes alive through the collective and mask the data instead.
• Block sizes that are not multiples of 32. A block of 100 threads occupies four warps, the last with only four live lanes, wasting most of a warp's issue slots.

Operational guidance and trade-offs

When profiling, check three SIMT-related numbers before touching the algorithm: average active threads per warp instruction, which reveals divergence; global load efficiency or sectors per request, which reveals poor coalescing; and achieved occupancy against the theoretical figure. Nsight Compute reports all three, and each points to a different fix. Coalescing is a SIMT effect too: the warp is the unit of memory access, so map adjacent lanes to adjacent addresses, as Coalesced Memory Access explains.

The main design trade-off is between giving each thread independent work, which is simple but vulnerable to divergence and imbalance, and giving each warp a shared task that lanes solve cooperatively with shuffles, which is uniform but needs warp-level code. Most production kernels for training, including fused softmax, layer norm and attention, use the cooperative pattern: one warp or block per row, lanes striding across the row, and shuffles for the reductions. High-level tools such as Triton generate these patterns for you, but the cost model is the same, and knowing it is what lets you predict when a kernel will be slow before you measure it.