Bank conflicts in shared memory — Hinglish notes

6.2.9 · Hardware › GPU Architecture

Ye kyun matter karta hai: Shared memory fast honi chahiye (~100× global memory se faster). Bank conflicts is advantage ko barbad kar dete hain parallelism ko serialize karke.

How shared memory banks work

First principles se derivation:

4 se divide kyun? Har bank 4 bytes wide hota hai (NVIDIA hardware par fixed bank width). Ek byte address ko word index mein convert karna padta hai. Address 0-3 → word 0, address 4-7 → word 1, etc.
Mod 32 kyun? 32 banks ke saath, hum words ko round-robin fashion mein distribute karte hain. Word 0 → bank 0, word 1 → bank 1, .., word 31 → bank 31, word 32 → bank 0 phir se.
Result: Words banks mein striped hote hain. Consecutive 32-bit words consecutive banks mein rehte hain.

Figure — Bank conflicts in shared memory

Types of access patterns

Ye kyun kaam karta hai:

Thread 0 reads shared[0] → word 0 → bank 0
Thread 1 reads shared[1] → word 1 → bank 1
...
Thread 31 reads shared[31] → word 31 → bank 31

Saare 32 threads alag-alag banks hit karte hain → 1 cycle access, perfect parallelism.

Analysis:

Thread 0: shared[0] → bank 0
Thread 1: shared[2] → bank 2
Thread 2: shared[4] → bank 4
...
Thread 16: shared[32] → word 32 → bank $(32 mod 32) = 0$ ← thread 0 se conflict!
Thread 17: shared[34] → bank 2 ← thread 1 se conflict!

Result: 16 pairs of threads collide (2-way conflicts). Access 2 cycles leta hai 1 ki jagah.

Ye step kyun? Stride-2 ke saath, har 16 threads baad hum bank 0 par wapas aa jaate hain. 16 positions ke antar par wale threads same bank request karte hain.

Thread 0: shared[0] → bank 0
Thread 1: shared[32] → word 32 → bank 0
Thread 2: shared[64] → word 64 → bank 0
Saare 32 threads bank 0 hit karte hain!

Access time: 32 cycles (fully serialized). Shared memory ab registers se bhi slow hai.

The mathematical pattern

Derivation:

Thread $i$ word index $i \cdot s$ access karta hai.
Thread $i$ ke liye bank: $(i \cdot s) mod 32$ .
Do threads $i$ aur $j$ conflict karte hain agar $(i \cdot s) mod 32 = (j \cdot s) mod 32$ .
Yeh simplify hota hai $(i - j) \cdot s \equiv 0 (mod 32)$ mein.
Sabse chhota $k = i - j$ jo yeh satisfy kare woh hai $k = 32/ g cd (s, 32)$ .
Toh har $32/ g cd (s, 32)$ threads par, banks repeat hote hain, matlab $g cd (s, 32)$ threads har bank par collide karte hain.

Examples:

Stride 1: $g cd (1, 32) = 1$ → koi conflict nahi
Stride 2: $g cd (2, 32) = 2$ → 2-way conflict
Stride 3: $g cd (3, 32) = 1$ → koi conflict nahi (odd strides safe hain!)
Stride 16: $g cd (16, 32) = 16$ → 16-way conflict
Stride 32: $g cd (32, 32) = 32$ → 32-way conflict

Common mistake: structure-of-arrays vs array-of-structures

Common mistake "Array-of-structures hamesha conflicts cause karta hai"

Galat idea:

struct Particle {
    float x, y, z;  // 12 bytes = 3 words
};
__shared__ Particle particles[32];
float px = particles[threadIdx.x].x;  // Accessing the .x field

Pehli nazar mein galat kyun lagta hai: "Threads 12 bytes apart hain, toh surely banks par collide karenge."

Sach mein kya hai: Chalte hain actually .x access ke liye banks compute karte hain.

Adjacent .x fields ke beech word stride 3 words hai (12 bytes / 4).
Thread $i$ word $3 i$ access karta hai, bank $(3 i) mod 32$ par land karta hai.
Kyunki $g cd (3, 32) = 1$ , map $i \mapsto (3 i) mod 32$ 32 banks ki ek permutation hai.
Thread 0→bank 0, T1→bank 3, T2→bank 6, ..., T10→bank 30, T11→bank 33 mod 32 = bank 1, ...

Toh ek warp ka .x (ya .y, ya .z) read actually CONFLICT-FREE hai, kyunki 3 aur 32 coprime hain. Yeh intuition ki "non-unit stride = conflict" galat hai; sirf woh strides jo 32 ke saath koi factor share karein (yaani even strides / powers of 2) conflicts cause karte hain.

Jab AoS sach mein problem karti hai: Conflicts tab aate hain jab struct size 32 ke saath koi factor share kare — jaise 4 floats ka struct (16 bytes = 4 words). Tab stride = 4, $g cd (4, 32) = 4$ → har field par 4-way conflict. AoS ki real danger yahi hai: aap stride control nahi karte, aur ek "achhi" power-of-two struct size silently conflicts create kar deti hai.

Robust fix: Structure-of-arrays

__shared__ float x[32], y[32], z[32];
float px = x[threadIdx.x];  // Stride-1, guaranteed conflict-free
float py = y[threadIdx.x];
float pz = z[threadIdx.x];

SoA har field access ko stride-1 banata hai, toh aap kabhi is par depend nahi karte ki struct size 32 ke saath coprime ho.

Padding to avoid conflicts

Case A — same row, alag columns (yeh conflict-free hai):

float value = matrix[row][tid];   // fixed row, tid = column

Thread $i$ word $row \cdot 32 + i$ access karta hai. Jab $i$ vary karta hai, stride 1 word hai.
Banks: $(row \cdot 32 + i) mod 32 = i$ (kyunki $32 \equiv 0$ ). Saare alag → koi conflict nahi.

Case B — same column, alag rows (yeh wala bura hai):

float value = matrix[tid][col];   // tid = row, fixed column

Thread $i$ word $i \cdot 32 + col$ access karta hai. Stride 32 words hai.
Banks: $(i \cdot 32 + col) mod 32 = col mod 32$ , har thread ke liye same → 32-way conflict!

Yeh column-walk pattern exactly wahi hai jo matrix transpose mein hota hai, isliye yeh bahut important hai.

Fix: inner dimension ko 33 tak pad karo

__shared__ float matrix[32][33];  // ek extra padding column
float value = matrix[tid][col];   // column access

Ab element $[r] [c]$ word $r \cdot 33 + c$ par hai.
Thread $i$ (row $i$ , fixed column) word $i \cdot 33 + col$ access karta hai.
Bank index (extra /4 nahi — hum already word units mein hain): $bank = (i \cdot 33 + col) mod 32 = (i + col) mod 32$ kyunki $33 \equiv 1 (mod 32)$ .
Jab $i$ 0..31 tak run karta hai, $(i + col) mod 32$ saare 32 distinct values leta hai → conflict-free!

Yeh step kyun? 33 tak padding karne se row stride $\equiv 1 (mod 32)$ ho jata hai, toh consecutive rows exactly ek bank shift hoti hain same bank par land karne ki jagah. Yahi poora trick hai.

Recall Ek 12-saal ke bacche ko explain karo

Socho tumhare school mein 32 cubbyholes hain jahan students apne lunch boxes rakhte hain. Rule yeh hai: student 1 cubby 1 use karta hai, student 2 cubby 2, aur student 32 tak cubby 32. Phir repeat hota hai: student 33 phir cubby 1 use karta hai.

Ab, 32 dost (ek "warp") ek saath apna lunch lene ki koshish karte hain. Agar unhone cubbies 1, 2, 3, .., 32 mein lunch rakha tha, toh sabko ek saath mil jaata hai — wait nahi karni! Yahi conflict-free hai.

Lekin kya agar unhone sab ne cubby 1 mein lunch rakha? Ab teacher ko ek ek karke dena padega — super slow! Yahi 32-way bank conflict hai.

Bank conflicts tab hote hain jab multiple threads (students) ek saath same bank (cubby) se data chahein. GPU ko unhe ek ek karke serve karna padta hai sab ek saath ki jagah. Trick yeh hai ki apna data aise organize karo ki har thread alag bank use kare — jaise ensure karo ki har student ka cubby number alag ho. Har row mein ek "khali spare cubby" add karna (padding to 33) sabko ek shift kar deta hai taaki woh same cubby par clash karna band kar dein.

Connections

Shared memory architecture—woh physical design jo banks ko enable karta hai
Memory coalescing—global memory ke liye similar parallelism concept
Occupancy—bank conflicts effective occupancy reduce karte hain warps stall karke
Matrix transpose—classic example jahan padding conflicts fix karti hai
Warp divergence—GPU execution mein ek aur serialization bottleneck

#flashcards/hardware

Shared memory banks kya hain aur ye kyun exist karte hain? :: Banks shared memory ke parallel sub-units hain (typically 32) jo multiple threads ko alag-alag addresses simultaneously access karne dete hain. Ye high-bandwidth parallel access enable karne ke liye exist karte hain — har bank ek cycle mein ek request serve kar sakta hai.

Bank conflict kya hota hai? :: Jab ek warp ke multiple threads alag-alag addresses access karein jo same bank pe map hote hain, toh parallel access ki jagah serialized access force hoti hai. Yeh latency ko conflict degree se multiply kar deta hai.

4-byte words ke liye bank index kaise calculate hota hai?

bank_index = (byte_address / 4) mod 32. Byte address ko word index mein convert karne ke liye 4 se divide karo, phir 32 banks mein se konsa bank us word ko hold karta hai ye pata karne ke liye mod 32 karo.

Stride-2 access ke liye conflict degree kya hai? :: 2-way conflict. gcd(2, 32) = 2, toh har do threads same bank par collide karte hain, access time double ho jata hai.

Stride-1 conflicts kyun avoid karta hai?

Consecutive threads consecutive words access karte hain, jo consecutive banks (0, 1, 2, ..., 31) pe map hote hain. Saare 32 threads alag-alag banks hit karte hain parallel mein.

Stride s (words mein) diya ho toh conflict degree ka formula kya hai?

conflict_degree = gcd(s, 32). Ye wo threads ki sankhya hai jo har bank par collide karenge.

Kya ek 12-byte struct (stride 3) bank conflicts cause karta hai jab ek thread ek field read kare?

Nahi. gcd(3,32)=1, toh field accesses 32 banks ki ek permutation banate hain — conflict-free. Odd/coprime strides safe hain; sirf woh strides jo 32 ke saath factor share karein (even/powers of 2) conflict karte hain.

Shared matrix[32][32] ke liye kaunsa bura hai: matrix[row][tid] ya matrix[tid][col]?

matrix[tid][col] (same column, alag rows) bura hai — stride 32, saare threads same bank hit karte hain → 32-way conflict. matrix[row][tid] (same row) stride-1 hai aur conflict-free hai.

[32][33] tak padding column access conflicts kaise fix karti hai?

Row stride 33 words ban jata hai. Kyunki 33 ≡ 1 (mod 32), bank = (i*33 + col) mod 32 = (i + col) mod 32, jo saare 32 distinct banks deta hai → conflict-free.

Kya 64-bit mode mein doubles ke liye bank_index /8 se recompute karte hain?

Nahi. Banks fixed 4 bytes wide hote hain. Ek 8-byte double simply do consecutive banks mein span karta hai; hardware phir bhi stride-1 warp of doubles ko conflict-free handle karta hai. Hamesha 4-byte banks mein reason karo.

Stride-32 access ke saath kya hota hai?

32-way conflict. Saare threads aise addresses access karte hain jo same bank pe map hote hain, access ko completely serialize kar dete hain (32 cycles 1 ki jagah).

Shared memory ke liye optimal access pattern kaunsa hai?

Stride-1 access jahan consecutive threads consecutive 4-byte elements access karein. Yeh threads ko alag-alag banks pe perfect parallelism ke saath map karta hai.