Occupancy and latency hiding

6.2.10 · Hardware › GPU Architecture

Occupancy Kya Hai?

Occupancy kyun matter karta hai?

DRAM se memory reads ~400-800 cycles lete hain
Arithmetic operations ~4-10 cycles lete hain (dependencies ke saath bhi)
Agar sirf 1 warp chal raha hai, toh SM in waits ke dauran stall karta hai → wasted cycles
32+ warps ke saath, scheduler har cycle ek ready warp pe switch karta hai → koi stalls nahi

GPU ye kaise achieve karta hai?

Context is free: Har warp ke registers on-chip rehte hain. Switching cost ~0 cycles.
Warp scheduler: Har cycle ek ready warp (koi data dependency nahi, koi memory wait nahi) choose karta hai.
Overlapping execution: Jab Warp A ld.global ka wait karta hai, Warps B, C, D compute karte hain.

Maano ek memory operation $L = 400$ cycles leta hai. Ek SM 1 warp-instruction per cycle issue kar sakta hai (single scheduler ke liye; modern SMs mein 4 schedulers hote hain, lekin simple shuru karte hain).

Goal: Pipeline full rakho—har cycle 1 instruction issue karo.
Problem: Jo warp memory load issue karta hai use result ready hone se pehle $L$ cycles wait karna padta hai.
Solution: Us wait ke dauran issue karne ke liye $L$ doosre warp-instructions ready rakho.

Isliye: $Warps needed \geq \frac{L}{Instructions per Warp per Cycle}$

Agar har warp scheduled hone par 1 instruction/cycle issue karta hai: $Warps \geq L$

$L = 400$ cycles ke liye → 1 scheduler ke saath latency fully hide karne ke liye 400 warp-instructions in flight chahiye. 4 schedulers aur pipelined execution wale modern SMs is concurrency ka bahut saara hissa har cycle supply karte hain. Practice mein, 16-32 active warps zyaadatar kernels ke liye kaafi hote hain kyunki:

Har warp mein multiple instructions in flight (ILP)
Compute aur memory ops ka mix
4 schedulers consecutive cycles mein alag-alag warps se issue karte hain (SM-wide MLP)

32 threads per warp kyun? Historical reason: SIMD width hardware complexity vs. parallelism ko balance karta hai. 32 lanes matlab 32-wide vector ALUs, jo transistor budgets aur memory coalescing patterns mein fit hote hain (cache lines 128B = 32× 4-byte floats hoti hain).

Occupancy Ko Limit Karne Wale Factors

Example (A100 GPU):

Max warps/SM: 64 (2048 threads)
Register file: 65,536 registers/SM
Shared memory: 164 KB/SM (ya L1 cache enabled ke saath 100 KB)
Max blocks/SM: 32

Scenario 1: Kernel 64 registers/thread use karta hai, koi shared memory nahi.

Registers per warp: $64 \times 32 = 2048$
Possible warps: $⌊ 65536/2048 ⌋ = 32$ warps
Occupancy: 32/64 = 50%

Scenario 2: Kernel 32 registers/thread, 48 KB shared mem/block, 256 threads/block (8 warps) use karta hai.

Register limit: $⌊ 65536/ (32 \times 32)⌋ = 64$ warps ✓
Shared mem limit: $⌊ 164/48 ⌋ = 3$ blocks → $3 \times 8 = 24$ warps
Occupancy: 24/64 = 37.5%

Ye kyun matter karta hai? Kam occupancy → latency hide karne ke liye kam warps → zyada stalls → kam throughput.

Step-by-Step Reasoning ke Saath Examples

Worked example Example 1: Memory-Bound Kernel

Ek kernel global memory se 1 float4 (16 bytes) per thread read karta hai, 10 FLOPS karta hai, 1 float write karta hai. Load latency ≈ 400 cycles; FLOP latency ≈ 4 cycles. Step 1: Arithmetic intensity compute karo.

Bytes moved: $16 + 4 = 20$ bytes
FLOPs: $10$
AI: $10/20 = 0.5$ FLOP/byte Step 2: Per thread critical-path memory time.
Load: ~400 cycles critical path par (result compute karne ke liye chahiye)
Store: asynchronously retire hota hai—global stores buffer hote hain aur background mein drain hote hain, isliye hum critical path mein ~400 cycles ka store latency nahi jodte
Compute: 10 FLOPs at ~1 FLOP/cycle per thread → ~10 cycles
Critical-path latency ≈ 400 cycles (load se dominate hota hai) Step 3: Latency hiding requirement.
~400 cycles ki load latency overlap karne ke liye enough concurrency chahiye
Naïve single-scheduler estimate: $400/ (IPC per warp)$ warp-instructions in flight

Naïve number overshoot kyun karta hai: Concurrency SM-wide provide hoti hai, per warp nahi. Ek single warp ek time mein ek memory instruction issue karta hai (jo uncoalesced hone par kuch transactions mein split ho sakta hai)—ye khud 32 independent outstanding requests sustain nahi karta. Balki, SM ke 4 schedulers kaafi saare warps ko loads issue karte rakhte hain, aur memory system inhe parallel mein service karta hai (SM-wide memory-level parallelism). ILP aur 4 schedulers ke saath, 16-32 active warps usually load latency hide karne ke liye enough outstanding requests supply karte hain.

Result: 50% occupancy (32 warps) zyaadatar latency hide kar deta hai. 100% (64 warps) tak jaana yahan diminishing returns deta hai.

Worked example Example 2: Compute-Bound Kernel

Matrix multiply tile: har thread 256 FMA operations karta hai, shared memory se 8 bytes read karta hai (latency ~30 cycles).

Step 1: Compute dominate karta hai.

256 FMAs at ~1 FMA/cycle (pipelined) → 256 cycles per thread
Shared mem reads: 30 cycles Step 2: Latency hiding.
30 cycles / (high ILP) → shared mem latency hide karne ke liye sirf ~4-8 warps chahiye
Compute naturally pipelining ke through overlap hoti hai

Result: Occupancy 25% (16 warps) aksar kaafi hoti hai. 100% tak badhane se registers waste hote hain bina performance gain ke.

Ye step kyun? Compute-bound kernels ALU throughput se limited hote hain, scheduling se nahi. High occupancy tabhi help karta hai jab scheduling bubbles hoon, jo compute-bound code loop unrolling aur ILP ke through avoid karta hai.

Common Mistakes

Common mistake Mistake 2: "Latency Hiding ke liye Warps = Latency Chahiye"

Ye sahi kyun lagta hai: Little's Law kehta hai $Concurrency = L \times Throughput$ . Agar $L = 400$ hai, toh 400 warp-instructions in flight chahiye.

Steel-man: Math directionally sahi hai—tumhe bahut saara outstanding work chahiye. Lekin do nuances ise soft karte hain:

Instruction-level parallelism (ILP): Har warp ke paas multiple independent instructions in flight ho sakti hain, jo har warp ki taraf se zyada concurrency contribute karta hai.
SM-wide memory-level parallelism (MLP) + 4 schedulers: Pura SM ek saath kaafi saare memory requests outstanding rakhta hai, isliye tumhe kam distinct warps chahiye.

Fix (careful reasoning): ILP, MLP, aur multiple schedulers tumhari di jaane wali concurrency badhate hain, isliye tum kam active warps ke saath required $L \times Throughput$ tak pahunch jaate ho. Ye kisi bhi single load ki intrinsic latency $L$ reduce nahi karte—wo ~400 cycles hi rehti hai. Ye likhna common galti hai $L_{eff} = L / (ILP \times MLP \times Schedulers)$ ; wo Little's Law ko galat apply karta hai. Balki: $Warps needed \approx \frac{L \times Throughput}{Concurrency per warp (ILP) \times ( schedulers )}$ Empirically ye typical kernels ke liye 16-32 warps par land karta hai.

Measuring and Optimizing

Tools:

CUDA Occupancy API: cudaOccupancyMaxActiveBlocksPerMultiprocessor()
Nsight Compute: Achieved occupancy, stall reasons (memory, ALU, sync) dikhata hai.
Heuristic: Kernel ko alag-alag block sizes (128, 256, 512 threads) ke saath run karo. Performance vs. theoretical occupancy plot karo.

Optimization Strategies:

Limiter	Fix	Trade-off
Registers	Live variables kam karo, `__restrict__` use karo, compiler flag	Computation badh sakti hai
Shared Memory	Chhota tile karo, L1 cache use karo (`cudaFuncSetCacheConfig`)	Kam data reuse
Block Size	Threads/block badhao (256→512)	Grid dim mein kam flexibility
Synchronization	`__syncthreads()` kam karo, warp-level primitives use karo	Code complexity
Divergence	Branching avoid karo; predication use karo ya kernels split karo	Zyada instructions

Example: Agar Nsight "Warp was stalled waiting for memory" dikhaye, lekin occupancy 30% hai, try karo:

Block size 512 threads tak badhao → 50% occupancy.
Agar registers limit kar rahe hain, precision reduce karo (FP32 → FP16 jahan safe ho).
Agar phir bhi stalling ho, kernel fundamentally memory-bound hai—memory access patterns optimize karo (coalescing, prefetching).

Active Recall

Recall 12-Saal-Ke-Bachche Ko Samjhao

Socho tum homework kar rahe ho. Har question padne mein 2 minute lagte hain, solve karne mein 10 second. Agar sirf 1 question hai, toh tum 2 minute idle baithe rehte ho uska wait karte. Boring!

Lekin agar 12 questions hain, toh tum Question 1 padhte ho, phir jab book Question 2 "load" kar rahi hoti hai, tum Question 1 solve karte ho. Tum kabhi wait nahi karte—hamesha kuch na kuch karte rehte ho. (Question 1 padhne mein phir bhi poore 2 minute lagte hain—tum bas us time ko doosre kaam se bharte ho, book ko faster load nahi karte.)

GPU ke paas 1000 questions aur 1000 brains hone jaisa hai. Har brain (thread) ek question pakadta hai. Jab ek brain "book" (memory) load hone ka wait karta hai, 999 doosre solve karte rehte hain. Jab tak tum cycle back karte ho, pehle brain ki book ready hoti hai.

Occupancy hai: "Tumhare kitne brains busy hain?" Agar 1000 mein se 500 brains ke paas questions hain, occupancy 50% hai. Jitne zyada busy brains, utna kam waiting.

Connections

6.2.1-SM-Architecture – Occupancy SM resources (regs, shared mem, warp slots) par depend karta hai
6.2.8-Warp-Scheduling – Scheduler jo warps ke beech switch karta hai latency hide karne ke liye
6.2.9-Memory-Coalescing – High occupancy uncoalesced access ki cost amplify karta hai
6.2.11-Tensor-Cores – Tensor Core ops ke fixed warp counts hote hain; occupancy kam critical hai
7.1.3-Roofline-Model – Occupancy determine karta hai ki tum compute ya memory ceiling hit karte ho
5.3.4-Littles-Law – Concurrency ke through latency hiding ka theoretical foundation

Flashcards

#flashcards/hardware

GPU occupancy kya hai?

SM par active warps aur maximum possible warps ka ratio, usually percentage mein express kiya jaata hai (jaise 32/64 = 50%).

High occupancy latency kyun hide karta hai?

Kaafi saare active warps ke saath, warp scheduler hamesha ek ready warp execute karne ke liye dhundh sakta hai jab doosre memory ya ALU results ka wait kar rahe hote hain, stalls eliminate ho jaate hain.

Occupancy limit karne wale teen main factors kya hain?

(1) Registers per thread, (2) Shared memory per block, (3) Maximum blocks per SM. Occupancy in constraints ka minimum hoti hai.

Kya warps add karne se single memory load ki latency kam hoti hai?

Nahi. Concurrency (Little's Law) tumhe kaafi saare loads overlap karne deti hai taaki SM idle na ho, lekin har load ki intrinsic ~400-cycle latency unchanged rehti hai.

Little's Law se latency hiding requirement derive karo

Latency

L

cycles hide karne ke liye throughput

T

ops/cycle par,

L \times T

operations in flight chahiye. 400-cycle memory latency aur 1 warp-instruction/cycle ke liye, ~400 warp-instructions outstanding chahiye (lekin ILP + SM-wide MLP + 4 schedulers distinct warps ko 16-32 tak reduce kar dete hain).

100% occupancy hamesha optimal kyun nahi hota?

~50% ke baad, gains plateau ho jaate hain. High occupancy cache pressure aur register spills badha sakta hai, performance hurt kar sakta hai. Goal zero stalls hai, maximum threads nahi.

A100 ka warps per SM mein maximum occupancy kitna hai?

64 warps/SM (2048 threads, kyunki har warp = 32 threads).

Agar ek kernel A100 (65,536 regs/SM) par 80 registers/thread use kare, toh occupancy kya hogi?

⌊ 65536/ (80 \times 32)⌋ = 25

warps → Occupancy = 25/64 ≈ 39%.

Ek single warp kitne outstanding memory requests sustain karta hai?

Ek time par ek memory instruction (jo uncoalesced hone par kuch transactions mein split ho sakti hai). Concurrency SM-wide kaafi warps mein aati hai, ek warp se 32 independent requests nahi aate.

Kya global stores apni poori latency critical path mein add karte hain?

Nahi. Global stores buffer hote hain aur asynchronously retire hote hain (background mein drain hote hain), isliye tum normally store latency ko critical-path memory time mein nahi jodte.

Low occupancy performance kyun hurt nahi karta?

Jab kernel compute-bound hota hai high instruction-level parallelism (ILP) ke saath, ya jab achieved throughput pehle se hardware saturate kar raha hota hai bawajood kam active warps ke.