Bank conflicts in shared memory: Worked examples — Hardware deep dive

This page is the drill floor for Bank conflicts in shared memory. We take the one master tool — the bank formula and its conflict degree — and throw every kind of access pattern at it until nothing can surprise you.

Before we compute anything, let us re-earn the two tools we will lean on the whole way.

The offset $b$ never changes the conflict degree — it only slides every thread by the same amount, shifting which banks are used, not how many threads share a bank. Keep that in your pocket.

The scenario matrix

Every access pattern this topic can throw at you falls into one of these cells. The worked examples that follow are each tagged with the cell they cover.

Cell	What makes it distinct	Covered by
A. Stride-1 (degenerate, best case)	$s = 1$ , gcd $= 1$	Ex 1
B. Odd stride ≠ 1	$s$ odd, coprime to 32, still conflict-free	Ex 2
C. Even stride, partial conflict	$s$ shares a factor $2^{k}$ with 32	Ex 3
D. Worst case (fully serialized)	$s$ a multiple of 32, gcd $= 32$	Ex 4
E. Nonzero offset $b$ (sign/shift check)	offset present — does it change the answer?	Ex 5
F. Column walk in a 2-D array + padding fix	derived stride from array shape	Ex 6 (figure)
G. AoS struct-size trap	stride implied by `sizeof(struct)`	Ex 7
H. Sub-word / broadcast (zero effective stride)	all threads → one address (special, NOT a conflict)	Ex 8
I. Real-world word problem	histogram / reduction access	Ex 9
J. Exam twist (multi-word / double, ≥32 threads)	8-byte type, reason in 4-byte banks	Ex 10

We will also hit the limiting behaviour: as $s$ climbs through powers of two, watch the conflict degree double at each step (Ex 3 → Ex 4), and as $s$ becomes odd again the conflicts vanish (Ex 2). The visual for that is below.

Figure — Bank conflicts in shared memory

Worked examples

Worked example Ex 1 — Cell A · stride-1, the degenerate best case

float s[128]; float v = s[threadIdx.x]; How many cycles for one warp?

Forecast: guess the cycle count before reading on. (Hint: this is the case banks were built for.)

Identify the line. Word index of thread $i$ is $i \cdot 1 + 0$ , so $s = 1, b = 0$ . Why this step? Everything downstream needs $s$ and $b$ ; read them straight off the index expression.
Conflict degree $= g cd (1, 32) = 1$ . Why this step? gcd is the only thing that decides serialization — degree 1 means each bank serves at most one thread.
Banks used: $bank (i) = i mod 32 = i$ for $i = 0..31$ → all 32 banks, each once. Why this step? Confirms the degree geometrically: 32 distinct banks, no pile-up.

Answer: 1 cycle.

Verify: 32 distinct bank values ${0, \dots, 31}$ , max threads on any bank $= 1$ . Cycles $=$ max degree $= 1$ . ✓

Worked example Ex 6 — Cell F · column walk + the padding fix (with figure)

__shared__ float M[32][32]; row-major, so M[r][c] is word $r \cdot 32 + c$ . Warp does M[tid][col] (fixed column, walking rows). Then repeat with the padded array M[32][33].

Forecast: guess the cycle count before padding, then after.

Unpadded stride. Thread $i$ reads word $i \cdot 32 + col$ → $s = 32, b = col$ . Why this step? Moving down one row jumps 32 words — the row length is the stride.
Degree $= g cd (32, 32) = 32$ → 32-way conflict, 32 cycles. All rows share bank $(col mod 32)$ .
Pad the inner dimension to 33. Now M[r][c] is word $r \cdot 33 + c$ , so thread $i$ reads $i \cdot 33 + col$ → $s = 33$ . Why this step? We deliberately break the "multiple of 32" stride.
New degree $= g cd (33, 32) = 1$ (33 is odd) → conflict-free, 1 cycle. Since $33 \equiv 1 (mod 32)$ , bank $(i) = (i + col) mod 32$ — consecutive rows shift by exactly one bank. Why this step? One wasted column per row turns a 32× slowdown into full speed; this is the standard transpose trick.

Answer: unpadded 32 cycles → padded 1 cycle.

Verify: unpadded banks all equal $col mod 32$ (degree 32); padded banks ${(i + col) mod 32}$ are all distinct (degree 1). ✓

This is exactly the pattern behind matrix transpose optimization, and it lives right next to memory coalescing in your kernel-tuning toolkit.

Worked example Ex 7 — Cell G · Array-of-Structures size trap

struct P4 { float a, b, c, d; };   // 16 bytes = 4 words
__shared__ P4 arr[32];
float x = arr[threadIdx.x].a;      // read the .a field

Conflict?

Forecast: you don't pick this stride — the struct size does. Guess the degree.

Struct size in words: $16/4 = 4$ . Adjacent .a fields are 4 words apart → $s = 4$ , offset $b = 0$ (the field is first). Why this step? In AoS the compiler-chosen sizeof(struct) becomes your stride whether you like it or not.
Degree $= g cd (4, 32) = 4$ → 4-way conflict, 4 cycles. Why this step? A "nice" power-of-two struct silently shares factor 4 with 32.
Contrast: a 3-word struct (float a,b,c;) gives $g cd (3, 32) = 1$ → conflict-free. Same code shape, opposite result, purely from struct size.

Answer: 4-way, 4 cycles for the 16-byte struct.

Fix: switch to structure-of-arrays (float a[32], b[32], …) → every field access is stride-1 → guaranteed 1 cycle, independent of struct size.

Verify: bank $(i) = (4 i) mod 32$ takes only 8 distinct values (0,4,…,28), each shared by 4 threads → degree 4. ✓

Worked example Ex 8 — Cell H · broadcast (all threads, one address — NOT a conflict)

float s[128]; int j = 7; float v = s[j]; (all 32 threads read the same word 7)

Forecast: all 32 hit bank 7 — surely 32 cycles like Ex 4?

Effective stride $s = 0$ : every thread's word index is $7$ regardless of $i$ . Why this step? This is the degenerate zero-stride case the gcd formula does not cover — same address, not different addresses.
Recall the definition of a conflict: two threads must want different addresses in the same bank. Here they want the same address. Why this step? The hardware broadcasts a single word to all requesting threads in 1 cycle — no serialization.
So this is 1 cycle, not 32. Zero stride is the opposite of the worst case.

Answer: 1 cycle (broadcast). The lesson: a bank conflict needs distinct addresses; identical-address reads are free.

Verify: all threads request word 7 → one unique address → broadcast → 1 cycle. ✓

Worked example Ex 9 — Cell I · real-world word problem (shared histogram)

A warp bins 32 samples into a shared histogram: atomicAdd(&hist[sample[tid]], 1);. In the worst input, every sample falls in bin 5. In the best input, samples map to bins $0..31$ (one each). Cycle cost of the shared-memory address phase in each case?

Forecast: which input is fast, which is slow, and by how much?

Worst input: all sample[tid] = 5 → all threads address word 5 → this is the same address (broadcast-style), but an atomicAdd must serialize the read-modify-write on that one bank → effectively 32 serialized updates. Why this step? Atomics on the same address serialize even though a plain read would broadcast — the write side forces ordering.
Best input: bins are a permutation of $0..31$ → 32 distinct banks → all atomics proceed in 1 cycle. Why this step? Distinct banks = independent atomics = full parallelism, the histogram ideal.
Design takeaway: heavily-skewed data drives 32× slowdown; this drives the choice of per-thread private sub-histograms then a merge.

Answer: worst-case 32 cycles, best-case 1 cycle for the address phase.

Verify: all-in-bin-5 → 1 bank, 32 serialized ops; permutation → 32 banks, degree 1. ✓

Worked example Ex 10 — Cell J · exam twist, doubles (8-byte type), reason in 4-byte banks

__shared__ double d[64]; double v = d[threadIdx.x]; A double is 8 bytes = 2 words. Does the natural per-thread access conflict? Reason in 4-byte banks only (there is no "8-byte bank").

Forecast: 8 bytes apart — does the stride-2-looking layout cause a 2-way conflict?

Lay out the words. Thread $i$ 's double occupies words $2 i$ (low half) and $2 i + 1$ (high half). Byte address $8 i$ . Why this step? Always convert the fancy type into raw 4-byte word indices — that is the only unit banks understand.
Naive fear: the low halves sit at words $0, 2, 4, \dots$ → stride 2 → $g cd (2, 32) = 2$ , looks like 2-way. Why this step? Shows where the trap intuition comes from.
What hardware actually does: it splits the warp's 64-word request into two phases. Phase 1 serves threads 0–15 (words 0–31, banks 0–31, all distinct); phase 2 serves threads 16–31 (words 32–63, banks 0–31, all distinct). Each phase touches all 32 banks with no repeat. Why this step? The two phases are how modern GPUs make d[tid] on doubles conflict-free despite the 2-word footprint.

Answer: conflict-free — the straightforward d[tid] double access costs the equivalent of no conflict (2 phases, each 1 cycle, no serialization within a phase).

Verify: phase-1 low-half words ${0, 2, \dots, 30} \Rightarrow$ banks ${0, 2, \dots, 30}$ distinct; high halves ${1, 3, \dots, 31}$ distinct; so words $0..31$ hit banks $0..31$ once each, degree 1. ✓

Recall Quick self-test (reveal after guessing)

Stride 12 — conflict degree? ::: $g cd (12, 32) = 4$ → 4-way conflict. Stride 15 — conflict degree? ::: $g cd (15, 32) = 1$ (15 is odd) → conflict-free. Does adding offset $+ 9$ change any conflict degree? ::: No — offset shifts which banks, never the pile-up count. All 32 threads read the same word — how many cycles? ::: 1 (broadcast), because it's the same address, not different ones. Padding [32][32]→[32][33] fixes a column walk because 33 mod 32 = ? ::: $1$ , so each row shifts by exactly one bank.

Where to go next: conflict-free shared memory feeds directly into occupancy and resource limits (shared memory per block caps how many warps fit), and the column-walk fix is the heart of matrix transpose optimization. Beware combining these lessons with warp divergence, which serializes for a different reason (branches, not banks), and with memory coalescing, which is the global-memory cousin of this same "make accesses line up" idea. The layout rules trace back to shared memory architecture.