Bank conflicts in shared memory: Exercises — Hinglish deep dive

6.2.9 · D4 · Hardware › GPU Architecture › Bank conflicts in shared memory

Intuition Is page kaise use karein

Neeche har problem ek hi fact pe tikaa hai jo tumne Bank conflicts in shared memory mein seekha tha: ek 4-byte word jiska word-index $w$ hai, woh bank $w mod 32$ mein rehta hai, aur ek warp of threads jinka word-stride $s$ hai, unka collision degree $g cd (s, 32)$ hota hai. Aur kuch nahi chahiye. Har problem ko pehle solution collapsed rakh ke try karo, phir open karo.

Hum teen anchor facts baar baar use karte hain, toh unhe ek baar plain words mein naam de lete hain:

Word. Ek 4-byte chunk. GPU unhe $0, 1, 2, \dots$ number karta hai — word $w$ bytes $4 w$ se $4 w + 3$ tak cover karta hai.
Bank. 32 parallel "cash registers" mein se ek. Word $w$ ko bank $== w mod 32 ==$ serve karta hai.
Stride $s$ . Gap, words mein, neighbouring threads jo padhte hain unke beech. Thread $i$ word $i \cdot s$ padhta hai (plus ek constant offset jo kabhi bank pattern nahi badlata).

Figure — Bank conflicts in shared memory

Level 1 — Recognition

L1.1 — Word index se bank padhna

Ek thread shared[70] read karta hai jahan shared ek float array hai. Yeh kis bank ko hit karta hai?

Recall Solution

WHAT: float = ek 4-byte word, toh word index $= 70$ directly hai (koi extra division nahi; hum already words mein count kar rahe hain, bytes mein nahi). WHY: bank $= w mod 32$ . $70 mod 32 = 70 - 64 === 6 ==$ Answer: bank 6.

L1.2 — Byte address se bank

Ek thread shared memory mein byte address $288$ access karta hai (4-byte banks). Kaun sa bank?

Recall Solution

WHY pehle 4 se divide karein: address bytes mein hai, toh word index nikalne ke liye $⌊ 288/4 ⌋ = 72$ karein. WHY mod 32: 32 banks pe round-robin. $72 mod 32 = 72 - 64 === 8 ==$ Answer: bank 8.

Level 2 — Application

L2.1 — Stride se conflict degree nikalna

Ek warp run karta hai float v = shared[tid * 8];. Conflict degree kya hai, aur access kitne cycles leti hai?

Recall Solution

WHAT: stride in words $s = 8$ hai. WHY gcd: do threads tab collide karte hain jab unke word indices 32 ke multiple se differ karein; har bank pe collide karne walon ki sankhya $g cd (s, 32)$ hoti hai. $g cd (8, 32) === 8 ==$ Answer: 8-way conflict → access 8 cycles mein serialize ho jaati hai.

L2.2 — Odd stride ka surprise

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

Recall Solution

$g cd (7, 32) = 1 \Rightarrow== no conflict ==$ Kyunki 7 ka 32 ke saath ( $2^{5}$ hai) koi common factor nahi hai, map $i \mapsto 7 i mod 32$ saare 32 banks ko hit karta hai — ek permutation. 1 cycle.

L2.3 — Ek `double` array

__shared__ double d[64]; double v = d[tid];. 4-byte banks mein sochte hue, kya yeh conflict-free hai?

Recall Solution

WHY "8 se divide" mat karo: banks physically 4 bytes wide hain — koi 8-byte bank exist nahi karta. Ek double do consecutive banks occupy karta hai (low half + high half). Thread $i$ ke do 4-byte halves word indices $2 i$ aur $2 i + 1$ pe hain. Warp bhar mein low halves banks $0, 2, 4, \dots$ mein aur high halves $1, 3, 5, \dots$ mein jaate hain; hardware do phases mein dono halves issue karta hai taaki ek phase mein har bank exactly ek thread ko serve kare. Answer: conflict-free seedhe d[tid] ke liye.

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;

.x read karne ke liye stride aur conflict degree compute karo.

Recall Solution

Struct size $= 3$ words, toh consecutive .x fields 3 words apart hain ⇒ $s = 3$ . $g cd (3, 32) = 1 \Rightarrow== conflict-free ==$ Thread $i$ bank $3 i mod 32$ hit karta hai, jo saare 32 banks ki permutation hai (T0→0, T1→3, …, T11→33 mod 32 = 1, …). 1 cycle.

L3.2 — Woh AoS jo actually daata hai

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

.a pe conflict degree?

Recall Solution

Struct size $= 4$ words ⇒ $s = 4$ . $g cd (4, 32) = 4 \Rightarrow== 4-way conflict ==$ Ek "nice" power-of-two struct silently har field pe 4-way conflict create karta hai. 4 cycles.

L3.3 — Ek square tile ka column walk

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

Dikhao ki har thread same bank hit karta hai.

Recall Solution

Thread $i$ word $32 i + col$ read karta hai. Stride $s = 32$ . $(32 i + col) mod 32 = (0 32 i mod 32 + col) mod 32 = col mod 32$ $i$ se independent ⇒ saare 32 threads bank $col mod 32$ mein land karte hain ⇒ 32-way conflict, 32 cycles. Figure dekho.

Level 4 — Synthesis

L4.1 — Conflict kill karne ke liye tile ko pad karo

L3.3 ko __shared__ float m[32][33]; declare karke fix karo. Dikhao ki column walk conflict-free ho jaata hai.

Recall Solution

WHAT badla: row stride ab 33 words hai (ek padding column). Element $[r] [c]$ word $33 r + c$ pe hai. Thread $i$ word $33 i + col$ read karta hai. WHY 33 kaam karta hai: $33 \equiv 1 (mod 32)$ , toh $(33 i + col) mod 32 = (i + col) mod 32$ $i = 0..31$ ke liye, $(i + col) mod 32$ saare 32 distinct banks sweep karta hai ⇒ conflict-free, 1 cycle. Padding ka cost ek wasted column per row hai (32 floats = 128 bytes).

L4.2 — 16-wide tile ke liye padding choose karo

Tumhare paas __shared__ float t[32][16]; hai aur tum columns walk kar rahe ho (t[tid][col], stride 16). Ek single padding value $p$ (row width ko $16 + p$ banate hue) kya hai jo sabse chhoti ho aur column walk ko conflict-free banaye?

Recall Solution

Row width ban jaati hai $16 + p$ words = naya stride $s$ . Humein chahiye $g cd (16 + p, 32) = 1$ , matlab $16 + p$ odd hona chahiye. $16$ even hai, toh use odd banane wala sabse chhota $p$ hai $== p = 1 ==$ (width 17). Check: $g cd (17, 32) = 1$ ✓. Answer: $p = 1$ , width 17.

Level 5 — Mastery

L5.1 — General stride cycle count

Ek warp stride $s$ words use karta hai. Cycles ki sankhya $s$ ke function ke roop mein likho, phir $s = 6$ aur $s = 24$ ke liye evaluate karo.

Recall Solution

General: cycles $= g cd (s, 32)$ (har colliding group serialize hota hai).

$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$ ek factor of 2 share karta hai $32 = 2^{5}$ ke saath; $24 = 2^{3} \cdot 3$ teen share karta hai.

L5.2 — Transpose kernel diagnosis

6.3.01-Matrix-transpose-optimization mein ek kernel tile[threadIdx.y][threadIdx.x] = in[...] (row write) aur baad mein out[...] = tile[threadIdx.x][threadIdx.y] (transposed read) karta hai tile[32][32] ke saath. Kaun sa access conflict karta hai, kitna, aur one-word fix kya hai?

Recall Solution

Write tile[ty][tx]: ek warp mein tx column pe vary karta hai → stride 1 → conflict-free.
Read tile[tx][ty]: ek warp mein tx row index pe vary karta hai → stride 32 → $g cd (32, 32) = 32$ → 32-way conflict, 32 cycles.
Fix: tile[32][33] declare karo. Row stride 33 ≡ 1 (mod 32) ⇒ read conflict-free ho jaata hai (L4.1). Total wasted shared memory: $32 \times 1 \times 4 === 128 ==$ bytes. Iseelie transpose optimization exist karta hai — yeh seedha 6.2.10-Occupancy-and-resource-limits se link karta hai kyunki woh extra column tumhara shared-memory budget khaata hai.

L5.3 — Broadcast ek conflict NAHI hai

float v = shared[0]; — har thread same word 0 read karta hai. Kitne cycles?

Recall Solution

Same address ≠ conflict. Bank conflict ke liye same bank mein threads ko different addresses chahiye. Jab saare 32 threads identical word request karte hain, hardware use ek single fetch mein broadcast kar deta hai. Answer: 1 cycle. (Compare karo shared[tid*32] se, jahan addresses differ karte hain lekin saare bank 0 pe map hote hain → 32-way conflict.)

Recall Do-line self test

Stride 8 conflict degree? ::: $g cd (8, 32) = 8$ → 8-way Padding jo float[32][N] ke stride-32 column walk ko fix kare? ::: koi bhi $N$ jiska $N mod 32$ odd ho; sabse chhota $N = 33$ hai

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