100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryMemory
CC 7.5

What is constant memory in CUDA?

A 64 KB read-only region of device memory with its own cache, fastest exactly when every lane of a warp reads the same address.

You declare it at file scope with __constant__, fill it from the host with cudaMemcpyToSymbol, and the kernel can only read it: __constant__ variables "are read-only in device code and can only be modified from the host using the CUDA Runtime API". The 64 KB is the same figure on every compute capability from 7.5 through 12.x. The region itself sits off chip, in the same DRAM everything else lives in, so the thing that makes it interesting is the cache in front of it and the odd way that cache bills.

Here is the bill, from the best practices guide: "Accesses to different addresses by threads within a warp are serialized, thus the cost scales linearly with the number of unique addresses read by all threads within a warp ... If all threads of a warp access the same location, then constant memory can be as fast as a register access." So a warp whose 32 lanes all read filter[k] asks for one address and pays once. A warp whose lanes read filter[0] through filter[31] asks for 32 and pays 32 times, and it makes no difference that those floats sit inside one 128-byte line. Global memory charges on the other axis: 32 addresses collapse into the sectors that cover them, four for that same line, however many distinct addresses there were. Constant memory charges for addresses, global memory charges for sectors, and every surprise about this space comes out of that one difference.

The other thing that has changed since the advice was written is that the rest of the memory system caught up. A 32-tap filter is 128 bytes, which the first warp to read it pulls into L2 and L1 on the way, so the global version is not going back to DRAM either. On Turing there is not even a separate read-only cache to appeal to: the guide says a const __restrict__ pointer compiles to a read-only load, and the Turing tuning guide says that path is L1. The honest question is whether the constant cache beats L1 holding the same 128 bytes, which is a much smaller gap than the folklore promises, and the answer is measured below.

Measured

Day 18 ran one 32-tap convolution with the filter in __constant__ and again in global memory, at five sizes, with a runtime mask deciding whether the lanes wanted one tap or 32. The filter used 128 bytes of the 65,536 available. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75.

Elements Every lane reads the same tap Each lane reads a different tap
4,096 1.05 5.21
32,768 0.90 9.31
262,144 0.75 12.65
2,097,152 0.69 13.12
8,388,608 0.68 15.08

Each cell is the constant-memory time divided by the global-memory time, so under 1.00 means constant won. Both kernels in a row read the same taps and the same inputs and write the same outputs, and the launch shape is identical, so neither can win by doing less. Two readings matter. The broadcast column settles near a third off rather than an order of magnitude, because L1 was already holding the filter. And the per-lane column does not converge, it widens the whole way to 15.08, because serialization scales with the element count while sector counts do not. The 4,096 row is the one to distrust: both kernels finish in microseconds there and it is measuring launch effects, not memory.

Code

From code/day18-constant-memory/constant_memory.cu. rotMask is a kernel argument so the compiler cannot fold it away, which is what lets one instruction stream produce both columns above.

__global__ void conv1dConstant(const float* __restrict__ in,
                               float* __restrict__ out, size_t n, int rotMask) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        const int rot = static_cast<int>(threadIdx.x) & rotMask;
        float acc = 0.0f;
        for (int k = 0; k < kFilterTaps; ++k) {
            acc += c_filter[(rot + k) & kTapMask] * in[i + k];
        }
        out[i] = acc;
    }
}

With rotMask at 0 every lane reads tap k at step k, which is one address. With rotMask at 31 the lanes fan out across all 32 taps, which is the access pattern behind a 2009 forum post complaining that constant memory ran ten times slower than global.

Diagram

memory-hierarchy-svg, with the 64 KB constant region drawn in the off-chip DRAM band and the constant cache drawn on the SM beside L1, and a broadcast fan from one cached address to 32 lanes.

Alt text: "Constant memory is 64 KB of off-chip DRAM with a small on-SM cache that serves one address to all thirty-two lanes at once. When the lanes want thirty-two different addresses, that cache serves them one after another."

Related terms

Where you meet this

Sources

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.