Shared Memory Banks: Eliminate a 32-Way Conflict

Each lane reads a different shared-memory word, but the starter spaces those words 32 apart. For 32-bit values, bank = word index mod 32, so all 32 addresses land in one bank and the warp request is serialized. Remap each warp to 32 adjacent words so lane 0 uses bank 0, lane 1 uses bank 1, and so on.

Languages: CUDA C++, Rust

Lesson

Shared memory is divided into 32 banks, and each bank can normally serve one 32-bit word per warp request. Consecutive words rotate through consecutive banks. A stride of 32 words wraps the bank index back to the same bank, so lane 0 reading word 0, lane 1 reading word 32, and lane 2 reading word 64 creates distinct addresses in bank 0. The hardware services those addresses as extra wavefronts instead of one parallel request. GPUFlight reports that cost as average N-way serialization: 1.00x is conflict-free and a deliberate stride-32 pattern approaches 32x.

A same-address read broadcast differs from different addresses mapping to the same bank. Preserve partial-warp checks, then apply the mapping to the transpose padding sequel.

What you will learn

Calculate a 32-bit shared-memory word's bank with wordIndex mod 32.
Distinguish a same-address broadcast from different addresses colliding in one bank.
Lay out each warp's values contiguously so adjacent lanes use adjacent banks.
Interpret GPUFlight's shared-memory N-way metric, where 1.00x is conflict-free.

Starter code

    extern __shared__ volatile float tile[];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;
    int lane = tid & 31;
    int warp = tid >> 5;

    // Every lane uses a distinct word, but a 32-word stride maps all of them
    // to the same bank: (lane * 32) mod 32 = 0.
    int sharedIndex = warp * 1024 + lane * 32;
    tile[sharedIndex] = static_cast<float>(lane + 1);
    __syncthreads();

    float sum = 0.0f;
    for (int repeat = 0; repeat < 256; ++repeat) {
        sum += tile[sharedIndex];
    }
    if (i < N) output[i] = sum;