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

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

Yeh page Bank conflicts in shared memory ka drill floor hai. Hum ek hi master tool lete hain — bank formula aur uska conflict degree — aur isko har tarah ke access pattern pe try karte hain jab tak kuch bhi surprise nahi kar sakta.

Kuch bhi compute karne se pehle, chaliye un do tools ko dubara earn karte hain jinpe hum poore time lean karenge.

Offset $b$ conflict degree ko kabhi nahi badalta — yeh sirf har thread ko ek hi amount se slide karta hai, yeh shift karta hai ki kaun se banks use ho rahe hain, na ki kitne threads ek bank share karte hain. Ise apne paas rakh lo.

Scenario matrix

Is topic ka har access pattern in cells mein se ek mein fit hota hai. Aage ke worked examples mein har ek ko ek cell ke saath tag kiya gaya hai.

Cell	Kya ise alag banata hai	Covered by
A. Stride-1 (degenerate, best case)	$s = 1$ , gcd $= 1$	Ex 1
B. Odd stride ≠ 1	$s$ odd, coprime to 32, phir bhi conflict-free	Ex 2
C. Even stride, partial conflict	$s$ 32 ke saath factor $2^{k}$ share karta hai	Ex 3
D. Worst case (fully serialized)	$s$ 32 ka multiple, gcd $= 32$	Ex 4
E. Nonzero offset $b$ (sign/shift check)	offset present — kya yeh answer badalta hai?	Ex 5
F. 2-D array mein column walk + padding fix	array shape se derived stride	Ex 6 (figure)
G. AoS struct-size trap	`sizeof(struct)` se implied stride	Ex 7
H. Sub-word / broadcast (zero effective stride)	saare threads → ek address (special, conflict NAHIN)	Ex 8
I. Real-world word problem	histogram / reduction access	Ex 9
J. Exam twist (multi-word / double, ≥32 threads)	8-byte type, 4-byte banks mein reason karo	Ex 10

Hum limiting behaviour bhi dekhenge: jaise $s$ powers of two se climb karta hai, har step pe conflict degree double hoti jaati hai (Ex 3 → Ex 4), aur jaise $s$ dobara odd ho jaata hai conflicts gayab ho jaate hain (Ex 2). Uska visual neeche hai.

Figure — Bank conflicts in shared memory

Worked examples

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

float s[128]; float v = s[threadIdx.x]; Ek warp ke liye kitne cycles?

Forecast: aage padhne se pehle cycle count guess karo. (Hint: yeh woh case hai jiske liye banks banaye gaye the.)

Line identify karo. Thread $i$ ka word index $i \cdot 1 + 0$ hai, to $s = 1, b = 0$ . Yeh step kyun? Aage ka sab kuch $s$ aur $b$ chahiye; inhe seedha index expression se padhlo.
Conflict degree $= g cd (1, 32) = 1$ . Yeh step kyun? Gcd hi ek cheez hai jo serialization decide karta hai — degree 1 matlab har bank zyada se zyada ek thread serve karta hai.
Banks used: $bank (i) = i mod 32 = i$ for $i = 0..31$ → saare 32 banks, har ek baar. Yeh step kyun? Degree ko geometrically confirm karta hai: 32 distinct banks, koi pile-up nahin.

Answer: 1 cycle.

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

Worked example Ex 2 — Cell B · odd stride ≠ 1 (the "surprisingly safe" case)

float s[512]; float v = s[threadIdx.x * 7]; Stride 7. Conflict hai?

Forecast: bahut log kehte hain "stride 7, zaroor conflict hoga." Kya hota hai?

Line padho: $s = 7, b = 0$ .
Degree $= g cd (7, 32)$ . Kyunki 7 odd hai aur 32 $2^{5}$ hai, inke beech koi common factor nahin → $g cd = 1$ . Yeh step kyun? 32 ka sirf prime factor 2 hai; koi bhi odd stride automatically coprime hota hai → conflict-free.
Banks: $i \mapsto (7 i) mod 32$ saare 32 banks ka ek permutation hai (T0→0, T1→7, T2→14, T3→21, T4→28, T5→3, …). $i = 32$ se pehle koi repeat nahin. Yeh step kyun? Dikhata hai ki map har bank ko ek baar hit karta hai — ek scrambled order, lekin phir bhi sab distinct.

Answer: conflict-free, 1 cycle.

Verify: 32 values ${(7 i) mod 32}$ sab distinct hain (ek permutation), max degree $= 1$ . ✓

Worked example Ex 5 — Cell E · nonzero offset (kya shift kuch change karta hai?)

float s[512]; float v = s[threadIdx.x * 6 + 5]; Stride 6, offset 5.

Forecast: kya 5 add karne se Ex 3 se better, worse, ya same result milta hai?

Line padho: $s = 6, b = 5$ .
Degree $= g cd (6, 32) = 2$ — $b$ se unchanged. Yeh step kyun? Collisions $(i - j) \cdot s$ pe depend karti hain; $b$ do threads ke beech cancel ho jaata hai. Offset shift karta hai ki kaun se banks use ho rahe hain, pile-up count kabhi nahin.
Banks ab 5 se start: bank $(0) = 5$ , bank $(16) = (101) mod 32 = 5$ — phir bhi collide kar rahe hain, bas bank 0 ki jagah bank 5 pe. Yeh step kyun? Confirm karta hai ki shift used banks ka ek rigid rotation hai.

Answer: phir bhi 2-way, 2 cycles. Offset conflict degree ke liye ek red herring hai.

Verify: bank(0) $= (0 \cdot 6 + 5) mod 32 = 5$ , bank(16) $= (96 + 5) mod 32 = 5$ collide karte hain; degree $= 2$ . ✓

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

__shared__ float M[32][32]; row-major, to M[r][c] word $r \cdot 32 + c$ hai. Warp M[tid][col] karta hai (fixed column, walking rows). Phir padded array M[32][33] ke saath repeat karo.

Forecast: padding se pehle aur baad mein cycle count guess karo.

Unpadded stride. Thread $i$ word $i \cdot 32 + col$ read karta hai → $s = 32, b = col$ . Yeh step kyun? Ek row neeche jaane pe 32 words jump hota hai — row length hi stride hai.
Degree $= g cd (32, 32) = 32$ → 32-way conflict, 32 cycles. Saari rows bank $(col mod 32)$ share karti hain.
Inner dimension ko 33 kar do. Ab M[r][c] word $r \cdot 33 + c$ hai, to thread $i$ reads $i \cdot 33 + col$ → $s = 33$ . Yeh step kyun? Hum deliberately "multiple of 32" stride ko tod rahe hain.
Naya degree $= g cd (33, 32) = 1$ (33 odd hai) → conflict-free, 1 cycle. Kyunki $33 \equiv 1 (mod 32)$ , bank $(i) = (i + col) mod 32$ — consecutive rows exactly ek bank shift karti hain. Yeh step kyun? Har row mein ek wasted column 32× slowdown ko full speed mein convert kar deta hai; yeh standard transpose trick hai.

Answer: unpadded 32 cycles → padded 1 cycle.

Verify: unpadded banks sab $col mod 32$ ke barabar hain (degree 32); padded banks ${(i + col) mod 32}$ sab distinct hain (degree 1). ✓

Yeh exactly woh pattern hai jo matrix transpose optimization ke peeche hai, aur yeh aapke kernel-tuning toolkit mein memory coalescing ke bilkul saath baithta hai.

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;      // .a field read karo

Conflict hai?

Forecast: yeh stride aap nahin chunte — struct size chunti hai. Degree guess karo.

Struct size in words: $16/4 = 4$ . Adjacent .a fields 4 words apart hain → $s = 4$ , offset $b = 0$ (field pehla hai). Yeh step kyun? AoS mein compiler-chosen sizeof(struct) aapka stride ban jaata hai chahe aap chaaho ya na chaaho.
Degree $= g cd (4, 32) = 4$ → 4-way conflict, 4 cycles. Yeh step kyun? Ek "nice" power-of-two struct silently 32 ke saath factor 4 share karta hai.
Contrast: ek 3-word struct (float a,b,c;) $g cd (3, 32) = 1$ deta hai → conflict-free. Same code shape, ulta result, purely struct size ki wajah se.

Answer: 4-way, 4 cycles 16-byte struct ke liye.

Fix: structure-of-arrays pe switch karo (float a[32], b[32], …) → har field access stride-1 hai → guaranteed 1 cycle, struct size se independent.

Verify: bank $(i) = (4 i) mod 32$ sirf 8 distinct values (0,4,…,28) leta hai, har ek 4 threads share karte hain → degree 4. ✓

Worked example Ex 8 — Cell H · broadcast (saare threads, ek address — conflict NAHIN)

float s[128]; int j = 7; float v = s[j]; (saare 32 threads same word 7 read karte hain)

Forecast: saare 32 bank 7 pe jaate hain — surely 32 cycles Ex 4 ki tarah?

Effective stride $s = 0$ : har thread ka word index $7$ hai regardless of $i$ . Yeh step kyun? Yeh degenerate zero-stride case hai jise gcd formula cover nahin karta — same address, alag-alag addresses nahin.
Conflict ki definition yaad karo: do threads ko ek hi bank mein alag-alag addresses chahiye. Yahan unhe ek hi address chahiye. Yeh step kyun? Hardware 1 cycle mein ek single word saare requesting threads ko broadcast karta hai — koi serialization nahin.
To yeh 32 nahin, 1 cycle hai. Zero stride worst case ka ulta hai.

Answer: 1 cycle (broadcast). Lesson: bank conflict ke liye distinct addresses chahiye; identical-address reads free hain.

Verify: saare threads word 7 request karte hain → ek unique address → broadcast → 1 cycle. ✓

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

Ek warp 32 samples ko ek shared histogram mein bin karta hai: atomicAdd(&hist[sample[tid]], 1);. Worst input mein har sample bin 5 mein jaata hai. Best input mein samples bins $0..31$ pe map hote hain (ek ek). Dono cases mein shared-memory address phase ka cycle cost kya hai?

Forecast: kaun sa input fast hai, kaun sa slow, aur kitna difference hai?

Worst input: sab sample[tid] = 5 → saare threads word 5 address karte hain → yeh same address hai (broadcast-style), lekin atomicAdd us ek bank pe read-modify-write serialize karna padta hai → effectively 32 serialized updates. Yeh step kyun? Same address pe atomics serialize hote hain chahe plain read broadcast hoti — write side ordering force karta hai.
Best input: bins $0..31$ ka permutation hain → 32 distinct banks → saare atomics 1 cycle mein proceed karte hain. Yeh step kyun? Distinct banks = independent atomics = full parallelism, histogram ideal.
Design takeaway: heavily-skewed data 32× slowdown drive karta hai; yahi per-thread private sub-histograms phir merge choose karne ki wajah hai.

Answer: worst-case 32 cycles, best-case 1 cycle address phase ke liye.

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), 4-byte banks mein reason karo

__shared__ double d[64]; double v = d[threadIdx.x]; Ek double 8 bytes = 2 words hota hai. Kya natural per-thread access conflict karta hai? Sirf 4-byte banks mein reason karo (koi "8-byte bank" nahin hota).

Forecast: 8 bytes apart — kya stride-2-jaisi layout 2-way conflict cause karti hai?

Words lay out karo. Thread $i$ ka double words $2 i$ (low half) aur $2 i + 1$ (high half) occupy karta hai. Byte address $8 i$ . Yeh step kyun? Fancy type ko hamesha raw 4-byte word indices mein convert karo — yahi ek unit hai jo banks samajhte hain.
Naive fear: low halves words $0, 2, 4, \dots$ pe baithte hain → stride 2 → $g cd (2, 32) = 2$ , lagta hai 2-way. Yeh step kyun? Dikhata hai trap intuition kahan se aati hai.
Hardware actually kya karta hai: warp ki 64-word request ko do phases mein split karta hai. Phase 1 threads 0–15 serve karta hai (words 0–31, banks 0–31, sab distinct); phase 2 threads 16–31 serve karta hai (words 32–63, banks 0–31, sab distinct). Har phase saare 32 banks ko bina repeat ke touch karta hai. Yeh step kyun? Yeh do phases hi woh tarika hai jisse modern GPUs d[tid] ko doubles pe conflict-free banate hain, 2-word footprint ke bawajood.

Answer: conflict-free — straightforward d[tid] double access no conflict ke equivalent cost karta hai (2 phases, har ek 1 cycle, koi serialization within a phase nahin).

Verify: phase-1 low-half words ${0, 2, \dots, 30} \Rightarrow$ banks ${0, 2, \dots, 30}$ distinct; high halves ${1, 3, \dots, 31}$ distinct; to words $0..31$ banks $0..31$ ko ek ek baar hit karte hain, degree 1. ✓

Recall Quick self-test (guess karne ke baad reveal karo)

Stride 12 — conflict degree? ::: $g cd (12, 32) = 4$ → 4-way conflict. Stride 15 — conflict degree? ::: $g cd (15, 32) = 1$ (15 odd hai) → conflict-free. Kya offset $+ 9$ add karne se koi conflict degree badlta hai? ::: Nahin — offset shift karta hai ki kaun se banks use ho rahe hain, pile-up count kabhi nahin. Saare 32 threads same word read karte hain — kitne cycles? ::: 1 (broadcast), kyunki yeh same address hai, alag-alag nahin. [32][32]→[32][33] padding ek column walk fix karta hai kyunki 33 mod 32 = ? ::: $1$ , to har row exactly ek bank shift karti hai.

Aage kahan jaayein: Conflict-free shared memory directly occupancy and resource limits mein feed karta hai (shared memory per block decide karta hai ki kitne warps fit hote hain), aur column-walk fix matrix transpose optimization ka dil hai. Inhe warp divergence ke saath combine karte waqt savdhan raho, jo ek alag wajah se serialize karta hai (branches, banks nahin), aur memory coalescing ke saath, jo is same "accesses ko line up karo" idea ka global-memory cousin hai. Layout rules shared memory architecture se trace back hote hain.