Bank conflicts in shared memory: Exercises — Hardware deep dive

Figure — Bank conflicts in shared memory

Level 1 — Recognition

L1.1 — Read the bank off a word index

A thread reads shared[70] where shared is a float array. Which bank does it hit?

Recall Solution

WHAT: float = one 4-byte word, so word index $= 70$ directly (no extra division; we are already counting in words, not bytes). WHY: bank $= w mod 32$ . $70 mod 32 = 70 - 64 === 6 ==$ Answer: bank 6.

L1.2 — Byte address to bank

A thread accesses byte address $288$ in shared memory (4-byte banks). Which bank?

Recall Solution

WHY divide by 4 first: the address is in bytes, so convert to a word index by $⌊ 288/4 ⌋ = 72$ . WHY mod 32: round-robin over 32 banks. $72 mod 32 = 72 - 64 === 8 ==$ Answer: bank 8.

Level 2 — Application

L2.1 — Conflict degree from stride

A warp runs float v = shared[tid * 8];. What is the conflict degree, and how many cycles does the access take?

Recall Solution

WHAT: stride in words is $s = 8$ . WHY gcd: two threads collide when their word indices differ by a multiple of 32; the number colliding per bank is $g cd (s, 32)$ . $g cd (8, 32) === 8 ==$ Answer: 8-way conflict → the access is serialized into 8 cycles.

L2.2 — Odd stride surprise

float v = shared[tid * 7];. Conflict degree?

Recall Solution

$g cd (7, 32) = 1 \Rightarrow== no conflict ==$ Because 7 shares no factor with 32 (which is $2^{5}$ ), the map $i \mapsto 7 i mod 32$ hits all 32 banks — a permutation. 1 cycle.

L2.3 — A `double` array

__shared__ double d[64]; double v = d[tid];. Reasoning in 4-byte banks, is this conflict-free?

Recall Solution

WHY not "divide by 8": banks are physically 4 bytes wide — there is no 8-byte bank. A double occupies two consecutive banks (low half + high half). Thread $i$ 's two 4-byte halves sit at word indices $2 i$ and $2 i + 1$ . Across the warp the low halves land in banks $0, 2, 4, \dots$ and high halves in $1, 3, 5, \dots$ ; hardware issues the two halves in two phases so that within a phase every bank serves exactly one thread. Answer: conflict-free for straightforward d[tid].

Level 3 — Analysis

L3.1 — Array-of-structures, `.x` field

struct P { float x, y, z; };      // 12 bytes = 3 words
__shared__ P p[32];
float px = p[tid].x;

Compute the stride and conflict degree for reading .x.

Recall Solution

Struct size $= 3$ words, so consecutive .x fields are 3 words apart ⇒ $s = 3$ . $g cd (3, 32) = 1 \Rightarrow== conflict-free ==$ Thread $i$ hits bank $3 i mod 32$ , a permutation of all 32 banks (T0→0, T1→3, …, T11→33 mod 32 = 1, …). 1 cycle.

L3.2 — The AoS that does bite

struct Q { float a, b, c, d; };   // 16 bytes = 4 words
__shared__ Q q[32];
float qa = q[tid].a;

Conflict degree on .a?

Recall Solution

Struct size $= 4$ words ⇒ $s = 4$ . $g cd (4, 32) = 4 \Rightarrow== 4-way conflict ==$ A "nice" power-of-two struct silently creates a 4-way conflict on every field. 4 cycles.

L3.3 — Column walk of a square tile

__shared__ float m[32][32];       // row-major, word(r,c)=32r+c
float v = m[tid][col];            // fixed col, tid = row

Show every thread hits the same bank.

Recall Solution

Thread $i$ reads word $32 i + col$ . Stride $s = 32$ . $(32 i + col) mod 32 = (0 32 i mod 32 + col) mod 32 = col mod 32$ Independent of $i$ ⇒ all 32 threads land in bank $col mod 32$ ⇒ 32-way conflict, 32 cycles. See figure.

Level 4 — Synthesis

L4.1 — Pad the tile to kill the conflict

Fix L3.3 by declaring __shared__ float m[32][33];. Show the column walk becomes conflict-free.

Recall Solution

WHAT changed: row stride is now 33 words (one padding column). Element $[r] [c]$ sits at word $33 r + c$ . Thread $i$ reads word $33 i + col$ . WHY 33 works: $33 \equiv 1 (mod 32)$ , so $(33 i + col) mod 32 = (i + col) mod 32$ As $i = 0..31$ , $(i + col) mod 32$ sweeps all 32 distinct banks ⇒ conflict-free, 1 cycle. The padding cost is one wasted column per row (32 floats = 128 bytes).

L4.2 — Choose padding for a 16-wide tile

You have __shared__ float t[32][16]; and walk columns (t[tid][col], stride 16). What single padding value $p$ (making the row width $16 + p$ ) is the smallest that makes the column walk conflict-free?

Recall Solution

Row width becomes $16 + p$ words = the new stride $s$ . We need $g cd (16 + p, 32) = 1$ , i.e. $16 + p$ must be odd. $16$ is even, so the smallest $p$ making it odd is $== p = 1 ==$ (width 17). Check: $g cd (17, 32) = 1$ ✓. Answer: $p = 1$ , width 17.

Level 5 — Mastery

L5.1 — General stride cycle count

A warp uses stride $s$ words. Write the number of cycles as a function of $s$ , then evaluate for $s = 6$ and $s = 24$ .

Recall Solution

General: cycles $= g cd (s, 32)$ (each colliding group serialized).

$s = 6$ : $g cd (6, 32) = 2 \Rightarrow== 2 ==$ cycles.
$s = 24$ : $g cd (24, 32) = 8 \Rightarrow== 8 ==$ cycles. Why: $6 = 2 \cdot 3$ shares one factor of 2 with $32 = 2^{5}$ ; $24 = 2^{3} \cdot 3$ shares three.

L5.2 — Transpose kernel diagnosis

In 6.3.01-Matrix-transpose-optimization a kernel does tile[threadIdx.y][threadIdx.x] = in[...] (row write) and later out[...] = tile[threadIdx.x][threadIdx.y] (transposed read) with tile[32][32]. Which access conflicts, by how much, and what is the one-word fix?

Recall Solution

Write tile[ty][tx]: within a warp tx varies over the column → stride 1 → conflict-free.
Read tile[tx][ty]: within a warp tx varies over the row index → stride 32 → $g cd (32, 32) = 32$ → 32-way conflict, 32 cycles.
Fix: declare tile[32][33]. Row stride 33 ≡ 1 (mod 32) ⇒ read becomes conflict-free (L4.1). Total wasted shared memory: $32 \times 1 \times 4 === 128 ==$ bytes. This is why the transpose optimization exists — it links straight to 6.2.10-Occupancy-and-resource-limits because that extra column eats into your shared-memory budget.

L5.3 — Broadcast is NOT a conflict

float v = shared[0]; — every thread reads the same word 0. How many cycles?

Recall Solution

Same address ≠ conflict. A bank conflict needs threads at different addresses in the same bank. When all 32 threads request the identical word, the hardware broadcasts it in a single fetch. Answer: 1 cycle. (Contrast with shared[tid*32], where the addresses differ but all map to bank 0 → 32-way conflict.)

Recall Two-line self test

Stride 8 conflict degree? ::: $g cd (8, 32) = 8$ → 8-way Padding that fixes a stride-32 column walk of float[32][N]? ::: any $N$ with $N mod 32$ odd; smallest is $N = 33$

Related: 6.2.08-Memory-coalescing · 6.2.07-Shared-memory-architecture · 6.2.11-Warp-divergence