100 DAYS OF GPU Course preview
Day 17Module 2
in-technical-review

Registers, local memory and spills

"Local memory" is the worst-named thing in CUDA. It sounds like storage close to the thread, but it is not. NVIDIA's own guide says so in three sentences:

"Local memory is so named because its scope is local to the thread, not because of its physical location. In fact, local memory is off-chip. Hence, access to local memory is as expensive as access to global memory."

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html section 10.2.4 (checked 2026-08-30)

So float window[64] inside a kernel is either the fastest storage on the card or the slowest, and nothing in the line that declares it tells you which. One compiler flag does. This page shows what that flag prints for four near-identical kernels and how to act when the answer is the slow one.

Where a variable actually lives

A thread's automatic variables have two possible homes. Registers are on the SM, private to one thread, and about as fast as storage gets.

Local memory is in the same DRAM as global memory and costs what global memory costs. The only thing local about it is who can see it.

The compiler decides, and the rule is narrower than people expect. A scalar goes in a register. An array goes in registers only if the compiler can turn every access into a named element, which means every index has to be known at compile time.

The best practices guide names the two cases that fail that test: "large structures or arrays that would consume too much register space and arrays that the compiler determines may be indexed dynamically" (same section, checked 2026-08-30).

That is the first way values leave the register file. The second is spilling: the compiler wanted registers, could not have them, and pushed values out to local memory, storing them on the way out and loading them back on the way in.

Both use the thread-local address space backed by device memory. Those accesses are cacheable in L1 and L2, so they do not guarantee a physical DRAM transaction. The compiler report tells the two cases apart, which is most of what this lesson is for.

The reason the compiler ever runs out is that the register file is a fixed pool on each SM, shared by every warp resident on it. On a Tesla T4 that pool is 65,536 32-bit registers and the SM holds at most 1,024 resident threads (Table 30, "Device and Streaming Multiprocessor (SM) Information per Compute Capability", https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-30).

Divide: 64 registers per thread is the most a kernel can use on a T4 and still fill the card. Every register count on this page is measured against that number, and the program reads both halves of it from the device rather than from a table.

On a T4, scalarizing 64 floats uses 72 registers and 0 local bytes; dynamic indexing allocates 256 local bytes, while launch bounds produce a 48-byte frame with 44-byte spills.

Why an array is not the stack you are used to

In C an array is memory. float v[64] has an address, &v[k] is a pointer, and the index can be anything, because the hardware does the arithmetic at run time. That model is right on a CPU and it is why GPU code surprises people here.

A register has no address. No instruction reads register number k for a runtime k, so an array lives in registers only if the compiler rewrites every access into a different named register, which it can do only when it can read the index off the source.

Give it a literal and the array stops being an array; give it a variable and the array becomes memory. These two lines compile to the same arithmetic and to two different kernels:

acc = acc * decay + window[k];                    // k is a literal after
                                                  // unrolling: a register
acc = acc * decay + window[(k + shift) & 63];     // shift is an argument:
                                                  // memory

Nothing else about the second line is slower. The multiply, add, and answer are the same. Only one operand moved from registers to local memory.

You can also force values out of registers. Fewer registers per thread means more threads resident, and more threads resident is one way a GPU hides latency. A limit increases register pressure, and a limit that is too low causes spills.

One filter, four places to put it

Full program in code/day17-registers/registers.cu. It runs a two-pass decay filter: a forward pass over a window of taps that keeps every intermediate step, then a backward pass over those steps in the opposite order. Reading the window backwards is the point, because it makes every tap stay live between the two loops instead of letting the compiler consume each one as it arrives.

Three rules hold the comparison together.

The three wide kernels compute identical numbers. The one that uses local memory is handed shift = 0, so it does the same arithmetic in the same order as the one that does not. The program compares them against each other as well as against a CPU reference, and reports how many elements differ.

If that count is not zero, they are no longer the same experiment.

The artifact comes out of the compiler, not out of the program. The register counts, the stack frames and the spill byte counts are printed by ptxas during the build, which is why the build line carries -Xptxas -v. No GPU is involved in producing them, which is also why they are the one thing on this page you have to generate yourself rather than read off the transcript below.

Every register count is checked against the device's own view. The program calls cudaFuncGetAttributes and cudaOccupancyMaxActiveBlocksPerMultiprocessor for each kernel, so the table it prints and the report nvcc printed have to agree. Two independent sources for one number is the only reason to trust either.

Here is the kernel the other three are variations of:

__global__ void decayWide(const float* in, float* out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        float window[kWideTaps];
        float carry = 0.0f;
#pragma unroll
        for (int k = 0; k < kWideTaps; ++k) {
            carry = carry * kDecay + in[i + static_cast<size_t>(k)];
            window[k] = carry;
        }
        float acc = 0.0f;
#pragma unroll
        for (int k = kWideTaps - 1; k >= 0; --k) {
            acc = acc * kDecay + window[k];
        }
        out[i] = acc;
    }
}

decayNarrow is that with four taps instead of sixty-four. decayWideIndexed changes one expression in the backward pass, and shift is a kernel argument:

        float acc = 0.0f;
#pragma unroll
        for (int k = kWideTaps - 1; k >= 0; --k) {
            acc = acc * kDecay + window[(k + shift) & (kWideTaps - 1)];
        }

And decayWideBounded is decayWide with a promise attached:

__global__ __launch_bounds__(kThreadsPerBlock, kMinBlocksPerSm) void
decayWideBounded(const float* in, float* out, size_t n) {

Note. __launch_bounds__(T, B) is not a description of how you 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." (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html section 5.4.3.2, checked 2026-08-30.) L here is the register file divided by T times B.

Results

Originally measured on a Tesla T4 with driver 595.84 and CUDA 12.6. The exact shipped program was re-verified with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02. Correctness, register counts, local-byte counts and occupancy were unchanged.

Timing moved: the dynamically indexed version was 7.2x slower than decayWide, against 9.5x in the original capture. Both transcripts remain in evidence; the original 12.6 table follows.

GPU: Tesla T4 (compute capability 7.5)
32-bit registers per SM: 65536
resident threads per SM: 1024
so 64 registers per thread is the most a kernel can use and still fill this card

n = 1049187 elements, 256 threads per block, 4 and 64 taps

kernel             regs   local B   maxTPB  blk/SM    occ %        ms
----------------- ----- --------- -------- ------- -------- ---------
decayNarrow          10         0     1024       4    100.0     0.038
decayWide            72         0      896       3     75.0     0.164
decayWideIndexed     69       256      896       3     75.0     1.556
decayWideBounded     64        48      256       4    100.0     0.227

all four kernels match the CPU reference; decayWideIndexed and decayWide
differ at 0 of 1049187 elements, so the storage change did not change the answer

In the CUDA 12.6 baseline, local memory cost 9.5x. decayWide uses 72 registers and no local memory, and runs in 0.164 ms. decayWideIndexed does the same arithmetic but indexes its taps dynamically, which forces the array out of registers into local memory, 256 bytes of it.

It gives the same answer and occupancy, but takes 1.556 ms. In the CUDA 13.0 re-run those rows were 0.205 and 1.474 ms, a still-large 7.2x.

"Local" is the most misleading word in CUDA. Local memory is not near the thread and it is not fast. It is device memory with a per-thread address, so a spill turns a register read into a global memory round trip, and this is what that looks like on a clock.

__launch_bounds__ trades registers for occupancy and it is not free either. decayWideBounded is capped to 64 registers, which restores 100 percent occupancy from 75.

It also picks up 48 bytes of local memory per thread where the unbounded version had none, and it runs in 0.227 ms, slower than the unbounded 72-register version at 0.164. More occupancy, worse time.

Occupancy is a means and this is the first place in the course where optimising for it directly makes things worse.

The table's local B column is cudaFuncGetAttributes's localSizeBytes, which is the whole stack frame and does not say how it got there. Only the ptxas report splits it, into a stack frame the compiler planned and spill stores and loads it did not want.

On decayWideIndexed the 256 bytes are the window itself, made addressable by a runtime index. On decayWideBounded the 48 bytes appeared when the bound took registers away, which is the shape of a spill, and -Xptxas -warn-spills confirms it.

The original run was captured without -Xptxas -v. The CUDA 13.0 evidence includes it: decayWideBounded used 64 registers with a 48-byte stack frame and 44-byte spill stores and loads, while decayWideIndexed used 69 registers and a 256-byte stack frame with no compiler-reported spills.

The register budget is arithmetic you can do before you compile: 65,536 registers per SM divided by 1024 resident threads is 64 per thread to fill the card. -Xptxas -v tells you where you are against that number, and it needs no GPU, which is why this page has a Compiler Explorer embed and no profiler.

Run it yourself

The report needs no GPU, so Compiler Explorer does the first half of this lesson for free. Paste registers.cu, pick an NVCC compiler rather than an NVRTC one, pin it to a version such as nvcc133 rather than trunk, and set the arguments to -std=c++17 -O3 -arch=sm_75 -Xptxas -v. The four ptxas info blocks land in the compiler output pane, not the program output pane, and they are there whether or not you press Execute.

Execute gets you the rest. That runner is a Tesla T4, so sm_75, and the program is sized to finish inside its 20 second cap. A free Colab T4 or any card you own does the same job with more room:

nvcc -std=c++17 -O3 -arch=sm_75 -Xptxas -v -o registers registers.cu

Exercise

Answer these three questions from the sections above.

  1. One kernel reports a stack frame and zero spills. Another reports a smaller stack frame and non-zero spills. Both put values in local memory. Which one is the compiler telling you it lost an argument, and what is the other one telling you?
  2. Your kernel reports 65 registers and you launch it with 256 threads per block on a T4. How many blocks fit on an SM? How many fit at 64? What did the 65th register cost?
  3. You add __launch_bounds__(1024) to a kernel you launch with 256 threads per block, and change nothing else. Predict what happens to its register count and to its speed.

Time: 20 to 30 minutes. Submit: three answers, plus the four register counts from your own build.

Check: the answers are in the reveal below, and each names the division or the documentation line that settles it rather than appealing to how the hardware feels. Question 2 is arithmetic. Questions 1 and 3 you check by making the edit and rereading the report, which is why both edits are written up in the repo's README.

Hint 1

All three use the same division. Write down the pool size and divisor before answering any of them.

Hint 2

For question 2, divide twice and round the block count down, then turn blocks into threads. For question 3, the first argument of a launch bound is not a description of your launch.

Solution

1. The one with spills. A stack frame with no spills means the compiler decided up front that something needs an address, which a dynamically indexed array does, so it never tried registers. Spills are "stores and loads done on stack memory which are being used for storing variables that couldn't be allocated to physical registers" (https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html section 8.4, checked 2026-08-30): it wanted registers and could not have them.

The tempting wrong answer is that both are spills, since both show up as local memory. The difference is which one more room would fix.

2. Three blocks, then four. A T4 has 65,536 registers per SM. At 256 threads and 64 registers a thread a block needs 16,384, so exactly four fit: 1,024 threads, every warp slot there is.

At 65 a block needs more than 16,384 and three fit: 768 threads. "A bit slower, it is one register" misses that the block is the unit of allocation, so the cost arrives in whole blocks.

3. Its register count falls to 64 or below, and its speed does not follow from that. The bound is a budget derived from the numbers you wrote, not from the launch you perform: one block of 1,024 threads has to fit, and on a T4 that is 65,536 over 1,024.

A kernel that wanted more gets cut and the remainder becomes spills. The obvious wrong answer is "nothing changes, I did not change the launch". The launch never was the input.

The rule to keep: a register count limits how many threads can be resident, and a spill trades register space for memory traffic.

Pitfalls

Your kernel got slower after you added a small array. A runtime index made it addressable, so it went to local memory, which is DRAM. Make every index a compile-time literal, usually by giving the loop a constant bound and unrolling it, or move the data to shared memory.

Day 13 covers the second option.

You added __launch_bounds__ and it got slower. The bound is a register budget and you set it too tight, so ptxas fitted the kernel by spilling. Build with -Xptxas -warn-spills and it says so: ptxas warning : Registers are spilled to local memory in function, followed by the mangled kernel name and the two byte counts.

Add -Xptxas -Werror and that becomes a build failure rather than something you notice next quarter.

The launch fails with too many resources requested for launch. Threads per block times registers per thread exceeded the register file, or you launched a kernel with more threads per block than its own __launch_bounds__ allows. The error page ranks the causes; the quick version is fewer threads per block, or a bound.

You capped registers with -maxrregcount and an unrelated kernel got worse. That option is file-wide: it "specifies the maximum amount of registers that GPU functions can use" (https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html section 4.2.7.6, checked 2026-08-30), all of them. __launch_bounds__ is per kernel.

You read a report from the wrong architecture. Register counts are per -arch, and a build with several -gencode targets prints one block per target. Compiled for sm_90, this program's four kernels report different register counts, stack frames and spill bytes than they do for sm_75.

You chase 100 percent occupancy. Occupancy is one of two ways to hide latency, and the other is more independent work per thread, which costs registers. Day 45 tunes one kernel two ways to the same speed at different occupancies.

Latency hiding is the thing you actually want.

Go deeper

Next

Day 18 is the last stop in module 2: the one read-only path that is fast for a reason none of the others are. Day 45 then spends the register budget on purpose instead of by accident, and day 46 opens the SASS so you can watch a spill happen instruction by instruction. The GPU kernel engineer hub maps where the rest of it goes.