# 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 [[warps-cuda|warp]] reading the same address at the same time.

```c
__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-cuda|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.