Shared Memory and Bank Conflicts

the short version

Shared memory is fast because it is physically split into independent banks that can be read simultaneously. Consecutive 4-byte words live in consecutive banks, cycling around. If every lane in a group hits a different bank, all of the accesses complete together.

If two lanes hit different addresses in the same bank, the hardware has no choice but to serialize them. An N-way conflict costs N times as long, and a pathological access pattern can turn the fastest memory on the chip into something slower than a cache hit.

how banking works

Think of shared memory as a set of parallel columns. Word 0 is in bank 0, word 1 in bank 1, and so on, wrapping back around after the last bank. A lane reading word 33 and a lane reading word 1 are hitting the same bank, because the bank is determined by the address modulo the bank count.

Two lanes reading the same address is not a conflict. The hardware broadcasts a single value to every lane that wants it, which is why reading a shared coefficient across a whole group is free. The conflict is specifically different addresses landing in the same bank.

the classic case

A tile of floats declared with a power-of-two row length is the textbook trigger, because a column walk then hits the same bank every time.


__shared__ float tile[32][32];

// Each lane reads a different row, same column.
// Row stride is 32 words, so every lane lands in the same bank.
float v = tile[threadIdx.x][0];     // 32-way conflict, fully serialized

Every row starts exactly one full bank cycle apart, so column c of every row sits in the same bank. Thirty-two lanes, one bank, thirty-two serialized accesses.

The fix is one character:


__shared__ float tile[32][33];   // pad the stride by one

float v = tile[threadIdx.x][0];      // conflict-free

With a 33-word stride, consecutive rows are offset by one bank, so a column walk now sweeps across all banks instead of hammering one. You waste a column of storage and recover the full width of the memory. This padding trick shows up in essentially every hand-written transpose and tiled matmul kernel.

the wave size complication

Bank conflicts are resolved per execution group, which means the group size participates. On hardware with 64-lane wavefronts the hardware handles the access in halves, so the conflict analysis is done over a half-wavefront rather than all 64 lanes at once.

The practical consequence is that a padding scheme tuned on 32-lane hardware is not automatically correct on 64-lane hardware, and vice versa. It is another item on the list of things that port silently and wrongly, alongside the shuffle and ballot issues.

when to care

Bank conflicts only matter in kernels that hammer shared memory in their inner loop, which in practice means tiled matrix kernels, transposes, stencils and reductions. A kernel that loads a tile once and reads it in a simple contiguous pattern is very unlikely to have a problem.

The ordering also matters: fix your global memory access pattern first. Coalescing failures cost you HBM bandwidth, which is the scarcest resource on the machine; bank conflicts cost you scratchpad bandwidth, which is plentiful. A kernel with both problems should have the global one fixed first.

see it on your own machine

ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum ./my_app | https://docs.nvidia.com/nsight-compute/ | counts shared memory bank conflicts directly; zero is the target and any large number in a hot kernel is worth the padding experiment |'bc_ncu'

ncu --set full ./my_app | https://docs.nvidia.com/nsight-compute/ | the memory workload section flags shared memory conflicts in context, alongside the global access analysis |'bc_full'

reference

CUDA C++ Best Practices Guide
CUDA Programming Guide: shared memory
AMD HIP documentation