Coalesced memory access: Foundations — Hinglish deep dive

6.2.8 · D1 · Hardware › GPU Architecture › Coalesced memory access

Is page par koi assumption nahi hai. Isse pehle ki tum samjho ki ek memory pattern doosre se 32× faster kyun hota hai, tumhe pata hona chahiye ki thread, warp, address, byte, transaction, aur stride actually kya hain — aur har ek ke peeche ka picture dekhna chahiye. Hum inhe order mein build karte hain, har ek pehle wale par lean karta hua.

1. Byte — woh unit jisme sab kuch measure hota hai

Ek row of lockers ki imagine karo. Har locker ek byte hai. Locker ke darwaaze par likha number row mein uska position hai — locker 0, locker 1, locker 2, aur aise hi aage.

Hume yeh kyun chahiye: is topic mein har quantity — float kitna bada hai, transaction kitna bada hai, thread kahan read karta hai — bytes mein count hoti hai. Yeh ruler hai.

2. Address — locker ke darwaaze par likha number

Tum addresses hexadecimal mein likhte dekkhoge, jaise 0x0004. Yeh bas ek normal number likhne ka aur tarika hai:

Hex	Plain number
`0x0000`	0
`0x0004`	4
`0x0020`	32
`0x0080`	128
`0x1000`	4096

Hex kyun? Kyunki memory sizes powers of two hoti hain (32, 128, 4096…), aur woh base 16 mein round, easy-to-spot numbers ke roop mein aati hain. 0x1000 = 4096 neat hai; "4096" base 10 mein roundness chhupa deta hai.

Question — 0x0020 plain number mein kya hai? ::: 32 Question — GPU engineers addresses hex mein kyun likhte hain? ::: Memory sizes powers of two hoti hain, jo base 16 mein round aur align karne mein aasaan lagti hain.

3. Element aur uska size $E$ — kai bytes mein spread ek value

Ek byte sirf 0–255 hold karta hai, jo bahut chhota hai. Ek real decimal number (ek float32) store karne ke liye computer 4 bytes ko glue karke ek element banata hai.

Toh data[i] ka address hai:

$address of data[i] = base + i \times E$

jahaan base pehle element data[0] ka address hai.

Hume $E$ kyun chahiye: ek array index (data[7]) ko ek address (byte 28) mein convert karne ke liye, jo woh language hai jo memory hardware actually bolti hai.

4. Thread — code run karne wala ek worker

32 identical workers imagine karo, har ek ke haath mein apna number likha slip hai: 0, 1, 2, …, 31.

Hume yeh kyun chahiye: coalescing poori tarah is baat ke baare mein hai ki kaunse threads ek saath kaunse addresses read karte hain. Threads nahi, story nahi. (Yeh page ek thread ko simply "ek worker with an ID" ki tarah treat karta hai; yahaan bas itna hi chahiye.)

5. Warp $W$ — 32 threads jo perfect lock-step mein move karte hain

Yeh poore topic ka dil hai.

Woh consequence jo coalescing ko possible banati hai: jab shared instruction "memory se read karo" hoti hai, toh sare 32 threads simultaneously memory hit karte hain. Hardware 32 addresses ek saath aate dekhta hai aur inhe cleverly serve karne ka decide kar sakta hai — ek batch ki tarah, 32 alag errands ki tarah nahi.

Question — SIMT kya force karta hai sabhi 32 threads ko karne ke liye? ::: Ek hi instruction ek hi time par execute karne ke liye, taaki unke memory requests ek saath aayein.

6. Transaction size $T$ (a.k.a. "sector") — memory sirf fixed-size bundles mein bechti hai

Memory hardware kabhi tumhe ek single byte nahi deta. Woh hamesha ek fixed-size burst deliver karta hai.

Aligned ka matlab hai ki burst hamesha ek address par start hota hai jo $T$ ka whole multiple ho. $T = 32$ ke saath: transaction 0 bytes 0x00–0x1F cover karta hai, transaction 1 0x20–0x3F cover karta hai, aur aise hi. Tum ek burst nahi paa sakte jo beech se start ho.

$T$ kyun exist karta hai: DRAM chips ek baar mein poori rows read karte hain (dekho 6.2.09-Memory-hierarchy). Ek byte grab karna almost utna hi cost karta hai jitna ek full burst grab karna, isliye controller bursts move karne ke liye bana hota hai. Yeh ek single fact hai jiski wajah se coalescing matter karta hai.

Base-alignment edge case

Kyunki transactions sirf $T$ ke multiples par start hote hain, array khud kahaan se shuru hota hai yeh matter karta hai.

Picture yeh hai: bytes ki ek perfectly packed run phir bhi ek extra burst mein spill kar sakti hai agar woh transaction boundary par start nahi karti. Isliye libraries 128 ya 256 bytes par aligned memory return karti hain — guarantee karne ke liye ki base ek boundary par baitha ho aur minimum transaction count achievable ho.

Question — ek perfectly sequential warp phir bhi 4 ki jagah 5 transactions kyun cost kar sakta hai? ::: Agar array base T ka multiple nahi hai, toh 128 useful bytes ek extra transaction boundary straddle karte hain.

7. Stride — ek thread ke address aur agley ke beech ka gap

Stride woh knob hai jo sab kuch decide karta hai:

Chhota stride (= element size): consecutive threads neighbouring lockers read karte hain, isliye unme se bahut saare ek hi transaction ke andar land karte hain. → coalesced, kam bursts.
Bada stride (≥ transaction size): har thread ek alag transaction mein land karta hai. → non-coalesced, 32 bursts tak.

Har case, spell out kiya gaya (assume base $T$ par aligned, $T = 32$ , $E = 4$ , toh 8 floats per transaction):

Stride	Kya hota hai	Warp ke liye Transactions
$= E$ (4 B)	8 floats fit per transaction; 32 floats 128 B span karte hain	4 (minimum)
$= 2 E$ (8 B)	threads 2 floats apart step karte hain; 4 threads per transaction; warp 256 B span karta hai	8
$= 4 E$ (16 B)	2 threads per transaction; warp 512 B span karta hai	16
$= T$ (32 B)	exactly 1 thread per transaction; har ek fresh one start karta hai	32 (worst case)
$> T$ (e.g. 4096 B)	threads aur bhi door, phir bhi ek transaction each	32 (worst case)
stride 0 (sab same address read karte hain)	ek transaction sab ko serve karta hai (broadcast)	1

Intermediate strides ka pattern ( $E \leq stride \leq T$ , base aligned): threads jo ek transaction share karte hain unki number $⌈ T / stride ⌉$ hai — tum round up karte ho, kyunki ek thread jiska element ek transaction ke andar partially bhi fall karta hai woh still udhar land karta hua count hota hai. Equivalently, threads-per-transaction $= ⌊ (T - 1) / stride ⌋ + 1$ . Toh transactions $\approx W \div ⌈ T / stride ⌉$ . Jaise stride $E$ se $T$ ki taraf badhta hai, transactions steadily 4 se 32 tak chadhte hain — coalescing gradually degrade hoti hai, yeh all-or-nothing cliff nahi hai.

Note karo ki worst case 32 par saturate hota hai: tumhe kabhi ek thread se zyada ek transaction nahi chahiye, kyunki sirf 32 threads hain.

Question — float32 ke liye perfect coalescing kaunsa stride deta hai? ::: Element size ke barabar stride, 4 bytes (har thread next element read karta hai). Question — 8 bytes (2 floats) stride ke saath, warp ko kitne transactions chahiye? ::: 8 — char threads har 32-byte transaction share karte hain, toh 32÷4 = 8. Question — stride s ke liye ek transaction kitne threads share karte hain (T-byte transactions)? ::: ceil(T/s) — round up karo, kyunki partially-overlapping element bhi count hota hai. Question — ek warp ko kabhi 32 se zyada transactions kyun nahi chahiye? ::: Sirf 32 threads hain; worst case ek transaction each hai.

8. Coalesced vs non-coalesced — definition, plainly stated

Ab ki transactions aur strides exist karte hain, hum woh term define kar sakte hain jiske naam par poora topic hai.

Note karo yeh minimum ke relative hai, koi fixed number nahi: 32 consecutive floats coalesced hain chahe woh 1 transaction cost kare (128-byte-line GPU) ya 4 transactions (32-byte-sector GPU) — dono cases mein yeh 128 useful bytes ke liye kam se kam possible bursts hain.

Question — coalesced access ko ek line mein define karo. ::: Warp ke addresses minimum possible number of transactions se serve hote hain, toh almost har fetched byte use hota hai. Question — kya "coalesced" ek fixed transaction count hai? ::: Nahi — iska matlab us architecture ke liye minimum hai; count ek 128-byte-line aur ek 32-byte-sector GPU ke beech differ karta hai.

9. Bandwidth $B$ aur Latency $L$ — memory trip ke do costs

Jab tum memory touch karte ho toh do alag clocks tick karte hain. Inhe confuse karna beginner ki classic galti hai.

Yeh exactly woh tension hai jo 6.3.02-Roofline-model capture karta hai.

10. $N_{transactions}$ aur Efficiency — scorecard

Pehle, woh counter naam karo jo hum poora time use karte rahe hain:

Ab hum kisi bhi pattern ko grade kar sakte hain. Fraction ko piece by piece build karo:

Actually used bytes — $W$ threads mein se har ek genuinely $E$ bytes ka ek element chahta hai, aur kuch nahi. Toh useful work $= W \times E$ . $W = 32$ , $E = 4$ ke saath woh $128$ bytes hai.
Actually fetched bytes — hardware $N_{transactions}$ poore bursts pull karta hai, har ek $T$ bytes, chahe har byte wanted ho ya na ho. Toh fetched $= N_{transactions} \times T$ .

Us fraction mein har symbol upar build kiya gaya tha: $W$ (§5), $E$ (§3), $T$ (§6), $N_{transactions}$ (§10). Yahaan kuch bhi naya nahi hai — tum ab poora formula own karte ho.

Question — "bytes used" W·E kyun hai? ::: W threads mein se har ek exactly ek E-byte element chahta hai, aur kuch nahi. Question — "bytes fetched" N_transactions·T kyun hai? ::: Hardware hamesha poore T-byte bursts pull karta hai, ek per transaction, wanted ho ya na ho. Question — 32-byte-transaction GPU par stride-32-byte access ki efficiency kya hai? ::: 12.5% — 1024 fetched mein se 128 useful bytes.

Pieces topic ko kaise feed karte hain

Neeche ka map is page ka exact build order trace karta hai: bytes hume addresses dete hain; addresses aur element size milkar har data[i] locate karte hain; threads ek warp mein group hokar sab ek saath read karte hain; unke addresses ke beech ka stride aur transaction size milkar $N_{transactions}$ decide karte hain; woh count define karta hai ki access coalesced hai ya nahi, aur latency/bandwidth ke saath yeh time cost deta hai aur, used-vs-fetched bytes ke saath, efficiency deta hai. Har arrow ko "isse pehle samajhna zaruri hai" ki tarah padho.

Vault mein related prerequisites: 6.2.04-Shared-memory-and-synchronization (ek alternative jab coalescing impossible ho), 6.2.09-Memory-hierarchy (jahaan transactions rehte hain), aur 5.1.05-Cache-coherence-protocols (cache-line ideas jo yahaan echo karte hain).

Equipment checklist

Right side cover karo aur out loud jawab do. Agar koi line surprise kare, toh parent note tackle karne se pehle us section ko dobara padho.

Byte hai... ::: ek memory box jo 0–255 number hold karta hai; memory inki ek numbered row hai. Address hai... ::: woh number jo identify karta hai ki tum kaunsa byte mean kar rahe ho. 0x0020 plain digits mein hai... ::: 32. Element hai... ::: ek logical data value jo program ek single number ki tarah treat karta hai, E bytes se bana. Float32 ke liye element size E hai... ::: 4 bytes; data[i] base + i×4 par start hota hai. Thread hai... ::: apna ID, threadIdx.x, wala ek independent worker. Warp (size W) hai... ::: 32 threads jo same instruction ek hi instant par run karne ke liye force hote hain (SIMT). Transaction / sector (size T) hai... ::: sabse chhota fixed, aligned burst jo memory controller fetch karta hai (T=32 B Volta-style par); dono words yahaan ek hi cheez mean karte hain. "Aligned" ka matlab hai... ::: burst T ke multiple par start hota hai; mid-burst se shuru hone wala burst nahi mil sakta. Ek misaligned array base kar sakta hai... ::: ek packed warp ko ek extra transaction mein push karna, kyunki useful bytes ek boundary straddle karte hain. Stride hai... ::: consecutive threads ki reads ke beech fixed address gap. Perfect coalescing ke liye stride hai... ::: element size ke barabar (4 B), taaki threads neighbouring lockers read karein. Stride s ke liye ek transaction share karne wale threads hain... ::: ceil(T/s) — round up, floor nahi. Coalesced access hai... ::: ek warp minimum possible number of transactions se serve hota hai, toh nearly har fetched byte use hota hai. Non-coalesced access hai... ::: scattered addresses jinhe minimum se zyada transactions chahiye, zyaadatar fetched bytes waste hote hain. N_transactions hai... ::: T-byte bursts ki kitni kitni number hardware ek warp ke read ke liye issue karta hai. Latency L hai... ::: pehla byte aane se pehle ka wait (ek representative ~300 ns; real GPUs 400–800 ns+), ek baar per transaction pay hota hai. Bandwidth B hai... ::: bytes flow shuru hone ke baad kitni tezi se stream karte hain (~900 GB/s); wait include nahi hai. Non-coalescing slow mainly isliye hai kyunki... ::: tum latency L 32 baar tak pay karte ho, ek baar har scattered transaction ke liye. Efficiency equal hai... ::: bytes used over bytes fetched = W·E / (N_transactions·T).