Constant memory (CUDA)

Constant memory (CUDA)

Constant memory is a small read-only region of device memory, declared with __constant__ and served through a dedicated cache. It is optimised for one specific pattern: every thread in a warp reading the same address at the same time.

__constant__ float coeff[16];
 
// host side, before the launch
cudaMemcpyToSymbol(coeff, h_coeff, 16 * sizeof(float));
 
__global__ void filter(const float *in, float *out) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;
    for (int k = 0; k < 16; k++)
        sum += coeff[k] * in[i + k];    // every thread reads coeff[k] together
    out[i] = sum;
}

When a whole warp reads one address, the constant cache broadcasts the value in a single operation, as cheaply as a register read. When threads in a warp read different addresses, the accesses serialise instead, one per distinct address, which makes constant memory a poor choice for anything indexed by thread ID.

The total budget is 64 KB per device, and the contents are set from the host with cudaMemcpyToSymbol before the launch. Kernels cannot write to it.

Typical uses are filter coefficients, lookup tables, transformation matrices, and configuration structs. In practice, a plain const __restrict__ pointer into global memory often performs just as well on recent hardware, since the compiler routes it through the read-only data cache. Constant memory is still worth reaching for when the broadcast pattern is exact and the data is genuinely small.

Table of Contents

Constant memory (CUDA)