Shared Memory Bank Conflicts

Definition

To sustain its huge bandwidth, shared memory is split into 32 banks (one per warp lane), 4 bytes wide, interleaved by word. The hardware services one access per bank per cycle, so a warp’s 32 accesses complete in one transaction iff they hit 32 distinct banks (or all hit the same address = broadcast). If $k$ threads hit different addresses in the same bank, they serialize into $k$ transactions — a $k$-way bank conflict costing $k\times$ the cycles.

Key math

$\text{bank}(\text{addr})=\left(\frac{\text{addr}}{4}\right)\mathrm{mod}32$

The classic offender: column access of a float tile[32][32]. Element tile[i][c] sits at word offset $32i+c$, so its bank is $(32i+c)\mathrm{mod}32=c\mathrm{mod}32$ — every row in a column maps to the same bank → a 32-way conflict, a 32× slowdown. Any stride sharing a common factor with 32 serializes (diagonal/transpose patterns, attention score tiles).

The fix: padding. Declare one element wider — float tile[32][33]. Now offset is $33i+c$, bank $(33i+c)\mathrm{mod}32=(i+c)\mathrm{mod}32$; as $i$ runs $0\dots 31$ this takes 32 distinct values → conflict-free. Cost: one wasted column ($32\times 4=128$ B/tile). A stride-2 access (s[2t]) hits 16 banks twice → 2-way conflict; making the stride coprime with 32 (stride 1) is the clean fix.

Why it matters

Bank conflicts are the on-chip analog of uncoalesced global access — same “serialize when addresses collide” principle, one level up. Every high-performance GEMM and matrix-transpose kernel pads its SMEM tiles to stay conflict-free; a 32-way conflict on an inner-loop SMEM read can dominate runtime. Conflict-free SMEM layout is the third leg (with coalescing and tiling) of taking matmul from a few % of peak to 80–90%+, and it matters wherever LLM kernels stage tiles through SMEM (FlashAttention score tiles, transpose in attention).

Taught in

• 02-gpu-memory-hierarchy — §5 banks, bank conflicts, and the padding fix.

See also

• gpu-memory-hierarchy — the SMEM tier
• memory-coalescing — the global-memory analog
• warp-simt — the 32 lanes hitting banks
• gemm — padded conflict-free tiles
• cuda-execution-model — __shared__ and __syncthreads()
• gpu-systems-for-llms