Memory hierarchy (global, shared, registers): Question bank

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

Shuru karne se pehle, hum woh teen ideas build karte hain jinpar is page ke saare traps tike hain — ek warp, shared-memory ka ek bank, aur global-memory ka ek segment. Pehle yeh padho; traps inhe assume karte hain.

Definition Shared-memory banks (traps se pehle yeh build karo)

Shared memory ek drawer nahi hai — yeh banks mein split hoti hai, parallel chhoti memories jo har ek ek cycle mein ek request answer kar sakti hain. Common GPUs par $B = 32$ banks hote hain, har ek 4-byte words hold karta hai, lekin yeh architectures mein alag ho sakta hai — 32 ko typical value maano, universal law nahi.

Ek byte address $a$ bank number mein aata hai $bank (a) = ⌊ \frac{a}{w} ⌋ mod B, w = 4 bytes, B = 32.$ $⌊ a / w ⌋$ kyun? Hum pehle byte address ko word index mein convert karte hain (konsa 4-byte slot), phir us index ko $mod B$ se 32 banks ke around wrap karte hain. Toh consecutive words 0,1,2,… banks 0,1,2,… mein jaate hain aur word 32 wapas bank 0 par wrap ho jaata hai.

Ek bank conflict tab hota hai jab ek hi warp ke do threads alag-alag words maangein jo ek hi bank par wrap ho jaayein — unhe requests baar baar karni padti hain. Figure s01 dekho: lal cells sab ek hi bank par collide kar rahi hain.

Figure — Memory hierarchy (global, shared, registers)

Warp, bank, aur segment apne dimag mein rakho — neeche ke traps inhi par ghoomte hain.

True or false — justify karo

True or false: Registers sabse fast memory hai, isliye thread pichhe zyada registers use karne se kernel hamesha fast hoti hai.

False. Har register access fast hai, lekin SM ke paas fixed register file hai; thread pichhe zyada registers matlab fewer warps fit hote hain, toh latency hiding collapse ho jaati hai aur kernel usually slow ho jaati hai. Speed-per-access aur total throughput alag-alag cheezein hain.

True or false: Shared memory ek tarah ka cache hai jo hardware tumhare liye automatically fill karta hai.

False. Shared memory ek ==scratchpad hai jo tum khud explicit loads aur __syncthreads() se fill karte ho==. L1 cache automatic wala hai; woh aksar ek hi physical SRAM share karte hain lekin alag-alag control hote hain.

True or false: Global memory slow hai kyunki iska bandwidth low hai.

False. Global memory ki bahut badi bandwidth hai (saikdon GB/s) lekin high latency hai (400–800 cycles). Yeh ek wide, lamba pipe hai: per second bahut saara data, lekin har individual byte ko pahunchne mein bahut time lagta hai. Latency aur bandwidth alag-alag axes hain.

True or false: Agar do threads ek hi shared-memory address padhein, toh woh bank conflict hai.

False. Same-address reads ek broadcast trigger karte hain: ek bank read saare requesting threads ko share ho jaata hai, bina kisi serialization ke. Conflict ke liye alag-alag words chahiye jo ek hi bank mein land karein (upar bank definition dekho).

True or false: Coalescing ka matlab hai threads program order mein consecutive addresses access karein.

False. Iska matlab hai ek warp ke 32 threads milke kam aligned segments touch karein (jaisa upar define kiya). Warp ke andar order matter nahi karta — ek 128-byte segment ke andar addresses ka permutation bhi ek transaction mein coalesce ho jaata hai.

True or false: Shared array ko [32][33] se pad karna memory waste karta hai aur isse bachna chahiye.

Context mein False. Extra column har row ki start ko ek word shift kar deta hai, toh

mod 32

bank mapping ab poori column ko ek bank par nahi bhejti (figure s01 ke caption mein work through kiya gaya hai). Chhoti memory cost ek badi serialization win dilati hai.

True or false: Kernel ke andar declare kiya gaya variable hamesha registers mein rehta hai.

False. Compiler koshish karta hai, lekin arrays jinhe statically index nahi kar sakta, ya jab register pressure zyada ho toh spill ho jaate hain, local memory mein jaate hain — jo physically slow global memory mein rehti hai, naam ke bawajood.

True or false: Shared memory poore GPU mein shared hoti hai.

False. Iska scope sirf ek thread block tak hai. Alag-alag blocks ke threads ek doosre ki shared memory nahi dekh sakte; jo storage sab blocks share karte hain woh sirf global memory hai (aur L2).

Error pakado

Claim: "Meri kernel 32 registers/thread use karti hai aur 100% occupancy milti hai, toh registers mujhe limit nahi kar rahe — main aur registers freely add kar sakta hoon." Kya galat hai?

Tum edge par ho. Exactly budget-fitting point par, ek bhi register add karne se puri ek block ke threads drop ho sakte hain (occupancy ek step function hai, smooth nahi). 100% hona matlab register pressure maximally tight hai, slack nahi.

Claim: "Maine har access shared memory par karwa di, toh bank conflicts ab mujhe hurt nahi kar sakte." Kya galat hai?

Bank conflicts shared memory ke andar rehte hain. Data ko shared memory mein move karne se global-latency problems khatam hoti hain lekin ek bilkul naya serialization problem ban sakta hai agar tumhara stride bahut saare threads ko ek hi bank par map kare (upar diye

bank (a)

formula ke hisaab se).

Code: transpose ke liye out[col*N + row] = in[row*N + col]; ko "coalesced kaha jaata hai kyunki read in[row*N + col] consecutive hai." Flaw pakado.

Read coalesce hoti hai, lekin write out[col*N + row] neighboring threads ke beech N se jump karta hai → ek segment per thread → uncoalesced. Sirf ek access par coalescing judge karna classic mistake hai.

Claim: "Occupancy 75% hai, toh main GPU ka 25% waste kar raha hoon — mujhe ise badhana chahiye." Error pakado.

Occupancy ek means hai, goal nahi. Agar 75% already itne warps deta hai ki saari memory latency hide ho jaaye, toh aur badhane se kuch nahi hoga — aur ise force karna (registers cut karke) variables spill kar sakta hai aur cheezein badtar kar sakta hai.

Code: __shared__ float As[16][16]; phir As[ty][tx] = A[...] — ek student kehta hai "yeh synchronization ki zaroorat ke bina load hota hai." Error pakado.

Cooperative load ke baad doosre threads ke likhe tiles koi thread padhne se pehle __syncthreads() call karni chahiye; warna ek fast thread stale/garbage values padh leta hai. Missing barriers silent, intermittent bugs hote hain.

Claim: "Strided global access sirf 2× slower hai coalesced se." Error pakado.

Bade stride ke saath 32 threads mein se har ek apne alag segment mein land kar sakta hai → 1 ki jagah 32 transactions, yaani 32× tak memory traffic. Penalty is baat par depend karti hai ki addresses segments mein kitne bikhar jaate hain, koi fixed 2× nahi.

Claim: "L1 cache aur shared memory bilkul alag hardware hai, toh ek use karne se doosra kabhi affect nahi hota." Error pakado.

Bahut saari architectures par woh ek hi SRAM pool se carved out hote hain. Zyada shared memory maangne se L1 cache shrink ho sakti hai (aur vice versa), toh woh ek doosre ke saath trade off karte hain.

Why questions

Kyun thread pichhe zyada registers dene se performance girti hai jab ki registers sabse fast storage hai?

Kyunki register file ek fixed pot hai jo saare resident threads mein split hoti hai. Har thread ko zyada registers → fewer resident warps → stall ke dauran swap in karne ke liye fewer standby warps → latency hiding fail ho jaati hai. Bottleneck access-speed se warp-availability par shift ho jaata hai.

GPUs latency CPUs ki tarah caches add karke nahi, balki threads add karke kyun hide karte hain?

GPUs massive parallelism par bet karte hain: hazaron threads ke saath, jab bhi ek warp memory ka wait karta hai, doosra kaam ke liye ready hota hai, arithmetic units busy rehti hain. Yeh "hide, don't shorten" strategy bahut saare warps maangti hai, isliye occupancy itna matter karta hai.

Shared-memory bank width usually 4 bytes aur bank count 32 kyun hoti hai?

Kyunki ek warp 32 threads ka hota hai aur ek natural word 4 bytes ka hota hai: agar 32 threads har ek consecutive 4-byte word uthayein, woh 32 alag-alag banks hit karte hain → zero conflict, ek cycle. Hardware common access pattern ko conflict-free banane ke liye tuned hai (lekin doosri generations alag widths choose kar sakti hain).

Memory coalescing concept sirf global memory ke liye kyun hai, registers ke liye nahi?

Registers thread ke liye private hain — combine karne ke liye koi shared transaction nahi hai. Global memory fixed aligned segments mein fetch hoti hai jo poore warp ko serve karte hain, isliye warp ke addresses ka segments mein grouping transaction count determine karta hai.

Do kernels jo identical arithmetic karte hain, unke runtimes wildly alag kyun ho sakte hain?

Kyunki runtime usually memory behaviour se dominate hoti hai — coalescing, bank conflicts, occupancy, reuse — FLOPs se nahi. Roofline view: zyatar kernels memory-bound hoti hain, toh data movement, math nahi, clock set karta hai.

Shared memory mein tiling matrix multiply ko speed up kyun karta hai jab ki total math same rehti hai?

Tiling arithmetic intensity badhata hai: global memory se ek baar load ki gayi har value ko fast shared memory se kaafi baar reuse kiya jaata hai, toh global traffic kaafi kam ho jaata hai jabki compute same rehta hai. Slow pantry ke fewer trips.

Edge cases

Kya hota hai occupancy ka agar ek kernel zero shared memory aur bahut kam registers use kare?

Tab registers/shared memory tumhe limit nahi karte — occupancy hardware ceilings se cap ho jaati hai (max warps ya max blocks per SM). Tum 100% hit kar sakte ho lekin aur nahi; hamesha koi na koi binding constraint hoti hai.

Edge case: ek warp ke saare 32 threads ek exact same global address padhein. Coalesced hai ya nahi?

Fully coalesced — ek segment, ek transaction, saare threads ko broadcast. Same-address global memory ke liye best case hai, jaise woh (broadcast ke roop mein) shared memory ke liye bhi hai.

Edge case: ek warp shared memory access karta hai saare 32 threads bank 0 hit karte hue. ==Serialization factor

S

== kya hai, aur yahan iska value kya hai?

S

kisi bhi single bank par land karne wale threads ki maximum number hai — formally

S = max_{i} (threads hitting bank i)

. Yahan saare 32 bank 0 par pile up hain, toh

S = 32

aur effective latency 32 se multiply ho jaati hai — worst possible shared-memory pattern.

Edge case: tumhari kernel registers "local memory" mein spill karti hai. Woh data physically kahan jaata hai, aur yeh dangerous kyun hai?

Global memory mein (L1/L2 backed), toh ek supposedly fast local variable ab hundreds of cycles per access cost karta hai. Yeh dangerous hai kyunki naam location chhupa leta hai — tum sochte ho yeh on-chip hai lekin nahi hai.

Edge case: ek thread block mein sirf 1 warp (32 threads) hai. Kya yeh memory latency achhe se hide kar sakta hai?

Ache se nahi. Sirf ek warp ke saath stall hone par switch karne ke liye kuch nahi hai, toh SM poori memory latency ke dauran idle rehta hai. Latency hiding ke liye bahut saare concurrent warps chahiye, isliye chhote blocks underperform karte hain.

Edge case: matrix column-major store hai lekin tum ise row-by-row consecutive threads se padh rahe ho. Coalesced hai?

Nahi — consecutive threads rows ke across step karte hain, jo column-major layout mein N elements apart hain, toh har ek alag segment mein land karta hai. Layout aur access pattern agree karni chahiye; wahi code row-major par coalesce hota hai aur column-major par scatter ho jaata hai.

Recall Quick self-test

Ek line ek idea — right side chhupao, jawab do, reveal karo. Registers fast per-access hain lekin limited pot ::: thread pichhe zyada → fewer warps → worse latency hiding. Broadcast vs bank conflict ::: same address broadcast karta hai (free); alag words same bank serialize karte hain. bank(a) formula ::: floor(a/4) mod 32 — word index 32 banks ke around wrapped. Serialization factor S ::: max threads hitting one bank; effective latency times S. Coalescing judged on ::: poore warp ka segment footprint, reads AUR writes dono. Occupancy ek ::: means hai latency hide karne ka, blindly maximize karne ka goal nahi. "Local memory" physically rehti hai ::: global memory mein (register spills), toh yeh slow hai.

Parent par wapas jaao: 6.2.7 Memory hierarchy.