What are registers in CUDA?

The fastest storage on the chip, private to one thread, handed out by the compiler and capped by a fixed pool each SM divides among its resident threads.

A register holds 32 bits with one thread's name on it. You never declare one and you cannot take its address. float carry = 0.0f; in a kernel body is a register unless the compiler decides otherwise, and that decision belongs to ptxas, which allocates them during the build and prints the count when you ask for it. One rule explains most of what follows: no instruction reads register number k for a value of k known only at run time. A scalar always fits. An array indexed by a variable never does, and where it goes instead is local memory, which is DRAM wearing a friendly name.

The count is charged per resident thread, not per launch, because every thread the SM is holding keeps its own copy for as long as its block sits there. A Tesla T4 has 65,536 32-bit registers on each SM and room for 1,024 threads, so 64 per thread is the most a kernel can ask for and still fill the card. Ask for the 65th and threads have to leave. Blocks are the unit the pool is spent in rather than threads, which is why the loss arrives in chunks: at 256 threads a block a 72-register kernel needs 18,432 registers per block, three of those fit inside 65,536, and 768 of the 1,024 slots get used.

The wrong lesson to draw is that fewer registers is better. Registers are where a thread does its arithmetic, so spending fewer means either doing less work per thread or putting values somewhere slower. Day 17 caps one kernel at 64 with __launch_bounds__, gets its occupancy back to 100 percent, and watches it get slower, because the values that stopped fitting went to memory. The target is not a low number. It is whatever register pressure says it should be for that kernel, and the report plus a clock is how you find it.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 17 compiled four kernels from one file and read every register count back off the device with cudaFuncGetAttributes, so the compiler's report and the runtime's view had to agree.

Kernel	Registers per thread	Occupancy	Time
decayNarrow, 4 taps	10	100.0 %	0.038 ms
decayWide, 64 taps	72	75.0 %	0.164 ms

Both ran at 256 threads a block over the same input. The narrow kernel does a fraction of the arithmetic, so the two times are not a race and nothing here says 10 registers beats 72. What the pair does settle is where the ceiling sits: 72 is above the T4's 64 and it costs a quarter of the SM's thread slots, and neither number appears anywhere in the source. The other two kernels in that build put values in local memory, and their rows are on day 17.

Diagram

occupancy-stepper, set to the register axis. One SM as a 1,024-slot grid with a register budget bar beside it. At 10 registers a thread the grid fills; at 64 it still fills and the budget bar reads full; at 72 the bar overflows on the fourth block and 256 slots grey out.

Alt text: "One SM's 65,536 registers against its 1,024 thread slots. At 64 registers a thread every slot is used. At 72 only three 256-thread blocks fit and a quarter of the slots stay empty."

Related terms

• register file
• register pressure
• register spilling
• local memory
• occupancy
• ptxas
• thread

Where you meet this

• Day 2, how a GPU differs from a CPU, where the register file first appears as the reason a GPU keeps thousands of threads in flight.
• Day 10, choosing threads per block, where the block size and the register count start pulling against each other.
• Day 17, registers, local memory and spills, the lesson that owns this term and produced the counts above.
• Day 23, __shfl_down_sync and the mask argument, where lanes read each other's registers without touching memory.
• too many resources requested for launch, the launch failure you get when threads per block times registers per thread runs past the file.

Sources

• CUDA C++ Best Practices Guide, 10.2.7.1 "Register Pressure", on what happens when a kernel wants more registers than are available: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
• CUDA Programming Guide, Table 30, for the 65,536 registers per SM and 1,024 resident threads on compute capability 7.5: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
• NVIDIA CUDA Compiler Driver NVCC, 8.4 "Printing Code Generation Statistics", for the report that names the count: https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html (checked 2026-08-30)
• CUDA Programming Guide, 5.4.3.2 "Launch Bounds", for the one lever that changes the allocation from the source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.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.