GPU Register Pressure

Registers are the compiler's budget, not yours

A thread’s registers are allocated by the compiler back end, ptxas, when it lowers PTX to SASS. It builds the kernel’s dataflow graph, computes a live range for every value — the span from its definition to its last use — and colors those ranges onto a finite set of physical registers. Values that are never live at the same moment share a register. The peak number of simultaneously live values, not the number of variables you typed, sets the count.

Two hard boundaries frame it. The instruction set caps a single thread at 255 32-bit registers. And the SM’s register file is a fixed on-chip array shared by every resident thread, so the count you land on divides directly into how many warps fit. That divisor arithmetic — granule rounding, block-granularity admission, the interplay with shared memory and warp slots — belongs to the occupancy discussion. Here it matters only as the consequence: the compiler’s allocation decision is also an occupancy decision, made for you, at compile time.

What actually drives demand up

Pressure rises whenever more values must be alive at once. Four sources cover most real kernels.

Long live ranges. A value loaded early and used late stays live across everything in between. Hoisting loads out of a loop ‘for efficiency’ often just extends live ranges.

Loop unrolling. The biggest single multiplier. Eight unrolled iterations means eight loads in flight and eight partial results live at once — that is the point, since it exposes instruction-level parallelism, but it is paid for in registers. An aggressive #pragma unroll can double a kernel’s count.

Inlining. Device functions inline by default, merging the callee’s live ranges into the caller’s. A deep call chain collapses into one function whose peak pressure is the sum of its parts.

Large accumulator tiles. A blocked GEMM or attention kernel that holds a register tile of results deliberately spends dozens of registers per thread to maximize reuse. That is the design, not an accident.

Spilling — local memory is global memory in disguise

When the allocator runs out it spills: it evicts values, stores them, and reloads them at their next use. The destination is what CUDA calls local memory, and the name is the most expensive piece of misleading terminology in the programming model. Local memory is not a separate on-chip resource beside shared memory. It is a per-thread private region carved out of the same device DRAM that backs global memory, with the same latency and the same bandwidth budget.

A spill is therefore an off-chip round trip wearing the costume of a local variable. In SASS it appears as LDL and STL rather than the LDG/STG of a global access, but the data path underneath is the same one.

Two things soften it. Local memory is laid out interleaved across threads, so the same spilled variable in consecutive threads is contiguous and coalesces naturally. And spills go through L1 and L2 like any other access, so a small, hot spill set may live in L1 and never reach DRAM. A few spill bytes in a cold path are harmless; spilling inside the inner loop is not.

Local arrays and the dynamic-index trap

There is one way to lose registers that has nothing to do with running out of them. Declare an array inside a kernel — float acc[8]; — and the compiler will happily keep it in registers, but only if it can rewrite every access into a distinct named register. That requires every index to be resolvable at compile time.

The moment you index it with a value the compiler cannot fold — a loop counter it declines to unroll, a thread-dependent offset, a value read from memory — the array goes to local memory in its entirety. The register file is not addressable at runtime; there is no indexed-register-read instruction to lower it to. One dynamic subscript demotes the whole array, and the kernel silently acquires traffic that never appears as a source-level memory operation.

The fixes follow directly: give the loop a compile-time trip count so it unrolls fully, keep bounds constant so indices fold, or accept that the array is genuinely dynamic and move it to shared memory, where indexed access is cheap and on-chip. Check this first whenever a small kernel reports a large stack frame.

Reading the numbers from the compiler

You do not have to guess at any of this. Ask ptxas to report and it prints its allocation for every kernel it compiles:

nvcc -arch=sm_90 -Xptxas -v -c kernel.cu

ptxas info : Used 72 registers, 96 bytes stack frame,
             84 bytes spill stores, 84 bytes spill loads,
             8192 bytes smem, 396 bytes cmem[0]

Read it in that order. Used N registers is the per-thread count that feeds the occupancy divisor. A non-zero stack frame means something landed in local memory — usually a demoted array. Spill stores and loads are the allocator admitting defeat.

One caveat trips people up: those spill figures are static byte counts over the compiled code, not dynamic execution counts. Eighty-four bytes of spill in a rarely taken branch costs nothing; the same eighty-four bytes inside a loop that runs a million times is your kernel. The compiler tells you spills exist. Only the profiler tells you whether they are hot.

Reading the numbers from the profiler

Nsight Compute closes the loop the static output cannot. Its launch statistics report the registers per thread actually used, alongside the theoretical and achieved occupancy that number produced — which is how you confirm registers, and not shared memory or block size, are the binding constraint.

Memory workload analysis is where spills become visible as traffic. Local loads and stores get their own row with their own sector counts and hit rates, so you can see how much of the spill traffic L1 absorbed and how much reached L2 and DRAM. That distinction is the whole question: spill traffic that stays in L1 is a modest instruction-count tax, while spill traffic that misses to DRAM competes with your real data for the same bandwidth.

Scheduler statistics complete the picture. A kernel that spills badly shows a large share of stalls on long-scoreboard memory dependencies even though the source contains no obvious global access; pair that with a non-zero stack frame and the diagnosis is clean.

The first lever — __launch_bounds__

__launch_bounds__ is a kernel attribute that tells the compiler what you intend to launch, so it can allocate accordingly:

__global__ void
__launch_bounds__(256, 4)   // maxThreadsPerBlock, minBlocksPerMultiprocessor
myKernel(const float* in, float* out) { ... }

The first argument is a promise: no launch of this kernel will exceed 256 threads per block. That alone helps, because without it the compiler must allocate conservatively enough that the kernel stays launchable at the architecture’s maximum block size — and a kernel whose register count makes the requested block size impossible fails at launch, not at compile time.

The second argument is the real lever. Asking for at least four resident blocks per SM hands the compiler a register budget derived from the register file, and it optimizes within that budget: rematerializing values, shortening live ranges, backing off unrolling, and spilling only as a last resort. Because it is per kernel and expressed as an occupancy goal rather than a raw number, it is almost always the lever to reach for first.

The second lever — -maxrregcount

-maxrregcount=N is the blunt instrument: a command-line flag that caps registers per thread for every kernel in the compilation unit. It predates __launch_bounds__ and behaves accordingly.

Its problems are scope and expressiveness. A translation unit rarely holds one kernel, and the count that suits a memory-bound elementwise kernel is usually wrong for a compute-bound GEMM sitting beside it — the flag caps both identically. It also asks the wrong question. You do not care about the number 64; you care about resident warps. __launch_bounds__ lets you state the goal and have the compiler compute the number for the architecture you are targeting, which stays correct when you retarget to a part with a different register file.

Where the two collide, the attribute wins: __launch_bounds__ overrides the global flag for that kernel. Reasonable uses for the flag remain — a quick experiment under a tighter cap, or a build that genuinely wants a uniform ceiling — but it is a sweeping default, not a tuning knob.

Why forcing registers down is often a net loss

The seductive thing about capping registers is that the benefit is arithmetic and immediate: halve the count, roughly double the resident warps, watch the occupancy number rise. The cost is not arithmetic and does not appear until you measure.

The asymmetry is enormous. A register operand is part of the instruction — no address, no transaction, effectively no latency. A spilled operand is a memory instruction with hundreds of cycles of potential latency, cache pressure, and bandwidth cost. Buying occupancy with spills trades the cheapest access in the machine for one of the most expensive, and the extra warps then have to hide a stall you created yourself.

This is also why high-performance GEMM and attention kernels deliberately run near the 255-register ceiling at low occupancy. They hide latency with instruction-level parallelism and a large register tile rather than with warp count, and capping them destroys the design.

A diagnosis order

Work the problem in this sequence and you will not waste a tuning cycle.

1. Get the facts. Compile with -Xptxas -v. Note registers per thread, stack frame, and spill bytes. No spills and no stack frame means register pressure is not your problem; stop here.

2. Confirm registers are the binding limit. Check achieved occupancy and which resource caps it. If shared memory or block size is the ceiling, cutting registers buys nothing.

3. Ask whether occupancy is the bottleneck at all. A compute-bound kernel with good ILP can sit near peak at low occupancy.

4. Attack demand at the source before reaching for a flag. Hunt the dynamically indexed local array first, then reduce the unroll factor, move loads closer to their uses, recompute cheap values instead of keeping them live, and consider splitting a fat kernel in two.

5. Only then use __launch_bounds__, targeting a block count rather than a register number.

6. Re-measure end to end. Occupancy is a means; kernel duration is the result.