Registers & spilling

01 · The fastest memory you have

Registers are memory with no memory instructions

The bottom of the memory hierarchy isn't shared memory (day 08) — it's the register file. Each thread's local variables live in registers: no addressing, no instructions to fetch them, ~1 cycle of latency. On the 4070 each SM has 65,536 32-bit registers — 256 KB of the fastest storage in the machine — and the compiler parcels them out to your variables at compile time.

You rarely declare registers. You get them by writing code: every live value in a kernel — loop counters, accumulated sums, staged tile elements — occupies one until its last use. The compiler counts them, and its count is a first-class output you should learn to read.

02 · The budget arithmetic

Registers are where occupancy is actually decided

Day 06 said registers and shared memory decide how many warps fit on an SM. Here is that sentence as arithmetic. A kernel using R registers per thread in blocks of T threads consumes R × T registers per block, and an SM hands out blocks until the 65,536-register file is full:

The hardware cap is 255 registers per thread — past that, the compiler has no choice about what to do next. Also worth remembering: blocks round up to whole warps (day 03), and the 4070's SMs hold at most 48 warps (1,536 threads) resident, so both budgets — warps and registers — must agree before a block moves in.

03 · Spilling

When a kernel asks for too much, the compiler quietly lies

What happens when a kernel has more live values than registers? The compiler doesn't fail — it spills: it parks overflowing values in local memory, which despite the name is ordinary off-chip DRAM with an L1/L2 cache in front. Your "register" access becomes a memory access. The kernel still runs, often only somewhat slower — which is exactly why spills go unnoticed for months.

The compiler tells you, if you ask. Compile with -Xptxas -v and read the report:

ptxas info : Compiling entry function 'normalize_row'
ptxas info : Used 96 registers, 0 stack, frame, spill stores, spill loads

// ...and the version that hurts:
ptxas info : Used 255 registers, 40 bytes stack, 24 spill stores, 28 spill loads

Spill stores and loads above zero deserve a raised eyebrow, not necessarily a fix — but rising spill counts alongside falling performance is a signature worth knowing on sight. Day 18's profiler shows the same numbers from the other direction.

04 · Living with the limit

Trade registers for speed on purpose, not by accident

When a kernel is over budget, the fixes are blunt but effective. Keep fewer values live: recompute instead of remembering, shrink the tile the block stages at once, split one do-everything kernel into two simpler passes. Or tell the compiler your priorities with __launch_bounds__:

// "this kernel launches with 256 threads/block,
//  and I need at least 4 blocks resident per SM"
__global__ void
__launch_bounds__(256, 4)
normalize(float* out, const float* in, int n){
    // compiler now caps registers ~64/thread instead of guessing
}

The honest tension: registers are also instruction-level parallelism. More live values means more independent work to interleave while waiting on memory (day 01's lesson at micro scale). Squeezing registers to boost occupancy can starve ILP; hoarding registers for ILP can crush occupancy. There's no constant answer — there's only your kernel, your measurement, and day 18's loop.

Next: day 10 — the caches above it all: L1 and L2 on a GPU, why bandwidth — not compute — is the usual ceiling, and how to tell which one is limiting you.