Memory hierarchy (global, shared, registers)

6.2.7 · Hardware › GPU Architecture

Overview

GPU memory hierarchy capacity aur speed ke beech trade-off karta hai, storage ko teen main levels mein organize karta hai: registers (sabse fast, sabse chhota), shared memory (medium speed, chhota), aur global memory (sabse slow, sabse bada). Is hierarchy ko samajhna GPU performance ke liye critical hai kyunki memory access latency computation se 100-1000× zyada lambi ho sakti hai.

Figure — Memory hierarchy (global, shared, registers)

Core Concepts

1. Registers

Key Properties:

Latency: ~1 clock cycle
Capacity: 32-256 KB per SM (saare active threads mein divide hota hai)
Scope: Ek single thread ke liye private
Allocation: Automatic (compiler local variables ko registers mein assign karta hai)

Worked example Register Usage Example

Problem: Aap ek kernel launch karte ho 512 threads/block ke saath ek SM pe jo 65,536 registers rakhta hai. Aapka kernel 40 registers/thread use karta hai. Kitne blocks concurrently run kar sakte hain?

Solution:

Step 1: Registers needed per block $R_{b l oc k} = 512 \times 40 = 20, 480 registers$

Ye step kyun? Capacity check karne se pehle hume total demand jaanni hai.

Step 2: Blocks jo fit hote hain $N_{b l oc k s} = ⌊ \frac{65 , 536}{20 , 480} ⌋ = 3 blocks$

Ye step kyun? Har block ko 20,480 registers chahiye; hum total capacity ko per-block demand se divide karte hain.

Step 3: Active threads $N_{t h r e a d s} = 3 \times 512 = 1, 536 threads$ Agar max 2048 hai, toh occupancy = 1536/2048 = 75%

Key insight: Agar hum register usage 32/thread tak reduce karein, toh 4 blocks (2048 threads, 100% occupancy) fit honge. Chhoti-chhoti register savings bhi matter karti hain.

2. Shared Memory

Key Properties:

Latency: ~30-100 clock cycles (global memory se roughly ~4-13× faster)
Capacity: 16-96 KB per SM (configurable, modern architectures pe aksar L1 cache ke saath reconfigurable)
Scope: Ek thread block ke andar share hoti hai
Access pattern: Banks mein organize hoti hai (typically 32 banks, 4-byte width)

Worked example Shared Memory Example: Tiling Matrix Multiply

Problem: Do 1024×1024 matrices ko multiply karo. Har element ko 1024 multiply-adds chahiye. Naive approach: har thread global memory se 2048 baar load karta hai.

Shared memory ke saath solution: Step 1: Matrices ko 16×16 tiles mein divide karo Kyun? Ek tile shared memory mein fit ho jaata hai (~1 KB float ke liye). Block ke saare threads ek hi tile data reuse karte hain.

Step 2: A ka ek tile aur B ka ek tile shared memory mein load karo (coalesced global reads)

__shared__ float As[16][16];
__shared__ float Bs[16][16];
As[ty][tx] = A[row*N + tileStart + tx];  // All threads cooperate

Ye step kyun? 256 threads, 256 elements ek baar load karte hain, instead of har thread same elements baar baar load kare.

Step 3: Har thread shared tile use karke apna partial result compute karta hai

for(int k=0; k<16; k++)
    Cvalue += As[ty][k] * Bs[k][tx];

Ye step kyun? Ab hum har loaded value ko fast shared memory se 16 baar reuse karte hain.

Step 4: Saare tiles ke liye repeat karo (1024/16 = 64 tiles)

Speedup calculation:

Naive: 2048 global reads/thread × 1M threads = 2B global reads
Tiled: 2×64 tile loads × 256 elements = 32K global reads per block × 4096 blocks = 131M global reads ~15× reduction in global memory traffic!

3. Global Memory

Key Properties:

Latency: 400-800 clock cycles
Capacity: GBs (e.g., modern GPUs pe 8-80 GB)
Scope: Poore device mein global
Bandwidth: 200-1000 GB/s (high throughput lekin high latency)

Worked example Coalescing Example: Transpose

Problem: Ek 4096×4096 matrix ko transpose karo. Naive approach coalesced read karta hai lekin strided write karta hai (ya vice versa). Naive transpose (bad):

out[col*N + row] = in[row*N + col];  // Read coalesced, write strided

Bad kyun? Thread $t$ , out[col*N + t] pe write karta hai. Jab $co l$ ek warp ke andar constant rehta hai, saare threads $N$ apart addresses pe write karte hain → uncoalesced. Better solution: Shared memory tile use karo

Step 1: Shared memory mein coalesced read

__shared__ float tile[32][33];  // Note: 33 avoids bank conflicts
tile[ty][tx] = in[row*N + col];

33 kyun? Agar hum [32][32] use karte, toh saare threads ke liye tile[i][i] access karna same bank hit karta. Padding banks ko shift karta hai.

Step 2: Shared memory mein transpose karo (free hai - bas differently index karo) Step 3: Transposed tile se coalesced write

out[col*N + row] = tile[tx][ty];  // Now coalesced

Result: Dono reads aur writes coalesced hain. Naive version se ~10× speedup.

Memory Hierarchy Summary Table

| Level | Latency (cycles) | Size | Scope | Managed By | |-------|------------------|-------|------------| | Registers | 1 | 32-256 KB/SM | Per-thread | Compiler | | Shared Memory | 30-100 | 16-96 KB/SM | Per-block | Programmer | | L1 Cache | 30-100 | 16-128 KB/SM | Per-SM | Hardware | | L2 Cache | ~200 | 0.5-6 MB | Global | Hardware | | Global Memory | 400-800 | GB | Global | Programmer |

Common Mistakes & Misconceptions

Common mistake Mistake 1: "Zyada registers per thread = faster code"

Ye sahi kyun lagta hai: Registers sabse fast memory hain, toh zyada use karna achha lagta hai.

Ye galat kyun hai: High register usage occupancy (concurrent threads) ko reduce karta hai. GPUs latency ko threads ke beech switch karke hide karte hain. Kam occupancy mein, SM ke paas switch karne ke liye kam threads hote hain → memory waits pe stalls aate hain.

Fix: Register usage aur occupancy ko balance karo. Occupancy measure karne ke liye profilers use karo. Kabhi kabhi thoda slower computation accept karna (kam registers) better latency hiding enable karta hai.

Example: 40 se 32 registers/thread tak jaana individual threads ko 10% slow kar sakta hai lekin occupancy 75% se 100% tak badhaa sakta hai, jo 20% overall speedup deta hai.

Recall Feynman Explanation (Ek 12-Saal-Ke Bachche Ko Samjhao)

Imagine karo tum homework kar rahe ho aur information dhundni hai:

Tumhari desk (registers): Woh paper jis pe tum likh rahe ho aur haath mein pencil. Use karna super fast hai — tumhe reach bhi nahi karna. Lekin sirf ek ya do sheets ki jagah hai.

Tumhara backpack (shared memory): Woh books aur notes jo tum apni desk pe laaye. Grab karne mein ek second lagta hai, lekin room ke across jaane se kaafi fast hai. Tum aur tumhare dost jo ek hi group project pe kaam kar rahe hain ek backpack ki cheezein share kar sakte hain.

Bookshelf (global memory): Room ke doosri taraf. Wahan jaane-aane mein 10 second lagte hain. Lekin usme saikdon books hain.

Smart strategy: Jab tum ek naya problem shuru karo, bookshelf tak ONCE jao, sab kuch jo chahiye le lo, backpack mein bharo, aur jo actively use kar rahe ho usse desk pe le aao. Agar tum har sentence ke liye bookshelf tak jaoge, toh poora din chalne mein guzar jayega aur actually kaam karne ka time nahi milega.

Exactly yehi GPUs karte hain! Woh data ko slow global memory → fast shared memory → super-fast registers mein move karte hain, aur jitna ho sake har level pe use reuse karte hain zyada ke liye jaane se pehle.

Connections

GPU Thread Hierarchy - Thread blocks shared memory scope determine karte hain
Warp Execution Model - Memory coalescing warp-level access patterns pe depend karta hai
Cache Architecture - L1/L2 caches explicit hierarchy ko complement karte hain
Memory Bandwidth Optimization - Har level pe throughput maximize karne ki techniques
CUDA Memory Types - Software abstractions (constant, texture memory)
Memory Latency Hiding - Occupancy performance ke liye kyun matter karta hai
Matrix Multiplication Optimization - Classic case study teeno levels use karke
Roofline Model - Memory vs compute bottlenecks analyze karna

Flashcards

GPU memory hierarchy ke teen main levels fastest se slowest tak kya hain?

Registers (per-thread, ~1 cycle), Shared Memory (per-block, ~30-100 cycles), Global Memory (device-wide, ~400-800 cycles).

Ek SM pe kitne threads concurrently run kar sakte hain, ye kya determine karta hai?

Register usage per thread (occupancy = total_registers / (threads × registers_per_thread)) aur doosre limits jaise shared memory aur thread block size.

Shared memory mein bank conflict kya hota hai?

Jab ek warp ke multiple threads alag-alag addresses access karte hain jo same memory bank pe map karte hain, jisse accesses serialize ho jaate hain aur latency conflicts ki number se multiply ho jaati hai.

Shared memory performance optimization kaise enable karta hai?

Ek block ke threads ko allow karke ki woh slow global memory se data cooperatively ek baar load karein, phir use fast shared memory se kai baar reuse karein, global memory traffic reduce karke.

Memory coalescing kya hai?

Jab ek warp ke consecutive threads consecutive memory addresses access karte hain, GPU ko requests ek single wide memory transaction mein combine karne deta hai instead of multiple narrow ones ke.

Zyada registers per thread use karna performance kyun hurt karta hai?

High register usage occupancy reduce karta hai (kam concurrent threads), jo GPU ki ability ko reduce karta hai ki memory latency ko threads ke beech switch karke hide kare.

Shared memory ka scope kya hai?

Per thread block — sirf ek hi block ke threads same shared memory access kar sakte hain, aur data alag blocks mein visible nahi hota.

Latency ke terms mein shared memory global memory se approximately kitna fast hai?

Roughly 4-13× faster (~30-100 cycles shared ke liye vs ~400-800 cycles global ke liye), 100× nahi.

Registers aur global memory ke beech typical size difference kya hai?

Registers: ~100 KB per SM, Global memory: tens of GB — roughly 100,000× bada capacity aur 400-800× zyada latency.