Bank conflicts in shared memory: Visual walkthrough

This is the visual walkthrough for the parent topic. We build the whole idea of bank conflicts from absolute zero, one picture at a time, until a single clean law — how many threads jam onto the busiest drawer — falls out on its own, and until you see why padding a matrix to width 33 magically fixes everything.

You need almost nothing to start. If a word like "thread" or "shared memory" is new, peek at 6.2.07-Shared-memory-architecture first, but I will re-explain everything you actually need below.

Step 1 — What is a bank, really?

WHAT. Shared memory is one block of fast storage sitting right next to the GPU cores. But it is not one single drawer — it is physically sliced into 32 separate drawers called banks. Each bank can hand out exactly one 4-byte word per clock cycle.

WHY this slicing? Because a group of 32 threads (a warp) all want data at the same instant. One drawer serving one item per tick would force 32 threads to wait in a line 32 long. Thirty-two drawers, each serving one item per tick, means all 32 threads can be served in a single tick — if they each reach into a different drawer.

PICTURE. Look at the figure. The horizontal strip is shared memory. It is chopped into 32 vertical slots, numbered 0 to 31. That number is the bank index. The accent-red slot is bank 0, where our story keeps returning.

Figure — Bank conflicts in shared memory

Step 2 — Which drawer does a word live in?

WHAT. We need a rule that tells us: given a word, which of the 32 banks holds it? The rule is round-robin — deal the words out like playing cards to 32 players sitting in a circle.

WHY round-robin? Because consecutive threads usually want consecutive words. If consecutive words sit in consecutive banks, then a normal shared[tid] access naturally spreads across all 32 drawers — the best possible case. Dealing cards round the table is exactly what achieves that.

PICTURE. Word 0 goes to bank 0, word 1 to bank 1, ..., word 31 to bank 31 — and then word 32 wraps back to bank 0 (accent-red). The wrap-around is the entire source of trouble later.

Step 3 — The happy case: everyone in a different drawer

WHAT. Let each thread $i$ (this is tid, running $i = 0, 1, \dots, 31$ ) read shared[i]. Thread $i$ touches word $w = i$ , so its bank is $i mod 32 = i$ .

WHY show this first? It is the reference point — the "1 cycle" ideal. Every later pattern is judged against it: how many extra cycles did we lose versus this?

PICTURE. Thirty-two arrows fan out, one per thread, each landing on its own bank. No two arrows share a drawer. All served in one tick.

Worked example Stride-1 access

int tid = threadIdx.x;       // this thread's number, 0..31
float value = shared[tid];   // each thread reads its own word

Thread $i \to$ word $i \to$ bank $i$ . All 32 banks distinct. 1 cycle — perfect.

Here "stride" means the gap in words between what thread $i$ and thread $i + 1$ read. Stride-1 = they read neighbours.

Step 4 — Break it on purpose: stride-2

WHAT. Now let thread $i$ read shared[2*i] (i.e. shared[tid*2]). The stride is $s = 2$ . Thread $i$ touches word $w = 2 i$ , so its bank is $(2 i) mod 32$ .

WHY do this? To catch the exact moment the wrap-around bites. We march $i$ upward and watch which thread first lands back on a drawer someone already claimed.

PICTURE. Threads 0..15 fan across the even banks $0, 2, 4, \dots, 30$ . Then thread 16 reads word 32, and $32 mod 32 = 0$ — it slams into bank 0 (accent-red), right on top of thread 0. Thread 17 collides with thread 1, and so on: 16 pairs, each pair fighting over one drawer.

(This is the first time we define "conflict degree" — every earlier line deliberately avoided the term until we could see it happen in a picture.)

Step 5 — The general law: $g cd (s, 32)$

WHAT. We now find one formula covering every stride $s$ . Two threads $i$ and $j$ collide when their banks are equal: $(i \cdot s) mod 32 = (j \cdot s) mod 32 ⟺ (i - j) \cdot s \equiv 0 (mod 32) .$ The right-hand side reads " $(i - j) \cdot s$ is a whole multiple of 32."

WHY $g cd$ — and why is $32/ g cd (s, 32)$ the smallest gap? We are asking: what is the smallest positive gap $k = i - j$ between two colliding threads? So we need the smallest positive $k$ making $k \cdot s$ a multiple of 32.

Here is the reason it comes out to $32/ g cd (s, 32)$ , spelled out. Let $g = g cd (s, 32)$ — the biggest number dividing both $s$ and 32. Split the pieces: $s = g \cdot s^{'}$ and $32 = g \cdot n$ , where now $s^{'}$ and $n$ share no common factor (we divided the shared part out, that's what $g cd$ does). The condition " $k \cdot s$ is a multiple of 32" becomes " $k \cdot g \cdot s^{'}$ is a multiple of $g \cdot n$ ," i.e. " $k \cdot s^{'}$ is a multiple of $n$ ." Since $s^{'}$ and $n$ have no common factor, $s^{'}$ contributes nothing toward the factor $n$ — so $n$ must come entirely from $k$ . The smallest such $k$ is therefore $k = n = 32/ g$ . That is exactly $k = 32/ g cd (s, 32)$ .

From "smallest gap" to "conflict degree" — the missing link. We just learned that colliding threads repeat every $k = 32/ g cd (s, 32)$ steps: threads $0, k, 2 k, 3 k, \dots$ all land on the same bank as thread 0, threads $1, 1 + k, 1 + 2 k, \dots$ all share another bank, and so on. So the 32 threads split into groups, each group sitting on one bank, and every group is the same size. How big is each group? The warp has 32 threads and they repeat every $k$ , so each bank collects $\frac{32}{k} = \frac{32}{32/ g c d ( s , 32 )} = g cd (s, 32)$ threads. That count — threads sharing the busiest (in fact every) occupied bank — is the conflict degree from Step 4. So the smallest gap $k$ and the conflict degree are two sides of one coin: $k$ tells you how far apart colliding threads are; $32/ k = g cd (s, 32)$ tells you how many of them pile up per bank.

$g cd (s, 32)$ means the greatest common divisor — the biggest number dividing both $s$ and 32. We use it because the collision condition is a divisibility question, and $g cd$ is the tool built precisely for divisibility.

PICTURE. A number line of strides. Under each stride sits its collision count. The red markers are the dangerous strides (share a factor of 2 with 32); the black ones are safe (coprime to 32). Notice: odd strides are always safe — an odd number shares no factor of 2 with $32 = 2^{5}$ , so $g cd = 1$ .

Step 6 — The special case the law gets wrong: stride 0 (broadcast)

WHAT. What if every thread reads the same word — stride $s = 0$ ? Plug it in naively: $g cd (0, 32) = 32$ , so the formula screams "32-way conflict, 32 cycles!" But that is wrong on real hardware. When all threads request the same address, the bank reads it once and broadcasts the value to every thread — 1 cycle, no conflict.

WHY the law misses this. Re-read the definition in Step 4: a conflict needs different words in the same bank. Stride 0 means every thread wants the same word, not different ones. The whole "one item per tick" bottleneck exists because a drawer can only pull out one distinct item — but here there is only one item to pull, and copying it out to 32 listeners is free. So the "different addresses" clause is the fine print that saves us.

PICTURE. All 32 arrows point at one word in bank 0 (accent-red), and a single fan of dashed lines carries that one value back out to all 32 threads — one read, one broadcast.

Step 7 — The degenerate extreme: stride-32

WHAT. Push the stride to $s = 32$ . Thread $i$ reads word $32 i$ , bank $(32 i) mod 32 = 0$ — for every $i$ . Each thread wants a different word ( $0, 32, 64, \dots$ ), but they all live in bank 0. That "different words" part is what makes this a genuine conflict (unlike Step 6's broadcast). All 32 pile onto bank 0.

WHY include this? It is the worst true conflict and a crucial edge: $g cd (32, 32) = 32$ , so 32 cycles. Shared memory becomes slower than doing nothing clever at all. Every thread waits in a 32-long line at one drawer.

PICTURE. All 32 arrows converge onto the single accent-red drawer (bank 0), each fetching a different word from it — a stack of 32, served one per tick.

Worked example Stride-32 = 32-way conflict

int tid = threadIdx.x;
float value = shared[tid * 32];   // 32 DIFFERENT words, all in bank 0

Every thread → bank 0, different word each → 32 cycles. This is the pattern that secretly appears when you walk down a column of a 32-wide shared matrix.

Step 8 — Where stride-32 hides: matrix columns

WHAT. Store a matrix row-major: __shared__ float m[32][32]. Element $[r] [c]$ sits at word $r \cdot 32 + c$ .

Row read m[row][tid]: thread $i$ → word $row \cdot 32 + i$ . Stride between threads = 1. Bank $= (row \cdot 32 + i) mod 32 = i$ (since $32 \equiv 0$ ). All distinct → free.
Column read m[tid][col]: thread $i$ → word $i \cdot 32 + col$ . Stride = 32. Bank $= col$ for every thread, and each wants a different word → 32-way conflict.

WHY care? Reading a column is exactly what matrix transpose does (see 6.3.01-Matrix-transpose-optimization). Innocent-looking m[tid][col] is secretly the stride-32 disaster from Step 7.

PICTURE. The matrix grid. A row highlighted → its cells scatter across all banks (safe, black). A column highlighted in red → all its cells map to one bank (the disaster).

Step 9 — The fix: pad to width 33

WHAT. Declare one extra column: __shared__ float m[32][33]. Now $[r] [c]$ lives at word $r \cdot 33 + c$ . The column read m[tid][col] gives thread $i$ the word $i \cdot 33 + col$ , so: $bank = (i \cdot 33 + col) mod 32 = (i + col) mod 32.$

WHY does $33$ work? Because $33 = 32 + 1$ , so $33 \equiv 1 (mod 32)$ . Multiplying the row index by 33 inside a mod-32 world is the same as multiplying by 1: $i \cdot 33 = i \cdot (32 + 1) = i \cdot 32 + i \equiv i (mod 32)$ , because that $i \cdot 32$ chunk is a whole multiple of 32 and vanishes under $mod 32$ . So the effective stride between consecutive rows collapses from 32 down to 1. Each next row shifts the bank by exactly one drawer instead of landing on the same one.

Now finish the count for a full warp of 32 threads. With effective stride 1, thread $i$ lands on bank $(i + col) mod 32$ . As $i$ runs through all 32 values $0, 1, \dots, 31$ , the quantity $i + col$ runs through 32 consecutive integers, and 32 consecutive integers taken $mod 32$ hit all 32 distinct remainders exactly once — just shifted (rotated) by $col$ . So every one of the 32 threads lands on a different bank: conflict degree $= g cd (1, 32) = 1$ , conflict-free, for the price of one wasted padding column.

PICTURE. The padded grid. A column read now walks diagonally across the banks — thread $i$ lands on bank $(i + col) mod 32$ , one step further each time. All 32 distinct → conflict-free. The single wasted padding column (red) is the small price paid.

Worked example Pad the inner dimension to 33

// BAD: column read is 32-way conflict
__shared__ float m[32][32];
float bad = m[threadIdx.x][col];    // stride 32 -> all bank col -> 32 cycles
 
// GOOD: one extra padding column
__shared__ float m[32][33];
float good = m[threadIdx.x][col];   // bank = (tid + col) mod 32 -> all distinct -> 1 cycle

The [33] costs one unused float per row (128 bytes for a 32-row tile) and buys back a full 32× on every column access.

The one-picture summary

Everything above is one story: round-robin dealing of words into 32 drawers, and how the stride decides whether threads share a drawer — with the one escape hatch that reading the same word is a free broadcast. The final figure compresses all of it — the bank rule, the collision law $g cd (s, 32)$ , the stride-0 broadcast exception, the stride-32 column disaster, and the $33$ -padding rescue.

Recall Feynman retelling — say it in plain words

Shared memory is 32 drawers. Words are dealt into them like cards around a circle of 32 people: word 0 to drawer 0, ..., word 31 to drawer 31, word 32 back to drawer 0. A group of 32 threads (a warp) can all be served in one tick only if each reaches into a different drawer for a different item. If two want different items from the same drawer, they take turns — that is a bank conflict. But if they want the same item, the drawer reads it once and copies it to everyone for free — a broadcast (that's why stride 0 is not a conflict, even though $g cd (0, 32) = 32$ would fib and say it is). Whether distinct-word threads collide depends only on the stride: how far apart in words consecutive threads reach. Colliding threads repeat every $32/ g cd$ steps, so exactly $g cd (stride, 32)$ of them jam into each drawer — that count is the conflict degree. Odd strides never collide. Stride 32 is the nightmare — everyone hits drawer 0 for a different item. That nightmare is exactly "reading down a column of a 32-wide matrix." The rescue: make the matrix 33 wide. Because 33 is one more than 32, each next row shifts by exactly one drawer, so the 32 rows land on 32 different drawers — conflict gone, for the price of one wasted column. (All numbers here are for 32 banks / 4-byte width; on 16-bank or 8-byte hardware, swap the constant and the same story holds.)

Recall Quick self-check

What is the conflict degree of stride 3? ::: $g cd (3, 32) = 1$ → conflict-free. Why does a column read of float m[32][32] cost 32 cycles? ::: The stride is 32 words; $(i \cdot 32 + c) mod 32 = c$ for all threads, each wanting a different word → all hit one bank. Why exactly 33 and not 34? ::: $33 \equiv 1 (mod 32)$ , giving effective stride 1. 34 gives $34 \equiv 2$ , so stride 2 → 2-way conflict — worse than 33. All 32 threads read shared[0] — how many cycles? ::: 1 — same address is a broadcast, not a conflict, even though $g cd (0, 32) = 32$ .

Where to go next: the padding trick in action → 6.3.01-Matrix-transpose-optimization; the sibling access rule for global memory → 6.2.08-Memory-coalescing; how shared-memory usage caps how many warps run → 6.2.10-Occupancy-and-resource-limits; and the cost of threads taking different paths → 6.2.11-Warp-divergence.