Warp Granularity: Why Thirty-Two Threads Share One Program Counter

A kernel launched with a thousand threads does not create a thousand independent instruction streams. The hardware gives the machine a fixed, small number of program counters, and each one drives a datapath that is thirty-two lanes wide. Threads are the programmer’s unit of expression; warps are the machine’s unit of execution, and the mapping between them is fixed by the hardware rather than chosen by the compiler.

Almost everything distinctive about GPU programming follows from that mapping. Divergence exists because one program counter serves thirty-two lanes. Occupancy is counted in warps because warps, not threads, occupy scheduler slots. Memory transactions are formed per warp because the lanes of a warp issue their addresses together. The model is called SIMT, for single instruction, multiple threads, and the name is precise: the instruction stream is single, the data and the register state are per-thread.

One program counter for thirty-two lanes

A block of threads is a three-dimensional grid of indices, but the hardware flattens it before anything executes. The linear index of a thread inside its block is threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y, so threadIdx.x varies fastest. Threads whose linear indices fall in the same run of thirty-two consecutive values form one warp, and the position within that run is the lane index.

/* Warp and lane are derived, not declared: the hardware fixes the mapping. */
__global__ void identity(int2 *out)
{
    unsigned linear = threadIdx.x
                    + threadIdx.y * blockDim.x
                    + threadIdx.z * blockDim.x * blockDim.y;
    out[linear].x = (int)(linear / 32);   /* warp index inside the block */
    out[linear].y = (int)(linear % 32);   /* position inside the warp   */
}

The consequence is that a block of 100 threads contains four warps: three full ones and one with four active lanes and twenty-eight idle. Nothing in the source distinguishes those lanes; the split happens at launch. A block dimension that is not a multiple of thirty-two therefore buys nothing and wastes issue capacity, which is why block sizes are conventionally chosen as multiples of the warp width.

Each lane keeps its own registers. When an instruction names register R4, the register file returns thirty-two separate 32-bit values, one per lane, and the datapath applies the same operation to all of them. The program counter, the instruction being decoded, the branch decision logic, and the scheduling state exist once per warp. A warp is the smallest entity the scheduler can start, stop, or leave resident, and there is no mechanism to advance one lane of a warp while the others stay put, at least not in the architectures that predate per-thread program counters, and even those keep the warp as the scheduling unit.

Why the machine amortizes over a group

The alternative design is a conventional multithreaded core: every thread has its own program counter, its own fetch, its own decode, and the core switches between them. That design scales to a handful of threads because fetch and decode are expensive in area and energy, and the cost is paid per thread. A GPU cannot afford that per-thread overhead if it wants thousands of threads in flight, so it makes the opposite trade: pay fetch and decode once per group and replicate only the wide, cheap parts: the ALU lanes and the register storage.

The trade has a clear axis. Wider groups amortize more control logic per unit of arithmetic, and they coarsen the granularity at which divergent control flow is resolved. Thirty-two is where those two pressures balance for the memory system the GPU is built around. A warp-wide 32-bit load from consecutive addresses covers exactly 128 bytes, built from four 32-byte sectors, and that is the width of an L1 cache line on the architectures in question, so a fully coalesced warp load touches a single line. The register file is built as banks wide enough to supply a full warp per cycle. Narrower groups would multiply scheduler state without improving either of those alignments.

The comparison with CPU SIMD is instructive because the two designs solve the same problem from opposite ends. An AVX-class instruction makes the vector width explicit in the instruction encoding: the compiler must pack eight or sixteen values into a vector register, and scalar code that fails to vectorize runs at scalar speed. SIMT keeps the instruction set scalar-looking. A thread computes one element, and the hardware supplies the vectorization for free by mapping thirty-two threads onto thirty-two lanes. The programmer writes a loop body for one element; the machine runs it for thirty-two. That is why GPU code can be written without intrinsics and still reach full datapath width, and why the same code degrades silently when the lanes disagree.

What the granularity costs at the edges

Granularity is not free at the boundaries. Three edges matter.

The first is the block boundary. Threads are grouped into warps within a block; warps never span two blocks. A block of 33 threads therefore consumes two warp slots and, in the second warp, issues instructions for one useful lane. Since the maximum block size is 1024 threads (exactly 32 warps), the block size also bounds how many warps a single block can contribute, and a grid of small blocks can hit the per-SM cap on resident blocks before it hits the cap on resident warps.

The second edge is the grid tail. A grid of 1000 threads with a block size of 256 fills three blocks and leaves the fourth with 232 threads, so the last block runs seven full warps and one warp of eight lanes. The work still completes; the machine simply issues instructions that mostly do nothing for the final warp.

The third edge is resource allocation. Registers, shared memory, and warp slots are all reserved at warp granularity. A warp whose lanes are mostly inactive still holds a full set of registers, because the register file cannot be partially returned mid-execution. Occupancy accounting, described in what occupancy actually counts, therefore counts a half-empty warp exactly like a full one.

The abstraction the scheduler sees through

Source code presents threads as independent scalar machines. The scheduler sees a coarser object: a warp is either resident in an SM or not, either ready to issue or stalled on a dependency, and its readiness is a property of the whole group. If one lane’s operand is not yet available, the instruction cannot issue for any lane, because there is only one instruction slot. Stall accounting is therefore warp-level even when the underlying cause is a single thread’s long-latency load.

This is also why a warp behaves as an all-or-nothing consumer of issue bandwidth. A warp scheduler selects one warp per cycle and issues one instruction to the datapath; a warp with a single active lane consumes exactly the same issue slot as a warp with thirty-two. The efficiency of that slot is the subject of active-mask accounting.

The one place where the abstraction leaks deliberately is control flow. Because the program counter is shared, a branch whose condition differs across lanes cannot be taken and not taken at the same time. The hardware resolves this by executing both sides in sequence with the inactive lanes masked off, a mechanism built on a per-warp stack of masks that is described in divergence and the reconvergence stack. Everything that makes divergent code expensive starts here: thirty-two lanes share a program counter, so any disagreement about where that counter should go costs time proportional to the number of distinct destinations.

Later architectures relaxed the sharing without abandoning the group. Independent thread scheduling gives each lane its own program counter and call stack, so lanes can be interleaved rather than strictly serialized, but the warp remains the unit that occupies a scheduler slot and the unit that issues. The granularity of execution is still thirty-two; what changed is how much freedom the hardware has in choosing which lanes to advance together. That change is treated in per-lane program counters.

Where the group size is visible to software

Most code never needs the number 32. Code that does needs it exactly. Warp-level primitives (the vote operations that produce a ballot mask, the shuffle operations that exchange register values between lanes) are defined in terms of the lanes of a warp and take a mask argument naming which lanes participate. Reductions that combine values within a warp are written against the assumption that lane i can read lane i+16, then i+8, and so on. Cooperative loading patterns assign each lane a slice of a tile. All of these are correct only if the mapping from thread index to lane is the one the hardware uses, and all of them break when a block size assumption changes the alignment.

The number is also visible in the occupancy arithmetic. An SM that supports 2048 resident threads supports 64 warps, and a kernel whose register use limits it to 32 warps runs at half the theoretical occupancy regardless of how many threads the launch requested. The register file, not the thread count, is often the binding constraint, as covered in register allocation and the occupancy ceiling.

Warp granularity is thus the fixed point of the model. Threads give the programmer a scalar machine to reason about; the warp gives the hardware a vector machine to build. The gap between those two views, visible in divergence, in occupancy, in memory coalescing, and in the cost of a partial warp, is where most GPU performance reasoning happens.