Register Pressure and the Spill into Local Memory

Every value a thread is holding at a given point in the program must be somewhere. On a GPU the intended somewhere is a register, and the number of registers per streaming multiprocessor is fixed at design time, in the same way that the size of a scratchpad is. The interaction between that fixed budget and the number of threads the programmer wants resident is the least visible of the memory hierarchy’s constraints, because the compiler resolves it silently by moving values into memory and the resulting kernel still computes the right answer.

The register file is a budget, not a resource

The register file of an SM holds on the order of 64K thirty-two-bit registers. A block of threads is allocated registers as a whole, in units finer than a block but coarser than a thread: the hardware divides the file per warp and then per thread, rounding the per-thread allocation up to a small multiple, commonly eight registers. A kernel compiled for 38 registers per thread therefore costs each warp 40, the next multiple of eight, times 32 threads.

Two consequences follow from that arithmetic. First, the number of resident threads is quantized by the register allocation, not by the thread limit alone. On a part with 64K registers and a 2048-thread limit per SM, full occupancy costs 32 registers per thread. A kernel using 33 registers cannot reach it, because 33 rounds up to 40 and 40 times 2048 exceeds the file; it lands at whatever the rounding allows. Second, the occupancy is decided at compile time and can be inspected before the kernel runs, which makes it a performance property that can be reasoned about statically.

The compile-time ceiling on registers per thread is 255, which exists because the encoding of the register operand is a byte. A kernel whose live values genuinely exceed that number cannot keep them all in registers regardless of what the programmer wants.

What happens when the budget is exceeded

When the allocation does not fit the kernel’s demand, the compiler keeps the frequently used values in registers and demotes the rest to local memory. Local memory is a region of device memory, private to each thread, addressed relative to a per-thread base. Loads and stores to it are ordinary memory instructions from the hardware’s point of view; they travel the same path as global accesses and are cached in the same levels. The compiler decides which values to demote by counting uses and weighing loop structure, so the choice is a heuristic rather than a proof. Values inside a loop tend to stay in registers; values that are live across a long stretch of code, or that participate in array indexing, tend to go.

// The index is known at compile time, so the array can live in registers even
// though it is an array. Only the loop that fills it touches memory.
__global__ void smooth(const float *in, float *out, int n, int off) {
    float window[8];
    for (int i = 0; i < 8; ++i) window[i] = in[off + i * n];  // 8 global loads
    float s = 0.0f;
    for (int i = 0; i < 8; ++i) s += window[i];               // unrolled, promoted
    out[off] = s;
}

The same code with a loop bound supplied at run time, or with an index computed from a value the compiler cannot see, turns the array into a local-memory object: the second loop issues eight loads rather than eight register reads, and the source still looks like an array walk. This is a common cause of unexpected local-memory traffic, and it is not register pressure at all. It is an array that the compiler could not promote.

Why the spill traffic still coalesces

If local memory were laid out so that each thread’s private slots occupied a contiguous block, a warp’s access to slot number j would touch thirty-two addresses spread across the width of that block, hundreds of bytes apart, and every one would land in a different sector. The hardware instead interleaves the per-thread regions so that slot j of thread 0 and slot j of thread 1 are adjacent. A warp reading the same slot of all its threads therefore reads one contiguous run, whose sectors are shared among the lanes.

The interleaving is what makes spilling tolerable at all. A spilled access by a full warp costs a few sectors, exactly like a well-formed global access, and the level-one cache can hold the spill region if the footprint per block is small. The traffic is redundant in the sense that the same address would have been a register, but it is not scattered, and the difference between the two is the difference between a modest slowdown and a collapse.

def sectors(addrs, sector=32):
    """Distinct 32-byte sectors touched by a set of addresses."""
    return len({a // sector for a in addrs})

slot = 4                                  # bytes per spilled 32-bit value
warp = range(32)
interleaved = [L * slot for L in warp]     # slot j of lane L at L*slot + j*128
blocked = [L * 64 * slot for L in warp]    # each lane owning 64 private slots

print(sectors(interleaved), sectors(blocked))   # 4 sectors versus 32

The interleaved layout costs a warp the same four sectors as a contiguous global read. A blocked layout, in which each lane owns a contiguous private region, would cost thirty-two, because the lanes would be separated by the size of that region.

Even with the interleaving, the cost is not zero. A spilled value is read at a latency of a cache access rather than at register latency, occupies a slot on the memory pipeline that a shared or global access would otherwise use, and consumes cache capacity. A kernel that spills heavily can become limited by a memory path that the algorithm never intended to use.

Reducing pressure without changing the answer

The levers are the ones the compiler exposes, plus a few source-level rearrangements.

Annotating the launch with the block size, and optionally the number of blocks the programmer wants resident per SM, gives the compiler a target. With a stated intent, it will allocate registers to fit that target and spill the remainder, rather than optimizing for register use alone. A compiler flag that caps registers per thread does the same thing globally for a translation unit. Both are blunt: they buy occupancy with spill traffic, and the exchange rate is not knowable without measuring on the actual part.

Source-level changes are usually better. The register cost of a loop is roughly the number of values live across its body, so splitting a loop that carries many accumulators into two passes over a scratchpad can reduce pressure at the price of one extra traversal. Hoisting loop-invariant computations out and recomputing values that are cheap to derive instead of holding them are the same trade in miniature. Moving per-block constants out of registers and into constant or shared memory is available when the values are uniform across the block, which is the case the constant path handles cheaply, as described in the constant and read-only paths. Choosing a wider register tile in a tiled kernel is a pressure decision as much as a locality decision, and it can push a kernel over the register ceiling that its scratchpad layout was sized against in shared-memory tiling.

One reduction technique is really a compression technique. When a value is only needed for a comparison or a boolean, keeping it as a wider type costs a register for no benefit; when several accumulators of the same kind exist, packing them can free a register at the cost of extra instructions to unpack. These are the same decisions a compiler makes when it rematerializes a value, done by hand because the compiler could not prove them safe.

Occupancy is not the objective

The intuition that more resident warps always help is incomplete. A warp that is waiting is not doing work, and the point of having many warps is to keep the execution units fed while some warps wait. But a kernel whose per-thread work is large and whose memory accesses are already in flight may need very few warps to saturate the machine, and forcing it to a higher occupancy by lowering the register allocation adds spill traffic that costs more bandwidth than the extra warps recover. The counterintuitive result, worked out in Volkov’s analysis of latency hiding, is that some workloads run fastest at deliberately low occupancy with a large register budget per thread.

The practical rule is to treat the register count as a design parameter alongside the tile size and the block size, to inspect the reported occupancy and spill counts the compiler provides, and to change one of the three at a time. The register count is the only one of the three that can silently change between compiler versions, which is why a kernel that is fast when built one way may be slow when built another without its source changing at all.