Constant Memory and Its Cache
Broadcasting read-only values cheaply.
Constant Memory and Its Cache is a free CUDA Academy lesson on CoddyKit — lesson 3 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.
A Space for Read-Only Values
Constant memory is a small region built for data that every thread reads but none of them ever writes during a kernel.
Declaring It
You mark a global-scope array with the __constant__ qualifier. It lives outside any function, visible to all your kernels.
__constant__ float weights[256];It Is Genuinely Small
Total constant memory is only 64 KB on typical GPUs. It is meant for coefficients and parameters, not for large datasets.
Filling It from the Host
You write to constant memory from the CPU using cudaMemcpyToSymbol, since kernels themselves cannot modify it.
cudaMemcpyToSymbol(weights, h_w,
256 * sizeof(float));Backed by a Special Cache
Reads flow through a dedicated constant cache on each multiprocessor. A cached value comes back almost as fast as a register.
The Broadcast Superpower
When a whole warp reads the same address, the hardware does one fetch and broadcasts it to all 32 threads in a single step. 📡
Same Address Is the Key
The win depends on uniform access. If threads in a warp read different constant addresses, those reads serialize and the speedup is lost.
Reading It Looks Normal
Inside a kernel you read constant memory like any array. The compiler quietly routes the access through the fast constant cache.
__global__ void k(float *out, int i) {
out[i] = weights[0] * 2.0f;
}A Perfect Fit for Filters
Convolution kernels and lookup tables shine here, because every thread reuses the same small set of coefficients over and over.
Set Once, Reuse Often
You typically upload to constant memory a single time, then launch many kernels that all read those unchanging values cheaply.
When Not to Use It
Skip constant memory if your data is large or if threads read scattered addresses. In those cases plain global memory serves you better.
Quick Check
When does constant memory give its biggest speedup?
Recap: Small, Cached, Broadcast
You learned constant memory is a tiny read-only space whose cache broadcasts uniform reads to a whole warp. Use it for small, shared values. ✨
Frequently asked questions
Is the “Constant Memory and Its Cache” lesson free?
Yes — the full text of “Constant Memory and Its Cache” 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 “Constant Memory and Its Cache”?
Broadcasting read-only values cheaply. 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 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Constant Memory and Its Cache” 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