Memory hierarchy (global, shared, registers): Exercises

6.2.7 · D4 · Hardware › GPU Architecture › Memory hierarchy (global, shared, registers)

Yeh page ek self-test hai. Har problem clearly likhi gayi hai, phir uska poora worked solution ek collapsible box mein chhupa hua hai. Pehle har ek ko paper pe try karo, phir reveal karne ke liye click karo.

Problems ek difficulty ladder pe chadhti hain:

L1 Recognition — kya tum cheezein naam le ke place kar sakte ho?
L2 Application — kya tum ek rule mein numbers plug kar sakte ho?
L3 Analysis — kya tum do options compare karke reason kar sakte ho ki kyun?
L4 Synthesis — kya tum kai rules ko combine karke ek design bana sakte ho?
L5 Mastery — kya tum invent, estimate, aur ek full answer defend kar sakte ho?

Yahan sab kuch parent topic pe build karta hai. Agar koi term unfamiliar lage, toh linked prerequisite notes hain: GPU Thread Hierarchy, Warp Execution Model, CUDA Memory Types, Memory Latency Hiding, Cache Architecture, Memory Bandwidth Optimization, Matrix Multiplication Optimization, aur Roofline Model.

Shuru karne se pehle, poore mental model ko fix karne ke liye ek picture — memory ki "distance ladder":

Figure — Memory hierarchy (global, shared, registers)

Yeh figure kaise padhein: har horizontal bar ek storage level hai, aur bar ki length uski latency ke saath badhti hai (log-scaled, taki door-door ke numbers fit ho sakein). Upar se neeche padho: green Registers bar chhota (~1 cycle) hai aur sabse upar hai — fastest, smallest, compute unit ke sabse paas. Uske neeche yellow Shared Memory bar (~65 cycles), phir blue L2 Cache bar (~200 cycles), aur finally neeche lamba red Global Memory bar (~600 cycles) — slowest, largest, sabse door. Upar wala green arrow aur neeche wala red arrow trend spell out karte hain: ladder pe upar jaana faster/smaller/closer storage deta hai, neeche jaana slower/larger/farther storage deta hai. Is page ke har exercise mein hot data ko ek high (green) rung pe rakhne ki koshish ki jaati hai.

L1 — Recognition

Exercise 1.1

In chaar storage levels ko fastest se slowest tak order karo: L2 cache, registers, global memory, shared memory.

Recall Solution 1.1

Hum kya karte hain: parent ki summary table ka latency column padhte hain aur ascending sort karte hain.

Level	Latency (cycles)
Registers	1
Shared memory	30–100
L2 cache	~200
Global memory	400–800

Answer: registers → shared memory → L2 cache → global memory.

Yeh order kyun: speed physical proximity se aati hai. Registers arithmetic unit ke andar baithe hain; global memory off-chip DRAM (VRAM) hai jahan tak light ko physically travel karna padta hai.

Exercise 1.2

Har level ke liye uska scope batao (per-thread, per-block, ya per-device): registers, shared memory, global memory.

Recall Solution 1.2

Registers → per-thread (private; koi doosra thread unhe read nahi kar sakta).
Shared memory → per-block (same block ke har thread ko share hoti hai).
Global memory → per-device (har block ka har thread ise dekh sakta hai).

Kyun: scope aur speed saath jaate hain. Private cheezein fast hain (koi coordination nahi); shared cheezein ko arbitration chahiye (banks); global cheezein ke liye lambi trip chahiye. Audience badhana time cost karta hai.

Exercise 1.3

Teeno levels ke liye — registers, shared memory, global memory — batao kaun ise manage karta hai (compiler-automatic ya programmer-explicit). Phir is student claim ko judge karo: "Compiler decide karta hai ki shared memory mein kya jaata hai."

Recall Solution 1.3

Har level kaun manage karta hai:

Registers → compiler (automatic; compiler tumhare local variables ko registers mein assign karta hai bina tumhare maange).
Shared memory → programmer (explicit; tum __shared__ likhte ho aur khud fill karte ho — isliye ise "user-controlled cache" kehte hain).
Global memory → programmer (explicit; tum ise allocate karte ho aur har read aur write choose karte ho).

Claim False hai. Compiler registers manage karta hai, shared memory nahi. Shared aur global memory dono programmer-controlled hain; sirf registers automatic hain.

L2 — Application

Formula Teen rules jo tum apply karoge

Register occupancy — threads jo ek SM run kar sakta hai given register budget $R_{t o t a l}$ (total registers per SM, upar define kiya) aur per-thread usage $r$ . Yahan $N_{ma x}$ SM ka fixed hardware thread limit hai (intuition box mein introduce kiya, typically 2048): $N_{t h r e a d s} = ⌊ \frac{R _{t o t a l}}{r} ⌋, Occupancy = \frac{N _{t h r e a d s}}{N _{ma x}}$

Warp-granularity caveat: real hardware pe, registers ek warp (32 threads) at a time distribute hote hain, toh schedulable thread count actually 32 ke multiple tak round down hota hai: $N_{t h r e a d s}^{hw} = ⌊ \frac{N _{t h r e a d s}}{32} ⌋ \times 32.$ Upar wala plain-floor formula idealized budget bound hai; warp-rounded version woh hai jo hardware actually schedule karta hai. Hum dono report karte hain jab woh differ karte hain.

Bank mapping — $B$ banks (har ek $w$ bytes wide) mein se kaunse bank mein address $a$ land karta hai: $bank (a) = (\frac{a}{w}) mod B$

Coalescing — ek warp ko kitne memory transactions chahiye, yeh un distinct aligned segments ki count ke barabar hai jhe woh touch karta hai.

Exercise 2.1

Ek SM ke paas $R_{t o t a l} = 65, 536$ registers hain. Ek kernel $r = 48$ registers per thread use karta hai. Dusre limits ko ignore karte hue, concurrent threads ki maximum number kya hai? Idealized budget bound aur warp-rounded hardware value dono do.

Recall Solution 2.1

Kya / Kyun: har thread fixed pot $R_{t o t a l}$ se $r$ registers le jaata hai, toh hum pot ko per-thread cost se divide karte hain aur floor operator se neeche round karte hain (tum fractional thread nahi run kar sakte). $N_{t h r e a d s} = ⌊ \frac{65 , 536}{48} ⌋ = ⌊ 1365.33 ⌋ = 1365 threads (idealized budget bound).$ Warp rounding: hardware 32 ke chunks mein allocate karta hai, toh $N_{t h r e a d s}^{hw} = ⌊ \frac{1365}{32} ⌋ \times 32 = 42 \times 32 = 1344 threads.$ Answer: 1365 threads plain-floor budget bound se, 1344 threads warp-granularity rounding apply karne ke baad — woh value jo SM actually schedule karta hai.

Exercise 2.2

$N_{ma x} = 2048$ threads/SM use karke, Exercise 2.1 ke kernel ke liye occupancy kya hai? Dono tarike se report karo.

Recall Solution 2.2

Idealized budget bound: $Occupancy = \frac{1365}{2048} = 0.6665 \dots \approx 66.7%.$ Warp-rounded (real hardware): $Occupancy^{hw} = \frac{1344}{2048} = 0.65625 = 65.6%.$ Occupancy kyun matter karta hai: GPUs 400–800 cycle global-memory wait ko doosre ready warp pe switch karke hide karte hain (dekho Memory Latency Hiding). Kam resident threads = switch karne ke liye kam warps = SM idly wait karta rehta hai. Yahan honest hardware number ~65.6% hai, toh roughly two-thirds hiding power available hai.

Exercise 2.3

Shared memory ke paas $B = 32$ banks hain, har ek $w = 4$ bytes wide. Byte address $a = 260$ kaunse bank mein padta hai?

Recall Solution 2.3

Kya / Kyun: hum pehle byte address ko word index mein convert karte hain (width $w$ se divide karo, kyunki banking per 4-byte word hoti hai), phir mod $B$ lete hain yeh dekhne ke liye ki kaunsa bank cycle around karta hai iske paas. $bank (260) = (\frac{260}{4}) mod 32 = 65 mod 32 = 1.$ Answer: bank 1.

L3 — Analysis

Exercise 3.1

32 threads ka ek warp float (4 bytes) ka ek array read karta hai. Segments 128 bytes ke hain. Do access patterns compare karo aur har ek ke liye transaction count do:

(a) thread $t$ address $base + 4 t$ read karta hai.
(b) thread $t$ address $base + 128 t$ read karta hai.

Phir (c) answer karo: case (a) mein kya hota hai agar $base$ 128-byte boundary pe aligned nahi hai?

Recall Solution 3.1

(a) Coalesced (aur aligned). Assume karo $base$ 128 ka multiple hai (aligned). Addresses $base + 0$ se $base + 124$ tak span karte hain, total $32 \times 4 = 128$ bytes = exactly ek 128-byte segment. $N_{t r an s a c t i o n s} = ⌈ 128/128 ⌉ = 1.$ (b) Strided. Har thread ka address 128 bytes apart hai, toh har thread apne khud ke segment mein land karta hai. $N_{t r an s a c t i o n s} = 32.$ (c) Misaligned base — hidden edge case. Maano $base = 64$ (128 ka multiple nahi). 128 requested bytes ab byte 64 se byte 191 tak run karte hain, 128 pe boundary straddle karte hue. Woh do segments touch karta hai (0–127 aur 128–255): $N_{t r an s a c t i o n s} = 2.$ Toh ek bilkul consecutive access bhi double transactions pay karta hai agar uska starting address misaligned ho. Lesson: coalescing ke liye dono chahiye — ek consecutive pattern aur ek aligned base — isliye CUDA allocators 128/256-byte-aligned pointers return karte hain.

Analysis: aligned coalesced-vs-strided ratio $32/1 = 32 \times$ hai; misalignment quietly tumhare best case ko $2 \times$ penalty mein badal deta hai. Same math, same 32 floats chahiye — lekin pattern (b) bus ke across 32 times zyada bytes drag karta hai, har fetch ka 31/32 waste karta hai. Yeh Memory Bandwidth Optimization mein sabse bada lever hai.

Yeh figure kaise padhein: green dots ki upar wali row coalesced aligned case dikhati hai — 32 thread requests ek single shaded green 128-byte segment mein pack hain (ek transaction). Yellow dots ki beech wali row misaligned case dikhati hai — same consecutive requests boundary line straddle karte hain aur doosre segment mein spill karte hain (do transactions). Red dots ki neeche wali row strided case dikhati hai — har request apne alag shaded red segment mein land karta hai, toh warp ko kai transactions chahiye. Horizontal axis byte address hai; vertical dotted white lines 128-byte segment boundaries mark karte hain.

Exercise 3.2

Kernel A 32 registers/thread use karta hai; kernel B 64 registers/thread use karta hai. Dono ek SM pe run karte hain jiske paas 65,536 registers hain, $N_{ma x} = 2048$ . Dono ke liye occupancy compute karo aur performance implication explain karo.

Recall Solution 3.2

Kernel A: $N = ⌊ 65536/32 ⌋ = 2048$ , $N_{ma x} = 2048$ se cap → occupancy $= 2048/2048 = 100%$ . Kernel B: $N = ⌊ 65536/64 ⌋ = 1024$ → occupancy $= 1024/2048 = 50%$ .

Analysis: registers double karne se occupancy half ho gayi. Kernel B ke paas latency hide karne ke liye sirf aadhe warps hain. Lekin — subtlety — B phir bhi jeet sakta hai agar woh extra registers memory traffic eliminate kar dein (global ke kam trips). Yeh classic register-pressure vs. occupancy tradeoff hai: tum kabhi occupancy ko akele optimize nahi karte, sirf us memory traffic ke saath jo woh tumhe dilaata hai.

Exercise 3.3

Do warps shared memory access karte hain. Warp X mein, sab 32 threads word index $5$ read karte hain (same address). Warp Y mein, thread $t$ word index $t$ read karta hai. Har ek ke liye serialization factor $S$ do.

Recall Solution 3.3

Warp Y (thread $t$ → word $t$ ): words $0..31$ banks $0..31$ pe map karte hain — sab distinct → $S = 1$ . Koi conflict nahi, full speed.

Warp X (sab → word 5): naïvely 32 threads ek bank hit karte hain, toh $bank (a) = count = 32$ , $S = 32$ suggest karta hai. Lekin shared memory mein ek broadcast rule hai: jab kai threads same address read karte hain, hardware ek hi baar mein broadcast karta hai → $S = 1$ .

Answer: dono ka $S = 1$ hai, lekin alag alag reasons se — Y isliye ki addresses banks mein spread hain, X isliye ki identical addresses broadcast trigger karte hain. Dangerous case in between hai: kai threads same bank pe different addresses hit karte hain (woh truly serialize karta hai).

L4 — Synthesis

Exercise 4.1

Tum ek matrix multiply ko $16 \times 16$ tiles se tile karte ho (dekho Matrix Multiplication Optimization). A aur B ka har float tile shared memory mein store hai. (a) Ek tile kitne bytes occupy karta hai? (b) Ek block ko dono tiles ke liye kitni shared memory chahiye? (c) Agar SM ke paas 96 KB shared memory hai, toh kitne blocks fit hote hain (sirf shared-memory limited)?

Recall Solution 4.1

(a) $16 \times 16 = 256$ floats $\times 4$ bytes $= 1024$ bytes = 1 KB per tile. (b) Do tiles (As aur Bs): $2 \times 1024 = 2048$ bytes = 2 KB per block. (c) $⌊ 96 KB /2 KB ⌋ = ⌊ 98304/2048 ⌋ = 48$ blocks. Synthesis point: shared-memory footprint registers ke saath ek doosra occupancy limiter hai. True block count registers limit, shared-memory limit, aur hardware block cap mein se minimum hai.

Exercise 4.2

Transpose kernel ke liye, parent __shared__ float tile[32][33] declare karta hai. (a) 32 ke bajay 33 columns kyun? (b) Padding ke liye per block kitne extra bytes cost hain?

Recall Solution 4.2

(a) [32][32] ke saath, row $i$ ka column $j$ word index $32 i + j$ pe baithta hai. Jab threads transpose ke dauran column read karte hain (varying $i$ , fixed $j$ ), unke indices 32 se differ karte hain → $(32 i + j) mod 32 = j$ sabhi $i$ ke liye → sab 32 same bank hit karte hain → 32-way conflict. 33 tak padding karne se column stride 33 ho jaata hai, aur $g cd (33, 32) = 1$ , toh consecutive column elements 32 different banks mein land karte hain → conflict-free. (b) Extra column = $32 rows \times 1 float \times 4 = 128$ bytes per block. Synthesis point: 128-byte "waste" transpose step pe 32× speedup kharidta hai — is page ka sabse sasta trade.

Exercise 4.3

$2048 \times 2048$ matrices ke saath $16 \times 16$ tiles ke liye parent ka tiled-vs-naive global-traffic count redo karo. (a) Naive global reads total. (b) Tiled global reads total. (c) Reduction factor, explanation ke saath.

Recall Solution 4.3

$N = 2048$ , tile $= 16$ , tiles per dimension $= 2048/16 = 128$ .

(a) Naive: $N^{2} = 204 8^{2} = 4, 194, 304$ output threads mein se har ek global se $2 N = 4096$ values read karta hai. $4, 194, 304 \times 4096 = 17, 179, 869, 184 \approx 1.72 \times 1 0^{10} reads.$ (b) Tiled: har block, per k-step, ek A-tile + ek B-tile $= 2 \times 256 = 512$ elements load karta hai, $128$ k-steps pe $= 512 \times 128 = 65, 536$ reads per block. Blocks ki number $= 128 \times 128 = 16, 384$ . $65, 536 \times 16, 384 = 1, 073, 741, 824 \approx 1.07 \times 1 0^{9} reads.$ (c) Reduction factor: $\frac{17 , 179 , 869 , 184}{1 , 073 , 741 , 824} = 16 \times .$ Explanation: reduction factor tile edge length ke barabar hota hai (yahan 16). Exactly 16 kyun? Naive version mein har output element har input value jo use chahiye woh global se seedha re-reads karta hai; tiled version mein, ek baar tile shared memory mein load hone ke baad, us row/column share karne wale sab 16 threads use 16 baar reuse karte hain pehle woh evict ho. Isliye har loaded byte 16 uses pe amortized hai, global traffic ko usi factor of 16 se cut karta hai.

Synthesis point: bade tiles = zyada reuse (bada reduction factor), lekin bade tiles zyada shared memory cost karte hain (Exercise 4.1) aur zyada registers — reuse, occupancy, aur capacity ke beech recurring three-way tension.

L5 — Mastery

Exercise 5.1

Ek GPU ka peak compute $P = 10$ TFLOP/s ( $10 \times 1 0^{12}$ FLOP/s) aur peak bandwidth $β = 500$ GB/s ( $500 \times 1 0^{9}$ B/s) hai. (a) Ridge-point intensity $I_{⋆}$ nikalo. (b) Ek naive kernel $8$ bytes read par $2$ FLOP karta hai (do freshly-loaded floats pe ek multiply-add). Kya woh memory-bound hai ya compute-bound? (c) 16× reuse dene ke baad tiling, naya intensity kya hai, aur kya woh ab compute-bound hai?

Recall Solution 5.1

(a) $I_{⋆} = P / β = \frac{10 \times 1 0 ^{12}}{500 \times 1 0 ^{9}} = 20$ FLOP/byte. (b) Naive intensity $I = 2 FLOP /8 B = 0.25$ FLOP/byte. Kyunki $0.25 < 20 = I_{⋆}$ , kernel ridge ke left side mein hai → memory-bound. Woh data ka wait karte karte almost sab compute waste karta hai. (c) Tiling har loaded byte ko 16× reuse karta hai, toh global-se-bytes 16× drop karte hain jabki FLOPs same rehte hain → intensity 16× badhti hai: $I_{n e w} = 0.25 \times 16 = 4$ FLOP/byte. Phir bhi $4 < 20$ , toh abhi bhi memory-bound hai, lekin ab memory roofline pe $16 \times$ faster run karta hai. Mastery point: tiling tumhe roofline pe rightward move karta hai. Ridge cross karke compute-bound hone ke liye tumhe aur bade tiles / register blocking chahiye taaki $I$ 20 se zyada ho.

Yeh figure kaise padhein: axes log–log hain — arithmetic intensity $I$ (FLOP/byte) neeche horizontal pe, attainable performance (FLOP/s) side pe. White line roofline hai: left mein ek rising blue-dashed part (memory-bound — performance bandwidth $β \cdot I$ se limited) aur right mein ek flat yellow-dashed ceiling (compute-bound — performance $P$ pe capped). Woh dotted vertical line pe milte hain, ridge $I_{⋆} = 20$ . Red dot ( $I = 0.25$ , naive) sloped part ke neeche far left mein baitha hai; green dot ( $I = 4$ , tiled) ridge ki taraf rightward slide kar gaya hai lekin abhi bhi slope ke neeche hai — abhi bhi memory-bound, bas faster.

Exercise 5.2

Design decision: tumhare paas ek kernel hai jo 100% occupancy pe bottlenecked hai lekin phir bhi slow hai. Profiling dikhata hai ki woh fully coalesced access ke saath memory-bound hai. Tumhare paas teen levers hain: (i) zyada data cache karne ke liye per thread registers badhao, (ii) shared-memory tiling add karo, (iii) block size badhao. Tum kaunsa choose karte ho aur kyun? 3–4 sentences mein argue karo.

Recall Solution 5.2

(ii) shared-memory tiling pick karo. Kernel memory-bound hai already-coalesced access ke saath, toh bus efficiently use ho rahi hai lekin bahut zyada data fetch ho raha hai — cure reuse hai, aur coalescing nahi. Lever (i) occupancy lower karne ka risk hai (Exercise 3.2) aur sirf tab help karta hai jab data ek single thread ke andar reuse ho; lever (iii) move kiye gaye total bytes bilkul kam nahi karta. Tiling arithmetic intensity raise karta hai (Exercise 5.1), tumhe roofline pe rightward slide karta hai — yahan sirf yahi lever hai jo total global traffic pe attack karta hai, jo actual bottleneck hai.

Exercise 5.3

Full estimation. Ek kernel 1,048,576 threads launch karta hai, jisme se har ek perfectly coalesced pattern mein global memory se exactly ek float (4 B) read karta hai. GPU bandwidth $500$ GB/s hai. (a) Total bytes moved. (b) Ise move karne ka minimum time (bandwidth-bound floor). (c) Agar access fully strided hoti (32× zyada transactions), toh new time estimate karo.

Recall Solution 5.3

(a) $1, 048, 576 \times 4 = 4, 194, 304$ bytes $\approx 4$ MB useful data. (b) $t = \frac{4 , 194 , 304 B}{500 \times 1 0 ^{9} B/s} = 8.388608 \times 1 0^{- 6} s \approx 8.39 μ s .$ (c) Strided access 32× zyada bytes fetch karta hai (single floats ke liye whole segments), toh effective transferred bytes $\approx 32 \times$ ho jaate hain aur time scale up karta hai: $8.39 μ s \times 32 \approx 268.4 μ s$ . Mastery point: coalesced floor ek hard physical limit hai — tum bandwidth allow karne se faster nahi ja sakte. 32× penalty poori tarah access pattern se self-inflicted hai, isliye coalescing kisi bhi optimization pass mein pehle check ki jaati hai.