Bank conflicts in shared memory

Why this matters: Shared memory is supposed to be fast (~100× faster than global memory). Bank conflicts destroy this advantage by serializing what should be parallel access.

How shared memory banks work

Derivation from first principles:

1. Why divide by 4? Each bank is 4 bytes wide (the fixed bank width on NVIDIA hardware). A byte address must be converted to a word index. Address 0-3 → word 0, address 4-7 → word 1, etc.
2. Why mod 32? With 32 banks, we distribute words in a round-robin fashion. Word 0 → bank 0, word 1 → bank 1, .., word 31 → bank 31, word 32 → bank 0 again.
3. Result: Words are striped across banks. Consecutive 32-bit words live in consecutive banks.

Types of access patterns

Why this works:

• Thread 0 reads shared[0] → word 0 → bank 0
• Thread 1 reads shared[1] → word 1 → bank 1
• ...
• Thread 31 reads shared[31] → word 31 → bank 31

All 32 threads hit different banks → 1 cycle access, perfect parallelism.

Analysis:

• Thread 0: shared[0] → bank 0
• Thread 1: shared[2] → bank 2
• Thread 2: shared[4] → bank 4
• ...
• Thread 16: shared[32] → word 32 → bank $(32 \mod 32) = 0$ ← conflict with thread 0!
• Thread 17: shared[34] → bank 2 ← conflict with thread 1!

Result: 16 pairs of threads collide (2-way conflicts). Access takes 2 cycles instead of 1.

Why this step? With stride-2, every 16 threads we wrap around to bank 0 again. Threads separated by 16 positions request the same bank.

• Thread 0: shared[0] → bank 0
• Thread 1: shared[32] → word 32 → bank 0
• Thread 2: shared[64] → word 64 → bank 0
• All 32 threads hit bank 0!

Access time: 32 cycles (fully serialized). Shared memory is now slower than registers.

The mathematical pattern

Derivation:

1. Thread $i$ accesses word index $i \cdot s$.
2. Bank for thread $i$: $(i \cdot s) \mod 32$.
3. Two threads $i$ and $j$ conflict if $(i \cdot s) \mod 32 = (j \cdot s) \mod 32$.
4. This simplifies to $(i - j) \cdot s \equiv 0 \pmod{32}$.
5. The smallest $k = i - j$ satisfying this is $k = 32 / \gcd(s, 32)$.
6. So every $32/\gcd(s,32)$ threads, we repeat banks, meaning $\gcd(s,32)$ threads collide per bank.

Examples:

• Stride 1: $\gcd(1,32) = 1$ → no conflict
• Stride 2: $\gcd(2,32) = 2$ → 2-way conflict
• Stride 3: $\gcd(3,32) = 1$ → no conflict (odd strides are safe!)
• Stride 16: $\gcd(16,32) = 16$ → 16-way conflict
• Stride 32: $\gcd(32,32) = 32$ → 32-way conflict

Common mistake: structure-of-arrays vs array-of-structures

Why it feels wrong at first glance: "Threads are 12 bytes apart, so surely they collide on banks."

The steel-manned truth: Let's actually compute the banks for the .x access.

• The word stride between adjacent .x fields is 3 words (12 bytes / 4).
• Thread $i$ accesses word $3i$, landing in bank $(3i) \mod 32$.
• Since $\gcd(3, 32) = 1$, the map $i \mapsto (3i) \bmod 32$ is a permutation of the 32 banks.
• Thread 0→bank 0, T1→bank 3, T2→bank 6, ..., T10→bank 30, T11→bank 33 mod 32 = bank 1, ...

So a warp reading .x (or .y, or .z) is actually CONFLICT-FREE, because 3 is coprime to 32. The intuition that "non-unit stride = conflict" is false; only strides sharing a factor with 32 (i.e. even strides / powers of 2) cause conflicts.

When AoS does bite you: conflicts appear if the struct size shares a factor with 32 — e.g. a struct of 4 floats (16 bytes = 4 words). Then stride = 4, $\gcd(4,32)=4$ → 4-way conflict on each field. That is the real danger of AoS: you don't control the stride, and a "nice" power-of-two struct size silently creates conflicts.

The robust fix: Structure-of-arrays

__shared__ float x[32], y[32], z[32];
float px = x[threadIdx.x];  // Stride-1, guaranteed conflict-free
float py = y[threadIdx.x];
float pz = z[threadIdx.x];

SoA makes every field access stride-1, so you never depend on the struct size being coprime to 32.

Padding to avoid conflicts

Case A — same row, different columns (this is conflict-free):

float value = matrix[row][tid];   // fixed row, tid = column

• Thread $i$ accesses word $\text{row} \cdot 32 + i$. As $i$ varies, the stride is 1 word.
• Banks: $(\text{row}\cdot 32 + i) \bmod 32 = i$ (since $32 \equiv 0$). All different → no conflict.

Case B — same column, different rows (this is the bad one):

float value = matrix[tid][col];   // tid = row, fixed column

• Thread $i$ accesses word $i \cdot 32 + \text{col}$. The stride is 32 words.
• Banks: $(i \cdot 32 + \text{col}) \bmod 32 = \text{col} \bmod 32$, the same for every thread → 32-way conflict!

This column-walk pattern is exactly what happens in matrix transpose, so it matters a lot.

The fix: pad the inner dimension to 33

__shared__ float matrix[32][33];  // one extra padding column
float value = matrix[tid][col];   // column access

• Now element $[r][c]$ lives at word $r \cdot 33 + c$.
• Thread $i$ (row $i$, fixed column) accesses word $i \cdot 33 + \text{col}$.
• Bank index (no extra /4 — we are already in word units): $\text{bank} = (i \cdot 33 + \text{col}) \bmod 32 = (i + \text{col}) \bmod 32$ because $33 \equiv 1 \pmod{32}$.
• As $i$ runs 0..31, $(i + \text{col}) \bmod 32$ takes all 32 distinct values → conflict-free!

Why this step? Padding to 33 makes the row stride $\equiv 1 \pmod{32}$, so consecutive rows shift by exactly one bank instead of landing on the same bank. That is the whole trick.

Now, 32 friends (a "warp") all try to grab their lunch at the same time. If they stored lunches in cubbies 1, 2, 3, .., 32, everyone gets theirs instantly—no waiting! That's conflict-free.

But what if they all stored lunch in cubby 1? Now the teacher has to hand them out one by one—super slow! That's a 32-way bank conflict.

Bank conflicts happen when multiple threads (students) want data from the same bank (cubby) at once. The GPU has to serve them one at a time instead of all together. The trick is to organize your data so each thread uses a different bank—like making sure each student has a different cubby number. Adding one "empty spare cubby" per row (padding to 33) shifts everyone by one so they stop crashing into the same cubby.

Connections

• Shared memory architecture—the physical design enabling banks
• Memory coalescing—similar parallelism concept for global memory
• Occupancy—bank conflicts reduce effective occupancy by stalling warps
• Matrix transpose—classic example where padding fixes conflicts
• Warp divergence—another serialization bottleneck in GPU execution

#flashcards/hardware

What are shared memory banks and why do they exist? :: Banks are parallel sub-units of shared memory (typically 32) that allow multiple threads to access different addresses simultaneously. They exist to enable high-bandwidth parallel access—each bank can serve one request per cycle.

What is a bank conflict? :: When multiple threads in a warp access different addresses that map to the same bank, forcing serialized access instead of parallel access. This multiplies latency by the conflict degree.

What is the conflict degree for stride-2 access? :: 2-way conflict. gcd(2, 32) = 2, so every two threads collide on the same bank, doubling access time.

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)