Bank conflicts in shared memory: Visual walkthrough

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

Yeh parent topic ka visual walkthrough hai. Hum bank conflicts ka poora idea bilkul zero se build karte hain, ek ek tasveer karke, jab tak ek simple clean law — kitne threads sabse busy drawer par jam karte hain — khud nikal ke nahi aata, aur jab tak tum yeh nahi dekh lete ki ek matrix ko width 33 tak pad karna magic ki tarah sab kuch kyun theek kar deta hai.

Shuru karne ke liye tumhe almost kuch nahi chahiye. Agar "thread" ya "shared memory" jaisa koi word naya hai, toh pehle 6.2.07-Shared-memory-architecture dekh lo, lekin jo bhi tumhe actually chahiye hoga woh main neeche re-explain kar dunga.

Step 1 — Bank actually hai kya?

KYA. Shared memory ek fast storage ka block hai jo GPU cores ke bilkul paas baitha hai. Lekin yeh koi ek single drawer nahi hai — yeh physically banks kehlaane wale 32 alag alag drawers mein cut hoti hai. Har bank exactly ek 4-byte word per clock cycle de sakta hai.

YEH slicing kyun? Kyunki 32 threads ka ek group (warp) sab ek hi instant par data chahta hai. Ek drawer jo ek item per tick serve kare woh 32 threads ko 32 lambi line mein wait karne par majboor kar deta. Battees drawers, har ek ek item per tick serve karte hue, matlab hai ki sab 32 threads ko ek hi tick mein serve kiya ja sakta hai — agar woh sab alag alag drawers mein pahunch rahe hon.

TASVEER. Figure dekho. Horizontal strip shared memory hai. Ise 32 vertical slots mein kaata gaya hai, 0 se 31 tak numbered. Woh number bank index hai. Accent-red slot bank 0 hai, jahan hamari story baar baar waapas aati hai.

Figure — Bank conflicts in shared memory

Step 2 — Koi word kaun se drawer mein rehta hai?

KYA. Humein ek rule chahiye jo humein bataye: ek word diya gaya, usे 32 banks mein se kaun sa hold karta hai? Rule round-robin hai — words ko 32 khiladiyon ko playing cards ki tarah baanto jo ek circle mein baithe hain.

Round-robin kyun? Kyunki consecutive threads usually consecutive words chahte hain. Agar consecutive words consecutive banks mein baithe hon, toh ek normal shared[tid] access naturally sab 32 drawers mein spread ho jaata hai — best possible case. Cards ko table ke aaround deal karna exactly yahi achieve karta hai.

TASVEER. Word 0 bank 0 ko jaata hai, word 1 bank 1 ko, ..., word 31 bank 31 ko — aur phir word 32 bank 0 par wrap back karta hai (accent-red). Yeh wrap-around hi aage trouble ka poora source hai.

Step 3 — Happy case: har koi alag drawer mein

KYA. Har thread $i$ (tid, running $i = 0, 1, \dots, 31$ ) ko shared[i] read karne do. Thread $i$ word $w = i$ ko touch karta hai, toh uska bank $i mod 32 = i$ hai.

YEH pehle kyun dikhayein? Yeh reference point hai — "1 cycle" ideal. Har baad ka pattern iske against judge hota hai: isके versus humne kitne extra cycles khoye?

TASVEER. Battees arrows fan out karte hain, ek per thread, har ek apne bank par utarta hai. Koi do arrows koi drawer share nahi karte. Sab ek tick mein serve ho jaate hain.

Worked example Stride-1 access

int tid = threadIdx.x;       // is thread ka number, 0..31
float value = shared[tid];   // har thread apna word read karta hai

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

Yahan "stride" ka matlab hai thread $i$ aur thread $i + 1$ jo read karte hain uske beech words mein gap. Stride-1 = woh neighbors read karte hain.

Step 4 — Ise purpose se break karo: stride-2

KYA. Ab thread $i$ ko shared[2*i] (yaani shared[tid*2]) read karne do. Stride $s = 2$ hai. Thread $i$ word $w = 2 i$ ko touch karta hai, toh uska bank $(2 i) mod 32$ hai.

YEH kyun karein? Exactly woh moment pakdne ke liye jab wrap-around bites karta hai. Hum $i$ ko upar march karte hain aur dekhte hain ki kaun sa thread pehle ek aise drawer par waapas utarta hai jis par kisi ne already claim kar rakha hai.

TASVEER. Threads 0..15 even banks $0, 2, 4, \dots, 30$ par fan karte hain. Phir thread 16 word 32 read karta hai, aur $32 mod 32 = 0$ — yeh bank 0 (accent-red) par slam karta hai, right on top of thread 0. Thread 17 thread 1 se collide karta hai, aur waise aage: 16 pairs, har pair ek drawer ke upar fighting.

(Yeh pehli baar hai jab hum "conflict degree" define kar rahe hain — har pehli line ne jaanboojhkar is term se parhez kiya jab tak hum ise ek tasveer mein hote nahi dekh sake.)

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

KYA. Ab hum ek formula dhundhte hain jo har stride $s$ ko cover kare. Do threads $i$ aur $j$ collide karte hain jab unke banks equal hote hain: $(i \cdot s) mod 32 = (j \cdot s) mod 32 ⟺ (i - j) \cdot s \equiv 0 (mod 32) .$ Right-hand side padho: " $(i - j) \cdot s$ 32 ka ek poora multiple hai."

$g cd$ kyun — aur $32/ g cd (s, 32)$ sabse chhota gap kyun hai? Hum pooch rahe hain: do colliding threads ke beech sabse chhota positive gap $k = i - j$ kya hai? Toh humein sabse chhota positive $k$ chahiye jo $k \cdot s$ ko 32 ka multiple banaye.

Yahan reason hai ki yeh $32/ g cd (s, 32)$ kyun nikalta hai, spell out karke. $g = g cd (s, 32)$ ho — woh sabse bada number jo dono $s$ aur 32 ko divide karta hai. Pieces split karo: $s = g \cdot s^{'}$ aur $32 = g \cdot n$ , jahan ab $s^{'}$ aur $n$ koi common factor share nahi karte (humne shared part divide kar diya, yahi $g cd$ karta hai). Condition " $k \cdot s$ 32 ka multiple hai" ban jaata hai " $k \cdot g \cdot s^{'}$ $g \cdot n$ ka multiple hai," yaani " $k \cdot s^{'}$ $n$ ka multiple hai." Kyunki $s^{'}$ aur $n$ ka koi common factor nahi, $s^{'}$ factor $n$ ki taraf kuch contribute nahi karta — toh $n$ poori tarah $k$ se aana chahiye. Isliye sabse chhota aisa $k$ hai $k = n = 32/ g$ . Yeh exactly $k = 32/ g cd (s, 32)$ hai.

"Sabse chhota gap" se "conflict degree" tak — missing link. Humne abhi sikha ki colliding threads har $k = 32/ g cd (s, 32)$ steps par repeat karte hain: threads $0, k, 2 k, 3 k, \dots$ sab thread 0 ki tarah ek hi bank par utarte hain, threads $1, 1 + k, 1 + 2 k, \dots$ sab doosra bank share karte hain, wagera. Toh 32 threads groups mein split ho jaate hain, har group ek bank par baitha hai, aur har group ek hi size ka hai. Har group kitna bada hai? Warp mein 32 threads hain aur woh har $k$ par repeat karte hain, toh har bank collect karta hai $\frac{32}{k} = \frac{32}{32/ g c d ( s , 32 )} = g cd (s, 32)$ threads. Woh count — threads jo sabse busy (aslmein har) occupied bank share karte hain — yahi hai Step 4 ka conflict degree. Toh sabse chhota gap $k$ aur conflict degree ek hi sikke ke do pehlu hain: $k$ batata hai colliding threads kitne door hain; $32/ k = g cd (s, 32)$ batata hai kitne unke per bank pile up hote hain.

$g cd (s, 32)$ ka matlab hai greatest common divisor — woh sabse bada number jo dono $s$ aur 32 ko divide karta hai. Hum ise isliye use karte hain kyunki collision condition ek divisibility question hai, aur $g cd$ precisely divisibility ke liye banaya gaya tool hai.

TASVEER. Strides ki ek number line. Har stride ke neeche uska collision count baitha hai. Red markers dangerous strides hain (32 ke saath 2 ka factor share karte hain); black ones safe hain (32 ke coprime). Notice karo: odd strides hamesha safe hain — ek odd number $32 = 2^{5}$ ke saath 2 ka koi factor share nahi karta, toh $g cd = 1$ .

Step 6 — Special case jo law galat get karta hai: stride 0 (broadcast)

KYA. Agar har thread same word read kare — stride $s = 0$ ? Ise naively plug in karo: $g cd (0, 32) = 32$ , toh formula chilla ke kehta hai "32-way conflict, 32 cycles!" Lekin yeh real hardware par galat hai. Jab sab threads same address request karte hain, bank use ek baar read karta hai aur value ko har thread tak broadcast karta hai — 1 cycle, koi conflict nahi.

Law yeh miss kyun karta hai. Step 4 ki definition phir se padho: ek conflict ko same bank mein different words chahiye. Stride 0 ka matlab hai har thread same word chahta hai, different nahin. Puri "ek item per tick" bottleneck isliye exist karti hai kyunki ek drawer sirf ek distinct item nikal sakta hai — lekin yahan sirf ek item nikalna hai, aur use 32 listeners tak copy karna free hai. Toh "different addresses" clause woh fine print hai jo humein bachata hai.

TASVEER. Sab 32 arrows bank 0 (accent-red) mein ek word par point karte hain, aur dashed lines ka ek single fan woh ek value wapas sab 32 threads tak le jaata hai — ek read, ek broadcast.

Step 7 — Degenerate extreme: stride-32

KYA. Stride ko $s = 32$ tak push karo. Thread $i$ word $32 i$ read karta hai, bank $(32 i) mod 32 = 0$ — har $i$ ke liye. Har thread ek different word chahta hai ( $0, 32, 64, \dots$ ), lekin woh sab bank 0 mein rehte hain. Woh "different words" part yahi hai jo ise ek genuine conflict banata hai (Step 6 ke broadcast se alag). Sab 32 bank 0 par pile up karte hain.

Yeh include kyun karein? Yeh sabse bura true conflict hai aur ek crucial edge: $g cd (32, 32) = 32$ , toh 32 cycles. Shared memory kuch clever na karne se bhi slower ho jaati hai. Har thread ek drawer par 32 lambi line mein wait karta hai.

TASVEER. Sab 32 arrows single accent-red drawer (bank 0) par converge karte hain, har ek usмен se ek different word fetch karta hua — 32 ka ek stack, ek per tick serve hota hua.

Worked example Stride-32 = 32-way conflict

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

Har thread → bank 0, har ek ka different word → 32 cycles. Yeh woh pattern hai jo secretly tab appear hota hai jab tum ek 32-wide shared matrix ka column walk karte ho.

Step 8 — Stride-32 chhup kahan jaata hai: matrix columns

KYA. Ek matrix row-major store karo: __shared__ float m[32][32]. Element $[r] [c]$ word $r \cdot 32 + c$ par baitha hai.

Row read m[row][tid]: thread $i$ → word $row \cdot 32 + i$ . Threads ke beech stride = 1. Bank $= (row \cdot 32 + i) mod 32 = i$ (kyunki $32 \equiv 0$ ). Sab distinct → free.
Column read m[tid][col]: thread $i$ → word $i \cdot 32 + col$ . Stride = 32. Bank $= col$ har thread ke liye, aur har ek alag word chahta hai → 32-way conflict.

Yeh kyun care karein? Ek column read karna exactly wahi hai jo matrix transpose karta hai (dekho 6.3.01-Matrix-transpose-optimization). Innocent-looking m[tid][col] secretly Step 7 ki stride-32 disaster hai.

TASVEER. Matrix grid. Ek row highlighted → uske cells sab banks mein scatter ho jaate hain (safe, black). Ek column red mein highlighted → uske sab cells ek bank par map karte hain (woh disaster).

Step 9 — Fix: width 33 tak pad karo

KYA. Ek extra column declare karo: __shared__ float m[32][33]. Ab $[r] [c]$ word $r \cdot 33 + c$ par rehta hai. Column read m[tid][col] thread $i$ ko word $i \cdot 33 + col$ deta hai, toh: $bank = (i \cdot 33 + col) mod 32 = (i + col) mod 32.$

$33$ kyun kaam karta hai? Kyunki $33 = 32 + 1$ , toh $33 \equiv 1 (mod 32)$ . Row index ko 33 se multiply karna mod-32 world ke andar 1 se multiply karne jaisa hai: $i \cdot 33 = i \cdot (32 + 1) = i \cdot 32 + i \equiv i (mod 32)$ , kyunki woh $i \cdot 32$ chunk 32 ka ek poora multiple hai aur $mod 32$ ke under gayab ho jaata hai. Toh consecutive rows ke beech effective stride 32 se 1 par collapse ho jaati hai. Aagla row exactly ek drawer shift karta hai instead of same par utarne ke.

Ab 32 threads ke poore warp ke liye count finish karo. Effective stride 1 ke saath, thread $i$ bank $(i + col) mod 32$ par utarta hai. Jaise $i$ sab 32 values $0, 1, \dots, 31$ ke through run karta hai, quantity $i + col$ 32 consecutive integers ke through run karti hai, aur 32 consecutive integers $mod 32$ lete waqt sab 32 distinct remainders exactly ek baar hit karte hain — bas $col$ se shift (rotated). Toh 32 threads mein se har ek ek different bank par utarta hai: conflict degree $= g cd (1, 32) = 1$ , conflict-free, sirf ek wasted padding column ki keemat par.

TASVEER. Padded grid. Ek column read ab banks ke across diagonally walk karta hai — thread $i$ bank $(i + col) mod 32$ par utarta hai, har baar ek step aage. Sab 32 distinct → conflict-free. Single wasted padding column (red) woh chhoti keemat hai jo chukani padti hai.

Worked example Inner dimension ko 33 tak pad karo

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

[33] har row ke liye ek unused float (32-row tile ke liye 128 bytes) ki cost leta hai aur har column access par pura 32× wapas khareedta hai.

Ek tasveer mein summary

Upar sab kuch ek hi story hai: words ko 32 drawers mein round-robin deal karna, aur stride kaise decide karta hai ki threads ek drawer share karte hain ya nahi — ek escape hatch ke saath ki same word read karna free broadcast hai. Final figure iska sab kuch compress karta hai — bank rule, collision law $g cd (s, 32)$ , stride-0 broadcast exception, stride-32 column disaster, aur $33$ -padding rescue.

Recall Feynman retelling — plain words mein bolo

Shared memory 32 drawers hain. Words unmein cards ki tarah deal hote hain 32 logon ke circle ke aaround: word 0 drawer 0 ko, ..., word 31 drawer 31 ko, word 32 wapas drawer 0 ko. 32 threads ka ek group (warp) sab ek tick mein serve ho sakta hai sirf tabhi agar har ek kisi alag drawer mein kisi alag item ke liye pahunche. Agar do alag drawer se different items chahte hain, toh woh turns lete hain — yeh bank conflict hai. Lekin agar woh same item chahte hain, toh drawer use ek baar read karta hai aur sab ko free mein copy karta hai — ek broadcast (isliye stride 0 ek conflict nahi hai, bhalе hi $g cd (0, 32) = 32$ keh kar jhooth bole). Distinct-word threads collide karein ya nahi yeh sirf stride par depend karta hai: consecutive threads kitne door words tak pahunchte hain. Colliding threads har $32/ g cd$ steps par repeat karte hain, toh exactly $g cd (stride, 32)$ unke har drawer mein jam hote hain — woh count conflict degree hai. Odd strides kabhi collide nahi karte. Stride 32 nightmare hai — sab drawer 0 par alag item ke liye hit karte hain. Woh nightmare exactly "32-wide matrix ka column neeche padhna" hai. Rescue: matrix ko 33 wide banao. Kyunki 33 ek zyada hai 32 se, aagla row exactly ek drawer shift hota hai, toh 32 rows 32 different drawers par utarte hain — conflict gone, ek wasted column ki keemat par. (Yahan sab numbers 32 banks / 4-byte width ke liye hain; 16-bank ya 8-byte hardware par, constant swap karo aur wahi story chalti rehti hai.)

Recall Quick self-check

Stride 3 ka conflict degree kya hai? ::: $g cd (3, 32) = 1$ → conflict-free. float m[32][32] ka column read kyun 32 cycles leta hai? ::: Stride 32 words hai; $(i \cdot 32 + c) mod 32 = c$ sab threads ke liye, har ek alag word chahta hua → sab ek bank hit karte hain. Exactly 33 kyun aur 34 kyun nahi? ::: $33 \equiv 1 (mod 32)$ , effective stride 1 deta hai. 34 deta hai $34 \equiv 2$ , toh stride 2 → 2-way conflict — 33 se bura. Sab 32 threads shared[0] read karte hain — kitne cycles? ::: 1 — same address ek broadcast hai, conflict nahi, bhalе hi $g cd (0, 32) = 32$ ho.

Aage kahan jaana hai: padding trick in action → 6.3.01-Matrix-transpose-optimization; global memory ke liye sibling access rule → 6.2.08-Memory-coalescing; shared-memory usage kitne warps run karne par cap lagata hai → 6.2.10-Occupancy-and-resource-limits; aur threads ke alag paths lene ki cost → 6.2.11-Warp-divergence.