Shared Memory and Bank Conflicts

Coalescing taught you to spend global-memory bandwidth wisely. This lesson is about not spending it at all. Every SM contains a block of SRAM — on the order of 100 KB — that the hardware will hand over to you: shared memory, a programmer-managed scratchpad with roughly 20× the bandwidth and a tiny fraction of the latency of global memory. No cache heuristics, no eviction surprises: you explicitly copy data in, every thread of the block can see it, and you explicitly decide when it is stale. It is the only memory in the machine whose behaviour you fully control — and the deal has exactly two clauses in the fine print: a barrier you must not misuse, and thirty-two banks you must not collide in.

The canonical use: tiled matrix multiply

Why does a scratchpad matter so much? Because of reuse. Recall from the roofline that a kernel's fate is set by its arithmetic intensity — FLOPs per byte of DRAM traffic. Naive matrix multiply is the textbook bandwidth victim: computing one output element reads a whole row of A and column of B, and neighbouring outputs re-read almost exactly the same data from global memory, again and again.

The fix is the most famous kernel structure in GPU computing. Carve the matrices into T \times T tiles, and let each block stage one tile of A and one of B in shared memory:

tiled multiply, C = A·B (T = 32; the block is T×T threads) for m in 0 .. K/T - 1: # march tile-pairs along the K dimension Asub[ty][tx] = A[row][m·T + tx] # each thread loads ONE element of each tile Bsub[ty][tx] = B[m·T + ty][col] # (coalesced!) into shared memory __syncthreads() # barrier: both tiles fully loaded for k in 0 .. T - 1: acc += Asub[ty][k] * Bsub[k][tx] # 2·T reads per thread — ALL from shared __syncthreads() # barrier: everyone done before reloading C[row][col] = acc

Count the traffic. Each element loaded into Asub is subsequently read by all T threads of its row; each element of Bsub by all T threads of its column. One global load now feeds T uses — with T = 32, global traffic drops 32× and arithmetic intensity rises 32×, which on the roofline is the difference between crawling along the bandwidth slope and reaching the compute plateau. This is the pattern to internalise: shared memory is an arithmetic-intensity multiplier — stage once, reuse many times, synchronise at the seams.

Thirty-two banks

Fast for thirty-two lanes at once — how? Shared memory is not one SRAM but 32 independent banks, interleaved by 4-byte word: word address w lives in bank w \bmod 32. Words 0–31 sit in banks 0–31, word 32 wraps back to bank 0, and so on. Each bank serves one word per cycle — so if the warp's 32 lanes hit 32 different banks, all 32 requests complete in a single cycle. That is the parallel jackpot the layout is designed for.

The classic self-inflicted wound: a 32 \times 32 tile stored row-major, accessed by column. Lane t reads tile[t][c] — word address 32t + c — and (32t + c) \bmod 32 = c for every lane: all 32 lanes pile into the same bank, wanting 32 different words. A 32-way conflict; the "one-cycle" scratchpad quietly runs 32× slower, precisely in the transpose-flavoured kernels where column access is natural.

The cure costs one wasted column: pad the row to 33 words. With tile[t][c] at address 33t + c, lane t maps to bank (33t + c) \bmod 32 = (t + c) \bmod 32 — thirty-two consecutive banks, conflict-free, for 3% more SRAM. One of the best exchange rates in performance engineering.

Seeing the banks

Here is the whole story on a mini shared memory with 8 banks (the real one is this times four). Each cell shows the bank its word maps to; watch which cells light up for each access pattern:

The padding trick is visible at a glance: adding one dead word per row shears the bank pattern diagonally, so columns — which used to stack in one bank — now sweep across all of them.

The conflict counter, executable

A bank conflict is just bucket-counting: map each lane's word address to its bank, and the cost in cycles is the largest number of distinct words any one bank must serve. Four patterns, graded:

// Bank-conflict counter: cost (cycles) = max distinct WORDS in any one bank. const BANKS = 32, WIDTH = 32; const lanes = [...Array(32).keys()]; function cost(addrs: number[]): number { const perBank = new Map<number, Set<number>>(); for (const a of addrs) { const bank = a % BANKS; if (!perBank.has(bank)) perBank.set(bank, new Set()); perBank.get(bank)!.add(a); // distinct words only — same word = broadcast } let worst = 1; for (const words of perBank.values()) worst = Math.max(worst, words.size); return worst; } const patterns: [string, (t: number) => number][] = [ ["row access tile[0][t] ", (t) => 0 * WIDTH + t], ["column access tile[t][0], width 32 ", (t) => t * WIDTH + 0], ["padded column tile[t][0], width 33 ", (t) => t * 33 + 0], ["stride 2 tile2[2*t] ", (t) => 2 * t], ["broadcast tile[0][0] ", () => 0], ]; for (const [name, f] of patterns) { const c = cost(lanes.map(f)); console.log(`${name} ${String(c).padStart(2)}x (${c === 1 ? "parallel — 1 cycle" : c + "-way conflict — " + c + " cycles"})`); }

Try stride 4 and stride 8 (degree = \gcd(s,32) — so 4 and 8), then stride 3 or 5: odd strides land on \gcd = 1, conflict-free. And check the broadcast line: 32 lanes, one word, one cycle — the hardware fans the value out for free.

The other occupancy currency

Shared memory is fast, but it is a fixed pot: the ~100 KB per SM is divided among all the blocks resident on that SM. Ask for 48 KB per block and at most two blocks fit; their warps are all the scheduler has to hide latency with. Shared memory thus joins registers as a currency that buys per-thread capability at the price of thread count — a tension this module resolves properly in the occupancy capstone.

Everywhere — banking is how hardware fakes a multi-ported memory on the cheap. A true 32-port SRAM (32 simultaneous arbitrary reads) would be enormous and slow; 32 single-ported banks with an interleaved address map deliver the same bandwidth as long as the addresses spread out. DRAM itself is banked, vector supercomputers interleaved their memory across banks in the 1970s (Cray programmers dreaded stride-8 arrays for exactly the gcd reason you just learned), and register files are banked too. The GPU's shared memory just puts the mechanism unusually close to the programmer: you get the multi-port illusion, and you get the address arithmetic that can break it.

__syncthreads() is a block-wide barrier: every warp of the block must arrive before anyone proceeds. Put it inside a divergent if — where some threads of the block take the branch and others don't — and the threads that skipped it never arrive; the ones waiting wait forever. Deadlock, or on some hardware, silent corruption as the barrier miscounts. The rule is absolute: a barrier must sit at a point every thread of the block reaches — hoist it out of the conditional, or restructure so the condition is uniform across the block. (This is divergence's nastiest interaction with shared memory: the tile pattern is load, sync, use, sync — and both syncs must be in straight-line code.)