6.2.9 · D1Bank conflicts in shared memory
Shared memory 32 parallel lanes mein split hoti hai jinhein banks kehte hain , aur ek warp ke saare 32 threads ek saath read kar sakte hain sirf tab jab har thread alag bank mein jaaye . Is topic mein sab kuch ek hi sawaal hai baar baar: "jis tarah mere threads array ko index kar rahe hain, kya unme se koi do ek hi bank mein jaate hain?"
Is sawaal ka jawab dene se pehle tumhare paas kuch words aur symbols hone chahiye. Parent note mein warp , bank , word , stride , mod, gcd, aur [r][c] addressing aise use hote hain jaise tum pehle se jaante ho. Yeh page inhe ek-ek karke, ek order mein build karta hai jahaan har piece pichle pe tikaa ho.
Ek thread ek worker hai jo tumhara program chalata hai. GPU pe tum hazaaron launch karte ho, aur sab ek hi line of code run karte hain lekin alag-alag data pe — thread number 5 element 5 handle karta hai, thread number 6 element 6 handle karta hai, aur aise hi aage.
Woh number jo ek thread ko batata hai "main kaun hoon" threadIdx.x likha jaata hai. Agar tumhare paas 32 threads hain, toh threadIdx.x 0 , 1 , 2 , … , 31 i
Ek warp 32 threads ka ek fixed bundle hai jo hardware lockstep mein run karta hai — sab ek hi instruction ek hi instant pe execute karte hain. Jab woh 32 threads shared-memory read pe pahunchte hain, saare 32 memory requests simultaneously fire hote hain. Yahi simultaneity woh reason hai kyun banks matter karte hain.
Intuition Kyun 32 baar baar aata haiEk warp mein 32 threads hain aur shared memory mein 32 banks hain. Yeh coincidence nahi hai — hardware aise bana hai ki ek perfectly-behaved warp apne 32 threads ko 32 banks pe one-to-one map kare. Har "conflict" usi one-to-one match ki failure hai.
Definition Byte aur byte-addressEk byte 8 bits hota hai — memory ka sabse chhota chunk jiska apna numbered slot hota hai. Woh number byte-address hai: byte 0, byte 1, byte 2, … Memory ko ek infinitely long ruler socho jahaan har millimetre numbered ho.
Definition Word (ek 4-byte word)Ek float (aur ek int) 4 bytes leta hai. Chaar consecutive bytes jo ek saath jod ke ek value ki tarah treat ki jaaye, usse 32-bit word kehte hain (32 bits = 4 bytes). Toh byte-addresses 0 , 1 , 2 , 3 0 andar hain; bytes 4 , 5 , 6 , 7 1
Byte-address ko word-number mein convert karne ke liye 4 se divide karo aur remainder phenk do:
word_index = ⌊ 4 byte_address ⌋
Woh ⌊ ⌋ floor symbol hain — "ek whole number pe round down karo." Humein iska zaroorat hai kyunki addresses 0, 1, 2, 3 sab ko word 0 pe collapse karna hai.
Intuition Words ki zaroorat kyun hai?Banks 4 bytes wide hain, toh bank hardware individual bytes ki parwah nahi karta — sirf yahi care karta hai ki tumne konsa 4-byte word maanga. Pehle word-index mein convert karna woh step hai jo agla formula ko messy byte details ignore karne deta hai.
Definition Modulo, likha jaata haimod ya mod a mod n remainder jo a n 7 mod 3 = 1 7 = 2 × 3 + 1 0 , 1 , … , n − 1
Picture ek clock face hai. Ek clock mod 12 mod 32
mod 32 aur kuch nahi?
Sirf 32 physical banks hone se, word number 32 ke paas koi naya jagah nahi — round-robin zaroor wrap karega. mod 32 exactly "32 piles mein cards deal karte raho, ek card per pile, har lap mein pile 0 se shuru" ki arithmetic hai. Koi bhi tool jo wrap-around capture kare kaam karega; mod iske liye standard naam hai.
s Stride woh gap hai, words mein measure kiya , thread i i + 1 shared[i] mein gap 1 word hai (stride 1). shared[i*2] mein gap 2 words hai (stride 2). matrix[i][col] mein 32 columns ke saath, ek row aage jaana 32 words jump karta hai (stride 32).
Toh thread i i ⋅ s
bank ( i ) = ( i ⋅ s ) mod 32.
Yeh akela line is poore topic ki engine hai. Parent note jo bhi compute karta hai woh is line mein specific s
Worked example Stride ko feel karo
Stride 1: banks hain 0 , 1 , 2 , … , 31 one-to-one , koi conflict nahi.
Stride 2: banks hain 0 , 2 , 4 , … , 30 , 0 , 2 , … aur thread 16 dono se hit hota hai. Collision.
Stride 32: har thread ka bank ( i ⋅ 32 ) mod 32 = 0
Definition Greatest common divisor,g cd( a , b ) g cd( a , b ) sabse bada whole number hai jo a b g cd( 2 , 32 ) = 2 g cd( 3 , 32 ) = 1 g cd( 16 , 32 ) = 16 g cd( s , 32 ) = 1 s 32 coprime hain (unka koi shared factor nahi sirf 1 ke).
sahi tool hai
Hum woh chhota step k k ( k ⋅ s ) mod 32 = 0 k ⋅ s s 32 shared factors ke baare mein pure sawaal hai. Shared-factors reasoning ka mathematical naam gcd hai. Koi aur tool "hum kitni jaldi ek used bank pe wrap karenge" ka jawab itna saaf nahi deta.
Woh payoff formula jo parent use karta hai:
conflict_degree = g cd( s , 32 )
Coprime stride (g cd= 1 i ↦ ( i ⋅ s ) mod 32 saare 32 banks hit karta hai — ek permutation, perfectly conflict-free. (Isliye odd-stride .x field access safe hai.)
g cd= d > 1 32/ d d d
Definition Row-major layoutEk 2-D array matrix[R][C] memory mein actually ek 1-D strip hai. Row-major ka matlab hai row 0 pehle laid out hoti hai, phir row 1, etc. Element [r][c] word pe hota hai:
word = r ⋅ C + c C c → c + 1 r → r + 1 C
Isliye ek column read karna (matrix[i][col], row i C C = 32 C = 33 33 mod 32 = 1
Stride s and bank of thread i
gcd predicts conflict degree
Row-major r times C plus c
Bank conflicts in shared memory
Banks ek wider 6.2.07-Shared-memory-architecture picture ke andar rehte hain.
Global memory mein wahi "different lanes at once" idea 6.2.08-Memory-coalescing hai.
Column-walk conflict 6.3.01-Matrix-transpose-optimization ka star villain hai.
Kitna shared memory afford kar sakte ho yeh 6.2.10-Occupancy-and-resource-limits se juda hai.
Threads ka lockstep mein kaam karna 6.2.11-Warp-divergence ke neeche bhi hai.
Parent pe wapas jaao: is topic ka Hinglish version .
Daayein taraf dhako aur reveal karne se pehle out loud jawab do.
Ek thread hai ek worker jo apne data pe same code run karta hai, threadIdx.x se identify hota hai.
Ek warp hai 32 threads ka ek bundle jo lockstep mein execute karta hai aur apne memory requests simultaneously fire karta hai.
Ek byte-address vs ek word byte ek numbered slot hai; ek 4-byte word chaar bytes ek float/int mein jode hue hain.
Byte-address ko 4 se kyun divide karte hain ise word-index mein convert karne ke liye, kyunki banks 4 bytes wide hain aur sirf words ki care karte hain.
a mod n a n 0 … n − 1
Master bank formula bank_index = word_index mod 32
Stride s thread i i + 1 i ( i ⋅ s ) mod 32
g cd( a , b ) woh sabse bada whole number jo a b g cd( s , 32 )
Coprime strides kyun safe hain agar g cd( s , 32 ) = 1 i ↦ ( i ⋅ s ) mod 32
matrix[r][c] ka word address (row-major, C r ⋅ C + c C
C 33 mod 32 = 1
