Registers and Local Memory
The fastest per-thread storage and its spills.
Registers and Local Memory is a free CUDA Academy lesson on CoddyKit — lesson 1 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
The Fastest Memory You Have
Every thread gets its own private registers, the quickest storage on the chip. They live right next to the math units, so reads cost almost nothing.
Plain Variables Become Registers
When you write a normal local variable in a kernel, the compiler usually keeps it in a register. No special syntax is needed at all. 🙂
__global__ void k() {
int x = 5; // x lives in a register
float y = x * 2.0f;
}Private to One Thread
A register value belongs to exactly one thread. Thread 7 cannot see thread 8's register, which is why per-thread work stays cleanly separated.
Registers Are a Shared Pool
Each streaming multiprocessor has a fixed register file. All resident threads split it, so using fewer registers per thread lets more threads run at once.
When You Run Out
If a kernel needs more registers than are available, the extra values get pushed elsewhere. This event is called a register spill.
Where Spills Go: Local Memory
Spilled values land in local memory, which despite the name is not on-chip. It actually lives in slow off-chip DRAM.
Local Is Still Per-Thread
Local memory is private to each thread just like registers. The difference is purely speed, since local memory is far slower to reach.
Big Local Arrays Spill
Declaring a large per-thread array often forces it into local memory, because there simply are not enough registers to hold every element.
__global__ void k() {
float buf[64]; // likely lives in local memory
}Spills Hurt Performance
Every spill turns a free register access into a slow DRAM trip. Cutting register pressure is one of the easiest optimization wins you can make.
Seeing Your Register Count
You can ask the compiler to report registers per thread. Pass --ptxas-options=-v to nvcc and it prints usage for each kernel.
nvcc --ptxas-options=-v kernel.cuCapping Registers on Purpose
The __launch_bounds__ qualifier hints the compiler to limit registers, trading a little per-thread speed for more threads running together.
__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }Quick Check
You declared a big per-thread array and performance dropped. What likely happened?
Recap: Fast and Private
You learned that registers are the fastest per-thread storage, that the pool is limited, and that overflow spills into slow local memory. Keep register pressure low. 🎯
Frequently asked questions
Is the “Registers and Local Memory” lesson free?
Yes — the full text of “Registers and Local Memory” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Registers and Local Memory”?
The fastest per-thread storage and its spills. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 1 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Registers and Local Memory” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Registers and Local Memory
- Global Memory Tradeoffs
- Constant Memory and Its Cache
- A Mental Model of the Hierarchy