Constant Memory: Broadcast Hot Coefficients

Every lane in a warp repeatedly requests the same entry from a small read-only coefficient table, but the starter uses a general global-memory path. Choose the CUDA memory path that matches a small, warp-uniform lookup and preserve the index pattern.

Languages: CUDA C++, Rust

Lesson

Constant memory is useful for small, read-only tables shared by many threads. Its best case is a warp-uniform address, which the constant cache can broadcast. Divergent constant addresses serialize, so the data size and access pattern must both fit the mechanism.

The baseline deliberately uses volatile global loads to expose the mechanism. Compare a normal cached-global implementation before generalizing this result to production.

What you will learn

Recognize a small read-only coefficient table that fits constant memory.
Choose a read-only storage path that can broadcast a warp-uniform address.
Preserve one coefficient address per warp to use the broadcast path.
Explain why divergent constant-memory indexing can be slower.

Starter code

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= N) return;

    int warp = tid >> 5;
    const volatile float* globalTable = table;
    float sum = 0.0f;
    for (int i = 0; i < kConstantIterations; ++i) {
        sum += globalTable[(warp + i) & (kConstantEntries - 1)];
    }
    output[tid] = sum;