Occupancy and latency hiding

What is Occupancy?

WHY does occupancy matter?

Memory reads from DRAM take ~400-800 cycles
Arithmetic operations take ~4-10 cycles (even with dependencies)
If you have only 1 warp running, the SM stalls during these waits → wasted cycles
With 32+ warps, the scheduler switches to a ready warp every cycle → no stalls

HOW does the GPU achieve this?

Context is free: Each warp's registers stay on-chip. Switching costs ~0 cycles.
Warp scheduler: Picks a ready warp (no data dependency, no memory wait) each cycle.
Overlapping execution: While Warp A waits for ld.global, Warps B, C, D compute.

Latency Hiding: The Math

Derivation from First Principles:

Suppose a memory operation takes $L = 400$ cycles. An SM can issue 1 warp-instruction per cycle (for a single scheduler; modern SMs have 4 schedulers, but let's start simple).

Goal: Keep the pipeline full—issue 1 instruction every cycle.
Problem: A warp that issues a memory load must wait $L$ cycles before its result is ready.
Solution: Have $L$ other warp-instructions ready to issue during that wait.

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

If each warp issues 1 instruction/cycle when scheduled: $\text{Warps} \geq L$

For $L = 400$ cycles → need 400 warp-instructions in flight to fully hide latency with 1 scheduler. Modern SMs with 4 schedulers and pipelined execution supply much of this concurrency per cycle. In practice, 16-32 active warps suffice for most kernels due to:

Multiple instructions in flight per warp (ILP)
Mix of compute and memory ops
4 schedulers issuing from different warps in consecutive cycles (SM-wide MLP)

WHY 32 threads per warp? Historical: SIMD width balances hardware complexity vs. paralelism. 32 lanes means 32-wide vector ALUs, which fit transistor budgets and memory coalescing patterns (cache lines are 128B = 32× 4-byte floats).

Factors Limiting Occupancy

Formula Occupancy Constraints

First compute the resource-limited active warps per SM, then divide by the hardware warp limit. Occupancy is the minimum of these (all floored to whole warps/blocks): $W_{\text{regs}} = \left\lfloor \frac{\text{Regs per SM}}{\text{Regs per Thread} \times 32} \right\rfloor$ $W_{\text{smem}} = \left\lfloor \frac{\text{Shared Mem per SM}}{\text{Shared Mem per Block}} \right\rfloor \times \text{Warps per Block}$ $W_{\text{blocks}} = \min(\text{Blocks Launchable per SM},\ \text{Max Blocks per SM}) \times \text{Warps per Block}$ $\text{Occupancy} = \frac{\min(W_{\text{regs}},\, W_{\text{smem}},\, W_{\text{blocks}},\, W_{\max})}{W_{\max}}$ where $W_{\max}$ is max warps per SM (e.g., 64). The active-blocks limit is set by how many blocks fit in resources, then capped by the hardware max blocks per SM—not the other way around.

Example (A100 GPU):

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

Scenario 1: Kernel uses 64 registers/thread, no shared memory.

Registers per warp: $64 \times 32 = 2048$
Warps possible: $\lfloor 65536 / 2048 \rfloor = 32$ warps
Occupancy: 32/64 = 50%

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

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

WHY does this matter? Lower occupancy → fewer warps to hide latency → more stalls → lower throughput.

Examples with Step-by-Step Reasoning

Worked example Example 1: Memory-Bound Kernel

A kernel reads 1 float4 (16 bytes) per thread from global memory, does 10 FLOPS, writes 1 float. Load latency ≈ 400 cycles; FLOP latency ≈ 4 cycles. Step 1: Compute arithmetic intensity.

Bytes moved: $16 + 4 = 20$ bytes
FLOPs: $10$
AI: $10 / 20 = 0.5$ FLOP/byte Step 2: Critical-path memory time per thread.
Load: ~400 cycles on the critical path (result is needed to compute)
Store: retires asynchronously—global stores are buffered and drain in the background, so we do not add ~400 cycles of store latency to the critical path
Compute: 10 FLOPs at ~1 FLOP/cycle per thread → ~10 cycles
Critical-path latency ≈ 400 cycles (dominated by the load) Step 3: Latency hiding requirement.
Need enough concurrency to overlap ~400 cycles of load latency
Naïve single-scheduler estimate: $400 / (\text{IPC per warp})$ warp-instructions in flight

WHY the naïve number overshoots: Concurrency is provided SM-wide, not per warp. A single warp issues one memory instruction at a time (which may split into a few transactions if uncoalesced)—it does not sustain 32 independent outstanding requests by itself. Instead, the SM's 4 schedulers keep many warps issuing loads, and the memory system services these in parallel (SM-wide memory-level parallelism). With ILP and 4 schedulers, 16-32 active warps usually supply enough outstanding requests to hide the load latency.

Result: Occupancy of 50% (32 warps) hides most latency. Going to 100% (64 warps) gives diminishing returns here.

Common Mistakes

Common mistake Mistake 2: "Latency Hiding Requires Warps = Latency"

Why it feels right: Little's Law says $\text{Concurrency} = L \times \text{Throughput}$ . If $L=400$ , need 400 warp-instructions in flight.

Steel-man: The math is directionally correct—you do need lots of outstanding work. But two nuances soften it:

Instruction-level parallelism (ILP): Each warp can have multiple independent instructions in flight, contributing more concurrency per warp.
SM-wide memory-level parallelism (MLP) + 4 schedulers: The SM as a whole keeps many memory requests outstanding, so you need fewer distinct warps.

The Fix (careful reasoning): ILP, MLP, and multiple schedulers increase the concurrency you supply, so you reach the required $L \times \text{Throughput}$ with fewer active warps. They do NOT reduce the intrinsic latency $L$ of any single load—that stays ~400 cycles. It is a common error to write $L_{\text{eff}} = L / (\text{ILP} \times \text{MLP} \times \text{Schedulers})$ ; that misapplies Little's Law. Instead: $\text{Warps needed} \approx \frac{L \times \text{Throughput}}{\text{Concurrency per warp (ILP)} \times (\text{schedulers})}$ Empirically this lands at 16-32 warps for typical kernels.

Measuring and Optimizing

Tools:

CUDA Occupancy API: cudaOccupancyMaxActiveBlocksPerMultiprocessor()
Nsight Compute: Shows achieved occupancy, stall reasons (memory, ALU, sync).
Heuristic: Run kernel with different block sizes (128, 256, 512 threads). Plot performance vs. theoretical occupancy.

Optimization Strategies:

Limiter	Fix	Trade-off
Registers	Reduce live variables, use `__restrict__`, compiler flag	May increase computation
Shared Memory	Tile smaller, use L1 cache (`cudaFuncSetCacheConfig`)	Less data reuse
Block Size	Increase threads/block (256→512)	Less flexibility in grid dim
Synchronization	Minimize `__syncthreads()`, use warp-level primitives	Code complexity
Divergence	Avoid branching; use predication or split kernels	More instructions

Example: If Nsight shows "Warp was stalled waiting for memory," but occupancy is 30%, try:

Increase block size to 512 threads → 50% occupancy.
If registers limit, reduce precision (FP32 → FP16 where safe).
If still stalling, the kernel is fundamentally memory-bound—optimize memory access patterns (coalescing, prefetching).

Active Recall

Recall Explain to a 12-Year-Old

Imagine you're doing homework. Each question takes you 2 minutes to read, 10 seconds to solve. If you only have 1 question, you sit idle for 2 minutes waiting for it. Boring!

But if you have 12 questions, you read Question 1, then while the book "loads" Question 2, you solve Question 1. You never wait—you're always working on something. (Reading Question 1 still takes the full 2 minutes—you just fill that time with other work, you don't make the book load faster.)

A GPU is like having 1000 questions and 1000 brains. Each brain (thread) grabs a question. While one brain waits for the "book" (memory) to load, 999 others are solving. By the time you cycle back, the first brain's book is ready.

Occupancy is: "How many of your brains are busy?" If 500 out of 1000 brains have questions, occupancy is 50%. The more busy brains, the less waiting.

Connections

6.2.1-SM-Architecture – Occupancy depends on SM resources (regs, shared mem, warp slots)
6.2.8-Warp-Scheduling – The scheduler that switches between warps to hide latency
6.2.9-Memory-Coalescing – High occupancy amplifies the cost of uncoalesced access
6.2.11-Tensor-Cores – Tensor Core ops have fixed warp counts; occupancy less critical
7.1.3-Roofline-Model – Occupancy determines if you hit compute or memory ceiling
5.3.4-Littles-Law – Theoretical foundation for latency hiding via concurrency

Flashcards

#flashcards/hardware

What is GPU occupancy?

The ratio of active warps to maximum possible warps per SM, usually expressed as a percentage (e.g., 32/64 = 50%).

Why does high occupancy hide latency?

With many active warps, the warp scheduler can always find a ready warp to execute while others wait for memory or ALU results, eliminating stalls.

What are the three main factors limiting occupancy?

(1) Registers per thread, (2) Shared memory per block, (3) Maximum blocks per SM. Occupancy is the minimum of these constraints.

Does adding warps reduce the latency of a single memory load?

No. Concurrency (Little's Law) lets you overlap many loads so the SM never idles, but the intrinsic ~400-cycle latency of each load is unchanged.

Derive the latency hiding requirement from Little's Law

To hide latency

L

cycles at throughput

T

ops/cycle, need

L \times T

operations in flight. For 400-cycle memory latency and 1 warp-instruction/cycle, need ~400 warp-instructions outstanding (but ILP + SM-wide MLP + 4 schedulers reduce distinct warps to 16-32).

Why is 100% occupancy not always optimal?

Beyond ~50%, gains plateau. High occupancy can increase cache pressure and register spills, hurting performance. The goal is zero stalls, not maximum threads.

What is the A100's maximum occupancy in warps per SM?

64 warps/SM (2048 threads, since each warp = 32 threads).

If a kernel uses 80 registers/thread on an A100 (65,536 regs/SM), what's the occupancy?

\lfloor 65536 / (80 \times 32) \rfloor = 25

warps → Occupancy = 25/64 ≈ 39%.

How many outstanding memory requests does a single warp sustain?

One memory instruction at a time (which may split into a few transactions if uncoalesced). Concurrency comes SM-wide across many warps, not from one warp issuing 32 independent requests.

Do global stores add their full latency to the critical path?

No. Global stores are buffered and retire asynchronously (drain in the background), so you normally don't add store latency to critical-path memory time.

When does low occupancy NOT hurt performance?

When the kernel is compute-bound with high instruction-level parallelism (ILP), or when achieved throughput already saturates the hardware despite fewer active warps.

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, GPU ka core magic ye hai ki wo slow memory ka wait nahi karta. Jab ek warp (32 threads ka group) memory se data lane ke liye ruka hota hai—jo ki 400-800 cycles le sakta hai—tab GPU dusre ready warps ko chalata rehta hai. Isko bolte hain latency hiding. Aur occupancy simply ye batata hai ki tum apne SM (Streaming Multiprocessor) ki capacity ka kitna hissa use kar rahe ho—yaani active warps divided by maximum possible warps. Restaurant kitchen wali example yaad rakho: agar tumhare paas 100 orders hain to har chef busy rehta hai, koi idle nahi khada, aur ingredients ka wait ho bhi to kaam ruकता nahi.

Ab ye matter kyun karta hai? Kyunki context switching GPU mein bilkul free hai—har warp ke registers on-chip padhe rehte hain, to warp badalne mein zero cycles lagte hain. Warp scheduler har cycle ek aisa warp uthata hai jo ready ho (jiska data aa gaya ho ya jo waiting nahi hai). Isliye agar tumhare paas kaafi warps hain, to SM kabhi stall nahi hoti aur pipeline hamesha bhari rehti hai. Little's Law bhi yahi kehta hai: latency $L$ ko hide karne ke liye tumhe $L \times$ throughput jitni concurrency chahiye. Practically 16-32 active warps zyadatar kernels ke liye enough hote hain, kyunki multiple instructions per warp (ILP) aur 4 schedulers milkar concurrency provide kar dete hain.

Lekin ek important cheez samajhna—occupancy tumhare registers per thread, shared memory per block, aur blocks per SM jaisi resources se limit hoti hai. Agar tumhara kernel bahut zyada registers use karta hai, to kam warps fit honge aur occupancy gir jayegi. Isliye final occupancy in saare constraints ka minimum hoti hai. Yaad rakho, zyada concurrency single operation ki latency ko chhota nahi karti—wo bas overlap karne deti hai taaki SM idle na baithe. Ye samajhna zaroori hai kyunki jab tum performance tune karoge, to tumhe balance karna padega ki registers aur shared memory kitna use karein taaki enough warps chalein aur latency effectively hide ho sake.

What is Occupancy?

Latency Hiding: The Math

Factors Limiting Occupancy

Examples with Step-by-Step Reasoning

Common Mistakes

Measuring and Optimizing

Active Recall

Connections

Flashcards

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections