Register Pressure and Spills

Registers are the fastest storage a thread has, but each SM holds only so many. How heavily a kernel uses them is called register pressure.

Every resident thread draws from one register file per SM. The more registers each thread needs, the fewer threads can stay resident together.

High register use directly caps how many warps fit on an SM. That drop in occupancy can leave the hardware with too little work to hide latency.

When a thread needs more registers than exist, the compiler moves extra values out. This overflow is a register spill into slower memory.

Spilled values go to local memory, which lives in off-chip DRAM. Every spill replaces a free register access with a slow round trip.

Ask the compiler to report usage. Adding --ptxas-options=-v to nvcc prints registers per thread and any spill bytes for each kernel.

The report lists spill stores and loads. Any nonzero spill count is a warning sign that your kernel is paying for slow local memory traffic.

You can set a ceiling at compile time. The --maxrregcount flag limits registers per thread, trading a little speed for higher occupancy.

A per-kernel hint is often better. __launch_bounds__ tells the compiler your block size so it can budget registers for the occupancy you want.

Often you can simply hold fewer values at once. Recomputing a cheap result or narrowing variable scope reduces how many registers stay live.

ILP and unrolling raise pressure, while caps raise occupancy. The art is finding the balance that runs fastest, and only the profiler can tell you where it is.

Where do values go when a kernel runs out of registers?

You learned that high register pressure cuts occupancy and can cause spills to slow DRAM. Measure usage, cap registers, and let the profiler guide you. 🎯