6.2.8 · D5 · HinglishGPU Architecture
Question bank — Coalesced memory access
6.2.8 · D5· Hardware › GPU Architecture › Coalesced memory access
Shuru karne se pehle, teen words jinhe hum poore page par use karenge — yeh sab parent note mein build hue the, lekin chaliye unhe re-anchor karte hain taaki kuch bhi use hone se pehle clear ho. (Is vault mein hum key terms ko ==double equals== mein wrap karte hain, jo Obsidian highlight ki tarah render karta hai aur flashcard cloze bana deta hai — har highlighted phrase ko "exactly yahi wording yaad karo" samjho.)
True or false — justify karo
True or false: "Coalesced" ka matlab hai ki saare 32 threads exactly ek memory transaction se read karte hain.
False. Ek 32-byte-sector GPU par, 32 floats (128 bytes) chaar aligned sectors mein span karte hain, phir bhi yeh perfectly coalesced hai. Coalescing ka matlab hai minimum transactions with no waste, na ki ek transaction.
True or false: Agar threads consecutive addresses read karte hain, toh access hamesha coalesced hoti hai chahe starting address koi bhi ho.
False. Alignment matter karta hai — agar 128-byte span kisi transaction ke beech mein start ho, toh woh ek extra sector ko straddle karti hai, toh tum ideal se ek zyaada transaction pay karte ho. Consecutive lekin misaligned hona, consecutive aur aligned hone se bura hai.
True or false: Har thread ke liye sirf 4 useful bytes read karna lekin 32-byte sectors fetch karna hamesha bandwidth waste karta hai.
Coalesced hone par False. Agar neighbouring threads usi sector ke baaki bytes consume karte hain, toh kuch bhi waste nahi hota — 8 threads ek 32-byte sector share karte hain aur saare 32 bytes use karte hain. Waste tab aata hai jab threads scatter ho jaate hain.
True or false: Exactly 2 ka stride (data[idx*2]) 1024 ke stride jitna bura hai.
False, lekin yeh phir bhi degraded hai. Stride-2 har sector ka aadha use karta hai, toh efficiency ~50% hai, na ki bade stride ka ~12.5%. Badness smoothly stride ke saath badhti hai jab tak har thread apne transaction mein land na kare.
True or false: Coalescing ek software optimisation hai jise tum compiler flag se turn on karte ho.
False. Yeh tumhare access pattern ki emergent property hai. Hardware automatically coalesce karta hai; tumhara kaam yeh hai ki indices aise likho taaki addresses contiguously land hon.
True or false: Ek modern sectored (Volta+) GPU par, non-coalesced access ab koi problem nahi hai.
False. Sectoring penalty ko soft karta hai (har stray thread ke liye 128-byte waste ki jagah 32-byte waste), lekin ek fully scattered warp phir bhi up to 32 transactions issue karta hai aur ~32× latency pay karta hai. Yeh better hai, solved nahi.
True or false: Agar ek kernel latency-bound hai, toh coalescing matter nahi karta.
Mostly false. Non-coalescing extra latency-bound transactions create karta hai (up to 32 separate round trips), toh poor coalescing often wahi kaaran hai jis wajah se tum pehli jagah latency-bound ho. Dekho 6.3.02-Roofline-model.
True or false: Coalescing aur cache hits same cheez hain.
False. Coalescing ek warp ke addresses ke baare mein hai jo abhi fewest transactions mein fit ho rahe hain. Caching time ke saath reuse ke baare mein hai. Ek coalesced load phir bhi cache miss kar sakta hai; ek cache hit phir bhi scattered pattern se aa sakti hai.
True or false: Do threads exactly same address read karna coalescing tod deta hai.
False. Identical addresses same transaction mein fall karte hain, toh hardware dono ko ek fetch se serve karta hai (ek broadcast). Yeh stride se sasta hai, bura nahi.
Error dhundho
"data[idx * 4] coalesced hai kyunki floats 4 bytes ke hote hain."
Error: index ko 4 se multiply karne par 4 floats = 16 bytes ka stride threads ke beech milta hai, har sector ke 3 mein se 4 floats unread reh jaate hain. 4-byte float size already indexing mein baked in hai; tumhe use dobara multiply nahi karna chahiye.
"Row-major matrix access ko coalesced banane ke liye, thread i ko M[i][0] read karna chahiye."
Error:
M[i][0] row-major storage mein ek column ke neeche step karta hai, toh consecutive threads ek poore row ke distance par hain — ek bada stride. Coalesced row-major access mein thread i, M[row][i] read karta hai, row ke saath chalta hai."Transaction size 128 bytes hai, toh coalescing ka matlab hai har warp exactly 1 transaction karta hai."
Error: yeh ek representative number ko universal law ke saath confuse karta hai. Transaction size architecture-dependent hai (128-byte line vs 32-byte sector), aur 128 bytes par bhi 32 floats ka ek warp exactly ek line hai sirf tab jab aligned ho. Dekho 6.2.09-Memory-hierarchy.
"Mera access strided hai lekin maine ek __shared__ buffer add kiya, toh ab yeh coalesced hai."
Error: shared memory global access ko coalesced nahi banata — tum use reorganise karne ke liye use karte ho. Tum global load ko contiguous addresses mein coalesce karte ho, use shared memory mein stage karte ho, phir threads ko scattered pattern shared memory se read karne dete ho (jiske alag rules hain — coalescing nahi, bank conflicts).
"Efficiency = bytes used / bytes fetched, aur mera kernel 1024 mein se 128 bytes use karta hai, toh efficiency 12.5% par theek hai."
Error: 12.5% terrible hai, theek nahi — iska matlab hai tumhara 87.5% memory bandwidth throw away ho raha hai. Formula sahi hai; interpretation good aur bad ko invert kar deta hai.
"Threads data[idx] aur data[idx+1] read karte hain to overlap hota hai, toh yeh double-counted aur slow hai."
Error: same sector ke andar overlapping reads free hain — sector already fetch ho chuka hai. Transaction ke andar reuse good case hai, slowdown nahi.
Why questions
Hardware poora 32-byte sector kyun fetch karta hai jab ek thread sirf 4 bytes maangta hai?
Kyunki DRAM physically isi tarah kaam karta hai: ek read pehle chip ki ek poori row ko row buffer mein activate karta hai (slow step), phir us buffer se consecutive columns ko burst mein stream karta hai. Costly part — row kholna — already paid ho chuka hota hai, toh 32 contiguous bytes pull karna 4 pull karne se almost kuch extra nahi costs. Isliye controller hamesha ek poora sector-sized burst transfer karta hai na ki ek single element; 4 bytes maangne par bhi row khulti hai aur poora burst stream hota hai.
Ek bada stride 32 transactions tak kyun cause karta hai, sirf kuch kyun nahi?
Jab stride transaction size se zyaada ho jaata hai, toh koi bhi do threads ke addresses ek sector share nahi karte, toh 32 threads mein se har ek apna alag fetch force karta hai — worst case.
Misalignment (kuch bytes offset) ek extra transaction kyun cost karta hai?
Ek contiguous 128-byte block jo kisi sector ke beech mein start hota hai next sector mein spill ho jaata hai, toh ab woh ek aur aligned transaction touch karta hai jo us block se zyaada hai jो sector boundary par start ho.
Speedup often "up to 32×" kyun quote kiya jaata hai, koi fixed number kyun nahi?
32× theoretical ceiling hai (1 transaction vs 32). Real speedup stride, alignment, cache behaviour, aur iss baat par depend karta hai ki tum bandwidth- ya latency-bound ho, toh yeh ~1× se ~32× ke beech kahin bhi land hota hai.
Array-of-Structs (AoS) layout Struct-of-Arrays (SoA) ke comparison mein coalescing kyun hurt karta hai?
AoS mein, thread
i struct i ka field.x read karta hai toh poore struct size se step karta hai, x values ko many sectors mein scatter karta hai. SoA saare x values contiguously store karta hai, toh same access ek clean stride-1 read hai.Coalescing kyun matter kar sakta hai jab do kernels mein total data moved identical ho?
Kyunki scattered access actually fetched bytes ko inflate karta hai (har 4 bytes use karne ke liye full sectors fetch karna). Same useful data, lekin kahin zyaada fetched data aur zyaada latency-bound transactions.
Ek purely compute-bound kernel mein coalescing fix karna kyun help nahi karta?
Agar arithmetic, memory nahi, bottleneck hai, toh memory subsystem ke paas already slack hai. Coalescing improve karna woh bandwidth free karta hai jo limiter nahi tha — roofline batata hai tum kis side par ho.
Edge cases
Edge case: Ek warp jahan saare 32 threads same single address read karte hain.
Fully coalesced (actually, ek broadcast ke saath ek transaction). Zero stride sabse friendly pattern hai, koi degenerate failure nahi.
Edge case: Sirf 8 of 32 threads (ek partial/predicated warp) ke saath active ek warp jo contiguous floats read karta hai.
Phir bhi coalesced — woh 8 addresses (32 bytes) ek sector mein fit hain. Inactive threads koi request issue nahi karte, toh woh na help karte hain na transaction count hurt karte hain.
Edge case: Stride exactly transaction size ke barabar (idx * 8 floats = 32 bytes).
Worst case — har thread ek distinct sector ke start par land karta hai, toh 32 sectors fetch hote hain aur har 32 bytes mein se sirf 4 use hote hain (≈12.5% efficiency).
Edge case: float ki jagah char (1 byte) contiguously read karna.
32 threads 32 bytes read karte hain = exactly ek 32-byte sector, phir bhi fully coalesced. Chhote elements zyaada threads ko ek sector mein pack karte hain, jo theek hai jab tak pattern contiguous rahe.
Edge case: Ek base pointer 8 bytes se misaligned lekin otherwise stride-1 access.
128-byte span ab chaar 32-byte sectors ki jagah paanch ko straddle karta hai, ek extra transaction costing. Allocations ko align karna (e.g.
cudaMalloc aligned memory return karta hai) ideal restore karta hai.Edge case: Array ek warp ke floats se chhota hai (sirf 5 elements).
5 in-bounds threads ek sector coalesced read karte hain; out-of-bounds threads ko mask off karna padega (bounds check) taaki woh koi illegal transaction issue na karein — warna tumhe coalescing issue nahi, out-of-bounds fault milta hai.
Recall Quick self-test
Kaun sa pattern coalesced hai: data[idx], data[idx*2], ya data[idx*1024]? ::: Sirf data[idx] — stride-1 contiguous. Baaki har fetched sector ka aadha, phir almost sab, waste karte hain.
Kya ek 32-byte-sector GPU scattered access ko acceptable banata hai? ::: Nahi — yeh har stray thread ke liye waste ko 32 bytes tak cap karta hai, lekin ek fully scattered warp phir bhi up to 32 latency-heavy transactions issue karta hai.
Kaun si ek property "coalesced" define karti hai? ::: Warp ke 32 addresses ke liye fewest possible transactions with near-zero wasted bytes — koi fixed transaction count nahi.
Related: Coalesced memory access · 6.2.09-Memory-hierarchy · 5.1.05-Cache-coherence-protocols · 6.3.02-Roofline-model