Coalesced memory access: Worked examples — Hinglish deep dive

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

Yeh deep-dive parent topic ke andar hai. Wahan humne seekha ki coalescing kya hoti hai. Yahan hum har tarah ke access patterns jo ek warp produce kar sakta hai — perfect, offset, strided, mixed, degenerate, aur ek real word problem — ke through grind karte hain, taaki wild mein koi bhi access pattern mile toh tumne already use paper par solve kar rakha ho.

Is poore page ke ground rules

Taaki har number checkable ho, hum ek fixed model architecture use karte hain saare examples ke liye (bilkul parent ka Volta-style model):

Definition Hamara fixed model (neeche har example mein use hoga)

Warp size $W = 32$ threads.
Element size $E = 4$ bytes (ek float).
Transaction / sector size $T = 32$ bytes. Toh ek transaction $T / E = 8$ floats le jaata hai.
Ek transaction aligned hoti hai: sector $k$ byte addresses $[32 k, 32 k + 31]$ cover karta hai. Address $A$ sector $⌊ A /32 ⌋$ mein rehta hai.
Useful bytes = $W \times E = 32 \times 4 = 128$ bytes — har thread ek $E$ -byte element request karta hai. Hum ise har distinct-data example mein numerator ke taur par use karte hain taaki efficiency numbers comparable rahen. (Single degenerate broadcast case, Example 7, ko ek footnote chahiye — wahan flag kiya hai — kyunki 32 threads ka same element request karna 32 distinct useful bytes nahi hai.)
Efficiency $= \frac{bytes used}{bytes fetched} = \frac{W \cdot E}{N _{tx} \cdot T} = \frac{128}{32 N _{tx}}$ , jahan $N_{tx}$ distinct sectors ki sankhya hai jo touch hue.

Sab kuch reduce hota hai "32 addresses kitne distinct sectors mein land karte hain?" par. Us number ko $N_{tx}$ kahte hain. Wahi ek count transactions, bytes fetched, aur efficiency decide karta hai.

Intuition "Transaction fetched" se yahan kya assume kiya hai — aur caching footnote

Neeche har example mein har warp load ko cold miss maana gaya hai: sector already cache mein nahi baitha, toh hardware ko actually transaction lower memory ko issue karni padti hai. Yeh ek fresh access pattern ke baare mein sochne ka sahi worst-case model hai. Real kernel mein, inme se kaafi transactions L1 ya L2 hit karti hain (cache lines 128 bytes hoti hain, chaar 32-byte sectors se bani) aur kaafi kam cost karti hain — dekho 6.2.09-Memory-hierarchy aur line/sector origins 5.1.05-Cache-coherence-protocols mein. Caching ek kharab pattern ki cost reduce karta hai lekin use kabhi fix nahi karta: ek scattered warp phir bhi zyada evict aur re-fetch karta hai, toh hamare examples ki relative ordering (kaun sa pattern better hai) cache hits ke saath ya bina dono mein hold karti hai. Hum cold transactions count karte hain taaki woh ordering clearly dikhaye.

Address-ramp metaphor (hamara "regime map")

Matrix se pehle, ek picture har example ko jodti hai. Thread index $t$ ko horizontal axis par plot karo aur uska byte address $A (t)$ vertical axis par. Har access pattern bas dots ki ek line (ya scatter) hai — ek address ramp. Us ramp ka slope stride hai, aur slope akela decide karta hai tum kitne coalesced ho. Yeh hamara stand-in hai "kaun sa quadrant / kaun sa sign" ke liye jo tum angles ke saath karte ho: quadrants ki jagah hamare paas slope regimes hain, figure mein labelled bands ki tarah.

Chaar slope regimes exactly hamare case split hain (example IDs neeche matrix aur worked sections mein fully introduce kiye gaye hain):

Slope $= 0$ (flat): broadcast — Example 7.
Slope $= E$ (gentle, unbroken): coalesced — Examples 1, 6; Example 2 mein ek shift se toota hua.
Slope strictly $E$ aur $T$ ke beech (medium): partial waste — Examples 3, 9.
Slope $\geq T$ (steep): ek transaction per thread, floor — Examples 4, 5, aur Example 8 ka column-access half (Example 8a, neeche padho is bullet par rely karne se pehle).
Koi single slope nahi (scatter): mixed/irregular — Example 8b.

Numbers kahan se aate hain (T = 32 B arbitrary nahi hai)

Common mistake

$T = 32$ B hamara chosen model hai, real hardware se tied lekin universal nahi Jaise parent note ne stress kiya, transaction aur cache-line sizes GPU compute capability par depend karti hain. Hamara $T = 32$ bytes Volta/Turing/Ampere-class memory subsystem ka 32-byte sector hai: physically, L1/L2 cache line 128 bytes hai, lekin woh line chaar 32-byte sectors mein split hai jo memory system independently request kar sakta hai. Toh hamare model mein "ek transaction" = "ek 32-byte sector fill." Purane Fermi/Kepler/Pascal GPUs par ek global load through L1 poori 128-byte line pull karta hai ( $T = 32$ ki jagah $T = 128$ rakh do aur arithmetic redo karo — neeche har result ki shape identical hai, sirf transaction counts scale hoti hain). Apna actual GPU hamesha profile karo; efficiency ratio portable lesson hai, absolute counts nahi.

Scenario matrix

Har access pattern jo ek warp throw kar sakta hai inhi cells mein se ek mein aata hai, cell order A → I mein listed. Jo examples follow karte hain woh har cell cover karte hain. Aakhri do cells woh hain jo students underestimate karte hain: Cell I "exam twist" hai — ek struct { float x, y, z; } ka ek field padhna (ek array-of-structs, AoS), jahan ek hidden 12-byte stride coalescing kharab kar deta hai — jise hum Example 9 mein fully build karte hain. Cell J ek mixed warp hai bina kisi single stride ke, Example 8b mein build kiya gaya.

Cell	Pattern class	Isme kya khaas hai	Covered by
A	Perfectly sequential	stride = 1 element, aligned start	Example 1
B	Sequential lekin misaligned start	consecutive, non-multiple of $T$ se offset	Example 2
C	Small stride (2 elements)	partial waste, kuch sharing	Example 3
D	Stride = exactly ek sector	boundary case, zero sharing	Example 4
E	Huge stride (worst case)	ek transaction per thread	Example 5
F	Reversed / permuted order	same addresses, shuffled threads	Example 6
G	Degenerate: sab threads same address	broadcast, minimum possible	Example 7
H	Real-world word problem (matrix rows vs cols)	layout choice kyun matter karta hai	Example 8a
I	Exam twist: `struct` of 3 floats (AoS)	element size ≠ 4, hidden stride	Example 9
J	Mixed / irregular warp (koi single stride nahi)	aadha sequential, aadha scattered	Example 8b

Har row upar ke ramp picture ke ek slope regime se map hoti hai: increasing vs decreasing (Cell F), aligned vs shifted (Cell B), zero-slope (Cell G), steep limiting case (Cell E), aur no-single-slope scatter (Cell J).

Example 1 — Cell A: Perfect sequential (baseline)

Pehle forecast karo: Aage padhne se pehle sector count guess karo. (Hint: 32 floats = kitne bytes? 32 se divide karo.)

Address function likho. $A (t) = 4 t$ bytes (har float 4 bytes ka hai, stride 1). Yeh step kyun? Baaki sab $⌊ A (t) /32 ⌋$ hai; pehle $A (t)$ explicitly chahiye.
Har thread ka sector nikalo. $sector (t) = ⌊ 4 t /32 ⌋ = ⌊ t /8 ⌋$ . Yeh step kyun? Yeh 32 addresses ko 32 sector numbers mein convert karta hai — wahi ek cheez jo matter karti hai.
Distinct sectors count karo. Jab $t$ $0 \to 31$ run karta hai, $⌊ t /8 ⌋$ values $0, 1, 2, 3$ leta hai. Toh $N_{tx} = 4$ . Yeh step kyun? $N_{tx}$ hamara master number hai.
Efficiency. $\frac{128}{32 \times 4} = \frac{128}{128} = 1.0 = 100%$ . Yeh step kyun? Zero wasted bytes confirm karta hai — coalesced ki definition.

Figure walkthrough: Upar ke map mein, 32 coral bars (thread reads) exactly 4 mint sectors ke upar baithe hain, koi gaps nahi aur koi grey (unused) sectors nahi. "Coral perfectly tiles mint" — yahi 100% efficiency dikhti hai.

Verify: 4 sectors × 32 bytes = 128 bytes fetched = 128 bytes used. Addresses 0x00–0x7F = exactly 128 bytes, exactly 4 aligned sectors. Units: bytes/bytes = dimensionless efficiency. ✓

Example 2 — Cell B: Sequential lekin misaligned start

Forecast: Same data size (128 bytes) — lekin kya window ko sector boundaries ke across shift karne se ek extra transaction ki cost aayegi?

Address function. $A (t) = 8 + 4 t$ . Kyun? Offset 8, $T = 32$ ka multiple nahi hai, toh 128-byte window ek extra sector boundary straddle karti hai.
Sectors touched. Sabse chota address $A (0) = 8 \Rightarrow$ sector $⌊ 8/32 ⌋ = 0$ . Sabse bada $A (31) = 8 + 124 = 132 \Rightarrow$ sector $⌊ 132/32 ⌋ = 4$ . Kyun? Span ab sector 0 se sector 4 tak pahunchta hai.
Distinct sectors count karo. Sectors $0, 1, 2, 3, 4$ sab touch hote hain $\Rightarrow N_{tx} = 5$ . Kyun? 128-byte window 5th sector mein thodi ghus gayi; hume usse bhi fetch karna padega.
Efficiency. $\frac{128}{32 \times 5} = \frac{128}{160} = 0.8 = 80%$ . Kyun? Humne 160 bytes fetch kiye lekin sirf 128 use kiye — misalignment ne hume 20% cost kiya.

Figure walkthrough: Example 1 ke map se compare karo. Coral bars 8 bytes right shift ho gaye hain, toh pehle do bars ab sector 0 mein hain jabki aakhri bar ek fifth sector (sector 4, mostly unused space ke saath grey-mint drawn) mein spill karta hai. Woh single extra half-empty sector hi woh 20% hai jo tumne khoya.

Verify: span = 0x08–0x84 (132 = 0x84). $⌊ 132/32 ⌋ - ⌊ 8/32 ⌋ + 1 = 4 - 0 + 1 = 5$ sectors. Efficiency $128/160 = 0.8$ . ✓

Example 3 — Cell C: Small stride (2 elements)

Forecast: Hum ab har doosra float skip kar rahe hain. Hum 128 bytes use karte hain lekin woh double span mein spread hain. $N_{tx}$ guess karo.

Address function. $A (t) = 4 \times 2 t = 8 t$ . Kyun? Stride 2 floats = consecutive threads ke beech 8 bytes.
Sectors. $sector (t) = ⌊ 8 t /32 ⌋ = ⌊ t /4 ⌋$ , $t : 0 \to 31$ ke saath $0, 1, \dots, 7$ run karta hai. Kyun? Ab sirf 4 threads ek sector share karte hain (8 nahi), toh hum zyada sectors mein spill karte hain.
Count. $⌊ t /4 ⌋$ $0..7$ leta hai $\Rightarrow N_{tx} = 8$ . Kyun? Distinct sector count hi transactions set karta hai.
Efficiency. $\frac{128}{32 \times 8} = \frac{128}{256} = 0.5 = 50%$ . Kyun? Har fetched sector ka aadha woh data hai jo humne skip kiya.

Verify: span 0x00–0xF8 (248 bytes); 8 sectors. 256 bytes mein se 128 use kiye = 50%. ✓

Example 4 — Cell D: Stride = exactly ek sector

Forecast: Har thread exactly ek sector jump karta hai. Kya koi do threads kabhi ek sector share karte hain?

Address function. $A (t) = 4 \times 8 t = 32 t$ . Kyun? Yeh tipping point hai jahan stride $T$ ke barabar hai.
Sectors. $sector (t) = ⌊ 32 t /32 ⌋ = t$ — har thread ke liye ek alag sector. Kyun? Kyunki stride exactly $T$ ke barabar hai, koi do threads kabhi ek hi sector mein land nahi kar sakte.
Count. Sectors $0, 1, \dots, 31 \Rightarrow N_{tx} = 32$ . Kyun? 32 distinct sectors yahan master number hai.
Efficiency. $\frac{128}{32 \times 32} = \frac{128}{1024} = 0.125 = 12.5%$ . Kyun? Har 32 fetched bytes mein se sirf 4 use hote hain.

Verify: 32 distinct sectors, 1024 bytes fetched, 128 used → 12.5%. Yeh parent note ka exact "strided" number hai. ✓

Example 5 — Cell E: Huge stride (worst case, limiting value)

Forecast: Jab stride ≥ $T$ ho, aur zyada stride se nuksaan nahi hota — tum already ek transaction per thread par ho. Confirm karo.

Address function. $A (t) = 4 \times 1024 t = 4096 t$ . Kyun? Hum test kar rahe hain ki $T$ se zyada extra stride cheezein aur kharab karti hai ya nahi.
Sectors. $sector (t) = ⌊ 4096 t /32 ⌋ = 128 t$ — sectors $0, 128, 256, \dots$ sab distinct. Kyun? Koi bhi stride $\geq T$ 32 distinct sectors guarantee karta hai. Usse aage, extra stride sirf unhe spread karta hai; count 32 par saturate ho jaata hai.
Count. $N_{tx} = 32$ (32-thread warp ke liye maximum possible). Kyun? 32-thread warp at most 32 distinct sectors touch kar sakta hai — yeh woh ceiling hai.
Efficiency. $\frac{128}{1024} = 12.5%$ — Example 4 se identical. Kyun? Same transaction count → same efficiency; floor reach ho gayi.

Verify: $N_{tx} = 32$ , efficiency $= 0.125$ , Example 4 se match → floor confirmed. ✓

Example 6 — Cell F: Reversed / permuted order

Forecast: Kya threads addresses visit karne ka order matter karta hai, ya sirf addresses ka set?

Address function. $A (t) = 4 (31 - t)$ . Thread 0 reads 0x7C, thread 31 reads 0x00. Kyun? Hum test karna chahte hain ki coalescing permutation ki parwah karta hai ya nahi.
Addresses ka set. Jab $t$ $0..31$ range karta hai, $31 - t$ $31..0$ range karta hai — Example 1 jaisa same set ${0, 4, \dots, 124}$ . Kyun? Coalescing warp ke touch kiye distinct sectors count karta hai, yeh nahi ki unhe kaun touch karta hai.
Sectors. Exactly ${0, 1, 2, 3}$ , toh $N_{tx} = 4$ . Kyun? Example 1 jaisa same sector set → same count.
Efficiency. $100%$ — Example 1 se identical. Kyun? Same $N_{tx}$ → same efficiency.

Recall Kya thread-to-address permutation coalescing change karta hai?

Warp ke andar permutation free hai — sirf sectors ka set matter karta hai ::: Sahi. Reverse, shuffle, ya interleave karo — jab tak 32 addresses same sectors cover karein, $N_{tx}$ unchanged hai. Hardware transactions issue karne se pehle addresses collect karta hai.

Verify: same sector set ${0, 1, 2, 3}$ Example 1 ki tarah → $N_{tx} = 4$ , 100%. ✓

Example 7 — Cell G: Degenerate — sab threads same address read karte hain

Forecast: 32 threads, ek address. Minimum possible $N_{tx}$ kya hai?

Address function. $A (t) = 0$ , constant (zero slope — degenerate "flat" ramp). Kyun? Yeh boundary case hai jahan address function ki stride 0 hai.
Sectors. $sector (t) = ⌊ 0/32 ⌋ = 0$ sab threads ke liye → single sector ${0}$ . Kyun? 32 saare addresses identical hain, toh woh ek se zyada sector mein spread nahi ho sakte.
Count. $N_{tx} = 1$ . Hardware sector 0 ek baar fetch karta hai aur value broadcast karta hai saare 32 threads ko. Kyun? Ek distinct sector → ek transaction, ek warp load ki absolute minimum cost.
Yahan efficiency ka matlab (broadcast footnote). Hamare page-wide numerator "bytes used $= W \cdot E = 128$ " assume karta hai ki 32 threads 32 distinct elements chahte hain. Broadcast mein woh same element chahte hain, toh actually sirf $E = 4$ distinct useful bytes hain, 128 nahi. Formula literally plug karne par $\frac{128}{32 \times 1} = 4.0 = 400%$ milega — ek nonsense >100% number jo sirf isliye aata hai kyunki numerator identical requests over-count karta hai. Toh is ek degenerate cell ke liye hum iske bajaaye honest operational headline report karte hain: $N_{tx} = 1$ , minimum possible transaction count, single fetched sector saare 32 threads reuse karte hain. Broadcast sabse achi pattern hai, koi 400%-efficient nahi. Yeh step kyun? Is footnote ke bina reader $400%$ figure trust kar leta ya sochta ki model toot gaya; numerator ka assumption (distinct elements) name karna exactly woh hai jo paradox resolve karta hai.

Verify: single distinct sector → $N_{tx} = 1$ , minimum. Distinct useful bytes $= 4$ ; literal formula flagged nonsense $128/32 = 4.0$ deta, isliye hum $N_{tx} = 1$ report karte hain. ✓

Example 8a — Cell H: Real-world word problem (row-major matrix)

Forecast: Kaun sa layout traversal — column down ya row across — warp ko consecutive addresses deta hai?

Column access address function. $A (t) = 4 \times (t \cdot 1024) = 4096 t$ bytes. Kyun? Row-major matlab consecutive rows 1024 floats = 4096 bytes apart hain.
Column sectors. $⌊ 4096 t /32 ⌋ = 128 t$ — sab distinct → $N_{tx} = 32$ , efficiency $12.5%$ . Yeh Example 5 disguise mein hai! Kyun? Row-major storage mein column ka stride = row length ≫ sector → worst case. (Yeh upar slope-regime bullets mein referenced column-access case hai.)
Row access address function. $A (t) = 4 t$ → Example 1 jaisa identical → $N_{tx} = 4$ , $100%$ . Kyun? Thread ke saath varying last index stride 1 deta hai → perfect coalescing.
Design conclusion. Row-major array ke liye, consecutive threads ko consecutive columns read karwao (last index threadIdx.x ke saath vary karta hai). Columns efficiently process karne ke liye, pehle shared memory mein transpose karo. Yeh step kyun? Steps 2 aur 3 ne prove kiya ki row traversal column traversal se $8 \times$ kam transactions karta hai; woh decide karna kya karna hai measured fact ke saath hi arithmetic ko engineering rule mein turn karta hai — warna derivation actionable takeaway ke bina khatam ho jaati.

Verify: column: $N_{tx} = 32$ , 12.5%; row: $N_{tx} = 4$ , 100%. Row-over-column speedup $= 32/4 = 8 \times$ transactions mein. ✓

Example 8b — Cell J: Mixed / irregular warp (koi single stride nahi)

Forecast: Warp aadha achha, aadha kharab hai. Kya achhe 16 kharab 16 ko rescue karte hain, ya total bas add up ho jaata hai?

Dono halves ke address functions.
- Threads $0..15$ : $A (t) = 4 t$ → bytes 0,4,...,60.
- Threads $16..31$ : $A (t) = 4 (1000 + 64 (t - 16)) = 4000 + 256 (t - 16)$ → bytes 4000, 4256, 4512, .... Kyun? Ek mixed warp mein koi closed-form single slope nahi hoti, toh hume har sub-pattern alag count karni padti hai, phir sectors ka union lena padta hai.
Contiguous half ke sectors. Bytes 0..60 → sectors $⌊ 0/32 ⌋ .. ⌊ 60/32 ⌋ = 0, 1$ . Yeh 2 sectors hain. Kyun? 16 floats = 64 bytes exactly two aligned sectors span karte hain — good half locally coalesced hai.
Scattered half ke sectors. $sector (t) = ⌊(4000 + 256 (t - 16)) /32 ⌋ = 125 + 8 (t - 16)$ $t - 16 = 0..15$ ke liye → sectors $125, 133, 141, \dots$ sab distinct → 16 sectors. Kyun? Yahan stride 256 bytes $\geq T$ hai, toh 16 mein se har ek apne sector mein land karta hai (locally Example 5 behaviour).
Union aur count. Dono sector sets disjoint hain (0–1 vs 125–245), toh $N_{tx} = 2 + 16 = 18$ . Kyun? Total transactions = poore warp ke distinct sectors; good half kharab half ko rescue nahi karta, lekin apni cost low rakhta hai.
Efficiency. $\frac{128}{32 \times 18} = \frac{128}{576} \approx 0.222 = 22.2%$ . Kyun? Sirf 16 contiguous reads dense the; 16 scattered reads ne efficiency neeche kheechi lekin poori tarah 12.5% floor tak nahi.

Verify: contiguous half → 2 sectors; scattered half → 16 sectors; disjoint union → 18; efficiency $128/576 \approx 0.2222$ . 12.5% floor aur 100% ceiling ke beech, jaisa ek mixed pattern hona chahiye. ✓

Example 9 — Cell I: Exam twist — Array of Structs (hidden stride)

Forecast: Struct 12 bytes ka hai, toh x fields 12 bytes apart hain — 4 nahi. Kya woh 32-byte sectors mein neat pack hota hai?

Address function. $A (t) = 12 t$ . Threads 0, 12, 24, 36, ... par read karte hain. Kyun? Struct size stride ban jaata hai jab tum ek thread per field touch karte ho.
Sectors touched. $sector (t) = ⌊ 12 t /32 ⌋$ . Span evaluate karo: $A (31) = 372 \Rightarrow$ sector $⌊ 372/32 ⌋ = 11$ . Smallest sector 0 hai. Kya $0..11$ sab hit hote hain? Step $12/32 = 0.375$ sector per thread, toh consecutive threads kabhi poora sector skip nahi karte → haan, sectors $0, 1, \dots, 11$ sab touch hote hain. Kyun? Count karne se pehle confirm karna padega ki range mein koi sector skip toh nahi hua.
Count. $N_{tx} = 12$ . Kyun? 12 distinct sectors master number hai.
Efficiency. $\frac{128}{32 \times 12} = \frac{128}{384} \approx 0.333 = 33.3%$ . Kyun? Humne y aur z fields fetch kiye jo humne kabhi use nahi kiye — har sector ka 2/3 waste.

Verify: span 0–372 bytes → sectors 0..11 = 12 sectors; efficiency $128/384 = 1/3$ . ✓

Ek nazar mein summary

Recall Har example ke liye

$N_{tx}$ fill karo (hamara fixed $T = 32$ B model) Ex1 seq → 4 · Ex2 misaligned → 5 · Ex3 stride2 → 8 · Ex4 stride8 → 32 · Ex5 huge → 32 · Ex6 reversed → 4 · Ex7 broadcast → 1 · Ex8a column → 32 / row → 4 · Ex8b mixed → 18 · Ex9 AoS → 12 ::: Shape yaad rakho: gentle stride = kam transactions; stride ≥ sector = 32 (floor); mixed = parts ka sum.

Recall Kisi bhi bade uniform stride ke liye efficiency 12.5% par capped kyun hai, isse neeche kyun nahi?

Kyunki warp mein sirf 32 threads hain, toh at most 32 distinct sectors hain ::: Ek transaction per thread ke saath tum $32 T$ bytes fetch karte ho aur $32 E$ use karte ho; ratio $E / T = 4/32 = 12.5%$ single-element-per-thread uniform-stride pattern ke liye hard floor hai.

Prerequisites & neighbours: 6.2.01-GPU-execution-model (warps & SIMT), 6.2.09-Memory-hierarchy (jahan sectors aur L2 rehte hain, aur caching ek bure pattern ko kyun soften karta hai but fix nahi karta), 6.2.04-Shared-memory-and-synchronization (transpose fix), 6.3.02-Roofline-model (bandwidth kyun matter karta hai), 5.1.05-Cache-coherence-protocols (line/sector granularity origins), aur Hinglish walkthrough 6.2.08 Coalesced memory access (Hinglish).