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

CUDA constant memory, measured against global

A post on NVIDIA's own forums, December 2009:

"I tryed to use different memory types for the same problem, so I stored the array in global memory, which was quite fast, then in shared mem, which was a bit faster, and in constant memory, which was really slow. [...] The kernel launch takes like ten times the processing time, when using constant memory, compared to global memory."

https://forums.developer.nvidia.com/t/really-slow-constant-memory-random-access-to-constant-memory/13568 (checked 2026-08-30)

Ten times slower, from the memory every tutorial calls the fast one. The post gives away the cause two paragraphs in: each thread read a different element of the array.

Nothing was broken. Constant memory is not always faster. It has one fast access pattern and one slow one.

That line uses the slow pattern. This page measures both patterns against the same filter in global memory across five sizes, and looks for the size where the question stops mattering.

What the constant cache charges for

__constant__ declares a file-scope array in a 64 KB region of device memory. You 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" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html , checked 2026-08-30). The 64 KB is the same figure on every compute capability from 7.5 to 12.x, from Table 31 of the compute-capability appendix (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-30).

The region is not the interesting part. The cache in front of it is, and what that cache bills you for. The Best Practices Guide is exact about it:

"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. As such, the constant cache is best when threads in the same warp accesses only a few distinct locations. If all threads of a warp access the same location, then constant memory can be as fast as a register access."

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

Read that as arithmetic. A warp whose 32 lanes all read filter[k] asks for one address and needs one access. A warp whose 32 lanes read filter[0] through filter[31] asks for 32 addresses and needs 32 accesses, one after another, and it makes no difference that all 32 floats sit inside the same 128 bytes.

Global memory bills on a different axis. Day 11 measured it: a warp's 32 addresses collapse into the set of 32-byte sectors that covers them, so 32 consecutive floats cost four sectors and one request, however many distinct addresses that is.

Global memory charges you for sectors. Constant memory charges you for distinct addresses. That difference explains the forum post above and this page's results.

Constant memory serves 32 lanes reading one tap in 1 access, but serializes 32 different taps; global memory covers those consecutive taps with 4 sectors.

The filter was already in a cache

The habit you arrive with is that read-only data belongs in constant memory, because that is what constant memory is for. It was better advice in 2009 than it is now, and the reason is that the rest of the memory system caught up.

The filter here is 32 taps, which is 128 bytes. Every global access goes through L2, a device-wide cache measured in megabytes; the program prints your card's l2CacheSize so you do not have to look it up.

The first warp to touch the filter pulls those 128 bytes into L2 and into its SM's L1. Every later warp on that SM hits in L1, and a warp on any other SM hits in L2. Constant memory cannot cut this DRAM traffic further.

There is a third path, and on this card it is not a third path at all. A const pointer parameter marked __restrict__ compiles to a read-only load: "Accesses to __global__ function const pointers marked with __restrict__ are compiled as read-only cache loads, similar to the PTX ld.global.nc or __ldg() low-level load and store functions instructions" (same programming guide page, checked 2026-08-30).

People call that the read-only data cache, or the texture path. On Turing it is L1: "Like Pascal and Volta, Turing combines the functionality of the L1 and texture caches into a unified L1/Texture cache" (https://docs.nvidia.com/cuda/turing-tuning-guide/index.html section 1.4.3.1, checked 2026-08-30).

It is not a cache next to L1. It is L1, told the data will not change.

Constant memory beats DRAM. The useful question is whether it beats L1 holding the same 128 bytes, which leaves a much smaller gap to measure.

Hardware. From compute capability 8.0 you can pin a global region in L2 instead: "Starting with CUDA 11.0, devices of compute capability 8.0 and above have the capability to influence persistence of data in the L2 cache" (best practices guide 10.2.2, checked 2026-08-30). You reserve part of L2 with cudaDeviceSetLimit and mark the region with a cudaAccessPolicyWindow. The T4 this course targets is 7.5, so persistingL2CacheMaxSize leaves nothing to reserve and the calls buy no residency. The program prints that field, so you can see what your own card reports. On a 7.5 card the way to keep something resident is the cache you manage yourself, shared memory, on day 13.

cudaStreamAttrValue attr;
attr.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(d_filter);
attr.accessPolicyWindow.num_bytes = filterBytes;
attr.accessPolicyWindow.hitRatio = 1.0f;
attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

One kernel pair, four rows

Full program in code/day18-constant-memory/constant_memory.cu. It runs a 32-tap 1D convolution four ways at each of five sizes: the filter in constant memory or in global memory, crossed with every lane reading the same tap or every lane reading a different one. The comparison follows three rules.

The lane pattern is a runtime argument, not a second kernel. rotMask is 0 for the broadcast rows and 31 for the per-lane rows, and it arrives as a kernel parameter, so the compiler cannot fold it away and all four launches execute the same instructions. Write two kernels instead and you are comparing two instruction streams as well as two access patterns, and you can no longer say which one moved the number.

Both kernels are checked before either is timed. Each is compared against the same CPU reference at all ten cases, and the program also checks that the per-lane row really is a different access pattern, by counting how many outputs it shares with the broadcast row.

If rotMask ever stopped changing addresses, the two rows would measure the same access twice and the table would look reasonable while proving nothing. Day 11 covers this benchmark bug.

Time with events, and warm up every kernel. Day 9 covers why a host clock around a launch measures the launch, and why warming one kernel does not warm the next.

The constant-memory kernel:

__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;
    }
}

The global-memory kernel is the same arithmetic with one more parameter:

__global__ void conv1dGlobal(const float* __restrict__ in,
                             const float* __restrict__ filter,
                             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 += filter[(rot + k) & kTapMask] * in[i + k];
        }
        out[i] = acc;
    }
}

The same 32 floats go to both places before anything is timed. Note what cudaMemcpyToSymbol takes:

    CUDA_CHECK(cudaMemcpyToSymbol(c_filter, h_filter.data(), filterBytes));
    CUDA_CHECK(cudaMemcpy(d_filter, h_filter.data(), filterBytes,
                          cudaMemcpyHostToDevice));

Note. The per-lane rows are not a convolution anyone wants. Rotating the taps by the lane index computes a different function for every lane, and no signal processing asks for that. It is the access pattern from the forum post at the top of this page, isolated so it can be timed, and it reads the same inputs and writes the same outputs as the broadcast row, which is what makes the two comparable. A benchmark kernel is named for what it measures, not for what it computes.

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. The access-pattern conclusion held, but the largest per-lane penalty moved from 15.08x to 29.98x.

Both transcripts remain in evidence; the original 12.6 table follows.

GPU: Tesla T4 (compute capability 7.5)
L2 4096 KiB, persisting L2 max 0 B, policy window max 0 B
This device sets aside no L2 for persisting accesses, so cudaAccessPolicyWindow has nothing to reserve.
Filter: 32 taps, 128 bytes, of 65536 bytes of constant
256 threads per block, 3 warm-ups, 10 timed runs, mean reported, no copies inside the measurement

        n  lanes      const ms  global ms   ratio  global GB/s
 --------  ---------  --------  ---------   -----  -----------
     4096  broadcast    0.0041     0.0039    1.05          8.5
     4096  per-lane     0.0198     0.0038    5.21          8.7
    32768  broadcast    0.0070     0.0078    0.90         33.7
    32768  per-lane     0.0706     0.0076    9.31         34.6
   262144  broadcast    0.0256     0.0340    0.75         61.8
   262144  per-lane     0.4290     0.0339   12.65         61.8
  2097152  broadcast    0.1769     0.2550    0.69         65.8
  2097152  per-lane     3.3542     0.2557   13.12         65.6
  8388608  broadcast    0.6912     1.0175    0.68         66.0
  8388608  per-lane    11.7892     0.7819   15.08         85.8

Both kernels in a row read the same 32 taps and the same
n + 31 inputs, and write the same n outputs, so neither can win
by doing less work. Only where the filter lives changes.

Constant memory wins when every lane reads the same address, and loses sharply when they do not. At the largest size the CUDA 12.6 run measured 0.68x and 15.08x the global-memory time; CUDA 13.0 measured 0.69x and 29.98x. The direction reproduced, while the divergent-access magnitude did not.

At 8.4 million elements the broadcast case runs at 0.68x the global-memory time. The constant cache serves one address to the whole warp in a single transaction, which is exactly the access pattern a convolution filter has: every thread wants tap zero, then tap one.

The per-lane column shows the bad case. It uses the same data and size, but indexes it so that each lane wants a different tap, and it costs 15.08x the global version. Constant memory has no cache line to spread across lanes; divergent addresses serialise.

A filter is a good fit and a lookup table indexed by thread is a terrible one.

At 4096 elements the ratio is noise-sized: 1.05x under CUDA 12.6 and 0.99x under CUDA 13.0. The reason is not the filter. Both kernels finish in a few microseconds and the original global row reports 8.5 GB/s against 66.0 at the top of the sweep, so that row is measuring launch and tail effects rather than the memory system.

The constant advantage arrives once the kernel is long enough to be about memory at all, and it is already there at 32,768 elements, well inside this card's 4 MiB of L2. Reaching for __constant__ to speed up a kernel that runs in microseconds is optimising the wrong thing.

This device reserves no L2 for persisting accesses, reporting 0 bytes for both the persisting maximum and the policy window. cudaAccessPolicyWindow is a compute capability 8.0 feature and there is nothing here to reserve, so the L2 residency control mentioned above is a forward pointer and not something you can try on a T4.

Run it yourself

Colab's free T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o constant_memory constant_memory.cu

There is no Compiler Explorer embed on this page. The program allocates 64 MiB of device buffers, makes 260 launches and runs a CPU reference over 8 Mi elements at 32 taps each, twice, which does not fit Compiler Explorer's 20 second run cap.

Cut kSizes to its first two entries and it does. If you have no GPU, read /setup/learn-cuda-without-a-gpu.

Exercise

Run the sweep, then answer in two sentences: why do the broadcast columns stay together as n grows, and why do the per-lane columns not?

Time: 25 to 40 minutes. Submit: your ratio column and the two sentences.

Check: three gates, each a real branch returning EXIT_FAILURE rather than an assert, because CI builds Release and NDEBUG deletes assert. Both kernels are compared against the same CPU reference at all ten cases, and the first mismatch prints the kernel, the index, the value it got and the value it wanted.

The broadcast reference is compared against the per-lane one and the run fails if more than half the outputs agree, which is how you learn that rotMask stopped doing anything. The row count is also checked against the ten in the table.

The ratio itself is not gated, because it is the claim the run exists to test; the T4 column above is the reference your own column is read against.

Hint 1

Count addresses, not bytes. At one instant, one warp, one iteration of the tap loop: how many different addresses does each variant ask its memory path for, and what does that path charge per address?

Hint 2

The filter is 128 bytes. Read your card's l2CacheSize off the second line of the output and work out how many copies of the filter fit in it. Then ask what work constant memory is saving you in the broadcast case.

Solution

The broadcast columns stay close because almost nothing separates them. Both kernels read the same 128 bytes, and after the first warp of the first block those bytes are cached on both paths, in the constant cache on one and in L1 on the other.

The ratio still settles at 0.68 rather than at 1.00. The constant path serves the whole warp from one fetch instead of issuing a load against a cache line.

The gap is about one third, not an order of magnitude. It reaches that size at about two million elements: 1.05, then 0.90, 0.75, 0.69, 0.68.

The per-lane columns separate because the paths bill on different axes. The constant path charges per distinct address, so 32 lanes wanting 32 taps cost 32 serialized accesses. The global path charges per sector, so the same 32 taps cost four sectors and one request.

Predict the direction, not the size: the gap should not close anywhere in the sweep, because nothing about a larger array gives the constant path fewer distinct addresses to serialize. The measured gap widens from 5.21 to 15.08.

The rule to keep: each memory space has its own access cost. Before moving anything into __constant__ or __shared__, ask how that space handles your access pattern.

Pitfalls

You put a lookup table in __constant__ and index it with threadIdx.x. This is the pattern from the post at the top of the page, and it is the worst case of the one thing constant memory is bad at. Leave the table in global memory and mark the pointer const __restrict__ so it takes the read-only path, or stage it in shared memory if every block reads all of it (day 13).

You assume constant memory is on chip. It is not. Table 1 of the Best Practices Guide, "Salient Features of Device Memory", puts constant memory's location off chip with its Cached column at Yes.

What sits on the SM is the cache, not the 64 KB, and a constant read that misses costs a read from device memory like any other.

You pass the address of the symbol. cudaMemcpyToSymbol(&c_filter, ...) compiles, because the parameter is a const void* either way, and then fails at runtime with invalid device symbol. The modern API takes the variable itself.

See /errors/invalid-device-symbol.

You split the __constant__ across translation units. A __constant__ declared in one .cu file is invisible to another unless you build with -rdc=true, and the symptom is the same invalid device symbol from a cudaMemcpyToSymbol that looks correct. That is separate compilation, and day 67 is the lesson.

You benchmark the two with different launch configurations. A constant kernel at 1024 threads a block against a global kernel at 32 is two launches doing different amounts of work, and whichever wins, the clock is answering a question about occupancy rather than about where the filter lives.

That is why rotMask here is a runtime argument and both rows share one launch shape, and day 11 covers this. For the sector and request counts behind these rows rather than the times, Nsight Compute is day 42.

Go deeper

Next

Day 19 leaves the on-chip caches and asks what managed memory costs when the pages themselves have to move. Day 20 is the module's capstone and puts this filter to work in two dimensions, with the halo, the tile and the constant filter together. Both sit in module 2, and the GPU kernel engineer hub maps where the rest of it goes.