What is register pressure in CUDA?

The tension between giving each thread more registers and keeping enough threads resident to hide latency.

Two things hide memory latency on a GPU and they want opposite things from the register file. Resident warps hide it by giving the scheduler somebody else to run, and every extra warp costs its threads' registers. Independent work inside one thread hides it by having several loads in flight at once, and that costs registers too, because each value in flight needs somewhere to land. Push either lever and the other gets shorter. Register pressure is that argument, and no rule settles it from the source; the kernel does.

You get two levers to intervene with and they have different blast radii. __launch_bounds__(T, B) is per kernel and it is not a description of your launch, it is a budget: the compiler "first derives the upper limit, L, on the number of registers that the kernel should use ... If the initial register usage exceeds L, the compiler reduces it until it is less than or equal to L. This usually results in increased local memory usage and/or a higher number of instructions." L is the file divided by T times B. -maxrregcount is the blunt version: it "specifies the maximum amount of registers that GPU functions can use" for the whole translation unit, so it reaches kernels you were not thinking about.

The belief worth losing here is that pushing occupancy up is progress. Occupancy is a means of hiding latency and not the thing being bought, so a change that raises it while moving values into local memory can trade a good deal for a bad one. That is exactly what happened on day 17: a bound took the kernel from 72 registers to 64, occupancy went from 75 percent to 100, and the kernel got slower. The 48 bytes of local memory it picked up on the way are the price, and they cost more than the extra warps returned.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Two builds of one kernel body, same input, 256 threads a block, from day 17:

Kernel	Registers	Local memory	Max threads per block	Occupancy	Time
decayWide	72	0 B	896	75.0 %	0.164 ms
decayWideBounded	64	48 B	256	100.0 %	0.227 ms

Same arithmetic, same answer, same launch. Occupancy went from 75 percent to 100 and the clock went the wrong way. The max threads per block column is the other thing the bound did: an unbounded 72-register kernel can still be launched with up to 896 threads a block, and the bounded one refuses anything over 256 with too many resources requested for launch. If you take one habit from this pair, take the habit of reporting occupancy and time together, because either one on its own can be made to look like an improvement.

Diagram

occupancy-stepper, with a time axis beside the slot grid. Sliding the register cap from 72 down to 64 fills the fourth block in the grid while the time bar beside it grows, and the local memory counter under the grid goes from 0 to 48 bytes.

Alt text: "Cutting a kernel from seventy-two registers to sixty-four fills the fourth block on the SM and takes occupancy from seventy-five to one hundred percent, while the kernel's time rises from 0.164 to 0.227 milliseconds and forty-eight bytes of local memory appear."

Related terms

• registers
• register file
• register spilling
• local memory
• occupancy
• latency hiding
• ptxas

Where you meet this

• Day 17, registers, local memory and spills, the lesson that owns this term and measured the trade above.
• Day 10, choosing threads per block, where the same budget is spent from the launch side instead of the compiler side.
• too many resources requested for launch, for the launch that fails because the budget was already spent.
• How to set up CUDA, whose device query prints the register file size the budget divides.

Sources

• CUDA Programming Guide, 5.4.3.2 "Launch Bounds", for the derivation of L and the warning about local memory: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
• NVIDIA CUDA Compiler Driver NVCC, 4.2.7.6, for what -maxrregcount covers: https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html (checked 2026-08-30)
• CUDA C++ Best Practices Guide, 10.2.7.1 "Register Pressure": https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.