What is register spilling in CUDA?
The compiler running out of registers and pushing values into local memory, so every later use of them costs a trip to device DRAM.
ptxas has a fixed number of registers to hand a kernel and a program that may want more. When it wants more, the values it cannot keep get written out and read back, and NVIDIA's compiler manual names them for what they are: spills are "stores and loads done on stack memory which are being used for storing variables that couldn't be allocated to physical registers". The store is a STL, the load is an LDL, and both go to local memory, which is the same DRAM as global memory with a per-thread address on it.
Two different things put bytes in local memory and the report tells them apart, which is the distinction worth carrying away. A stack frame with zero spills means the compiler decided up front that something needs an address, which a dynamically indexed array does, so it never tried registers at all. Spill counts mean it tried and lost. Both show up as one number in cudaFuncGetAttributes's localSizeBytes, and only the -Xptxas -v report splits the total into a planned stack frame and the spill store and spill load byte counts. The practical difference is what a fix would be: more registers helps the second case and does nothing for the first.
Getting the build to tell you costs nothing, and this is the part most people skip. Add -Xptxas -warn-spills and a spilling kernel prints ptxas warning : Registers are spilled to local memory in function, followed by the mangled name and the two byte counts. Add -Xptxas -Werror and that warning fails the build, which turns a silent regression into a red pipeline. No GPU is involved in producing any of it, so Compiler Explorer will do it for a file you paste in.
Measured
One kernel body, three builds, one input, 256 threads a block, all from day 17. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75.
| Kernel | Registers | Local memory | Occupancy | Time |
|---|---|---|---|---|
decayWide |
72 | 0 B | 75.0 % | 0.164 ms |
decayWideBounded, capped at 64 |
64 | 48 B | 100.0 % | 0.227 ms |
decayWideIndexed, runtime index |
69 | 256 B | 75.0 % | 1.556 ms |
Only the middle row is a spill. decayWideBounded is the unbounded kernel with __launch_bounds__ attached, and its 48 bytes appeared the moment the bound took registers away, which is the shape a spill has. The bottom row's 256 bytes are the tap array itself, made addressable by an index the compiler could not read at compile time, so it is a stack frame rather than a spill and giving it more registers would not move it. The local memory column comes from cudaFuncGetAttributes and cannot distinguish the two on its own; the run behind this table was captured without -Xptxas -v, so the split is the artifact you generate in your own build.
Code
From code/day17-registers/registers.cu. The bound is the whole edit that produced the 48 bytes, and it is attached to a kernel body that is otherwise character for character the unbounded one.
__global__ __launch_bounds__(kThreadsPerBlock, kMinBlocksPerSm) void
decayWideBounded(const float* in, float* out, size_t n) {
kThreadsPerBlock is 256 here. The pair (T, B) is a promise about residency, and ptxas turns it into a register ceiling of the file divided by T times B. Write a bound your kernel cannot afford and the spills are the compiler keeping your promise.
Related terms
Where you meet this
- Day 17, registers, local memory and spills, the lesson that owns this term and produced the three rows above.
- Day 10, choosing threads per block, where a larger block quietly tightens the same budget.
too many resources requested for launch, the failure at the other end of the same budget, when ptxas could not fit the kernel at all.- Learn CUDA without a GPU, because the spill report is a build artifact and needs no card.
Sources
- NVIDIA CUDA Compiler Driver NVCC, 8.4 "Printing Code Generation Statistics", for the definition of a spill and a worked example of the report: https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html (checked 2026-08-30)
- CUDA C++ Best Practices Guide, 10.2.4 "Local Memory" and 10.2.7.1 "Register Pressure", for where spilled values land and what they cost: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- CUDA Programming Guide, 5.4.3.2 "Launch Bounds", which states that a tight bound "usually results in increased local memory usage": 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. Two different named people have to sign this entry before it leaves draft.