11.5b Latency and bandwidth at each level — why kernel optimization targets L1/shared¶
🧠 Context Introduction¶
When you write a CUDA kernel or optimize an AI workload on an NVIDIA GPU, you are essentially managing a memory hierarchy — a pyramid of storage levels, each with different speeds, sizes, and costs. Understanding the latency (how fast data can be accessed) and bandwidth (how much data can be moved per second) at each level is critical. This knowledge explains why kernel optimization almost always targets L1 cache and shared memory — the fastest, most engineer-controlled memory tiers closest to the compute cores.
⚙️ The GPU Memory Hierarchy — A Quick Overview¶
The GPU memory hierarchy, from fastest/smallest to slowest/largest:
- Registers — Private to each thread, zero-latency, very small (256 KB per SM typically)
- L1 Cache / Shared Memory — On-chip, configurable, low latency (~30 cycles), 48–128 KB per SM
- L2 Cache — On-chip, shared across SMs, moderate latency (~200 cycles), several MB
- HBM (High Bandwidth Memory) — Off-chip, high latency (~400–800 cycles), large capacity (16–80 GB)
📊 Latency and Bandwidth at Each Level¶
| Memory Level | Typical Latency (cycles) | Typical Bandwidth | Size per SM / GPU | Engineer Control |
|---|---|---|---|---|
| Registers | 0–1 | Unlimited (per thread) | 256 KB per SM | Compiler-managed |
| L1 / Shared Memory | ~30 | ~10 TB/s (aggregate) | 48–128 KB per SM | Explicit (shared) or implicit (L1) |
| L2 Cache | ~200 | ~4–8 TB/s | 4–8 MB (GPU-wide) | Hardware-managed |
| HBM (VRAM) | ~400–800 | ~1–2 TB/s (e.g., A100: 2 TB/s) | 16–80 GB | Explicit (cudaMemcpy) |
Key insight: The gap between L1/shared memory and HBM is enormous — 10–20x in latency and 5–10x in bandwidth. Every time a thread reads from HBM instead of L1/shared, it stalls for hundreds of cycles.
📊 Visual Representation: Memory Latency vs. Bandwidth Performance Impact¶
This flowchart demonstrates how latency (the delay before data transfer begins) and bandwidth (the data transfer rate) limit performance.
🕵️ Why Kernel Optimization Targets L1 / Shared Memory¶
1. 🚀 Latency Reduction¶
- Accessing HBM takes 400–800 cycles. During that time, the SM can do nothing else for that thread.
- Accessing L1/shared takes ~30 cycles — a 10–20x speedup per memory operation.
- By keeping frequently reused data in shared memory, you eliminate repeated long-latency HBM accesses.
2. 📈 Bandwidth Amplification¶
- HBM bandwidth is shared across all SMs. If every thread constantly reads from HBM, you saturate the bus quickly.
- L1/shared memory bandwidth is per-SM and much higher (aggregate ~10 TB/s vs HBM's ~2 TB/s).
- Using shared memory effectively multiplies your effective bandwidth — you load data once from HBM, then reuse it many times from shared memory.
3. 🧩 Data Reuse Patterns (The "Tiling" Strategy)¶
- Many AI kernels (matrix multiply, convolution, attention) have high data reuse.
- Example: In a matrix multiply, each element of input matrices is used multiple times.
- Without shared memory: Each reuse requires a new HBM read.
- With shared memory: Load a tile (block) into shared memory once, then all threads in the block reuse it from L1/shared — reducing HBM traffic by the tile size factor.
4. 🔧 Engineer Control vs Hardware Control¶
- L1 cache is hardware-managed (you can't explicitly control what stays).
- Shared memory is explicitly managed — you decide exactly what data to load and when.
- This gives engineers fine-grained control over the memory access pattern, which is essential for achieving peak performance.
🛠️ Practical Example: Matrix Multiply Optimization¶
Consider a simple matrix multiply kernel:
Without shared memory (naive): - Each thread reads elements from global memory (HBM) for every multiplication. - Latency: 400–800 cycles per read. - Bandwidth: Quickly saturates HBM bus.
With shared memory (tiled): - Each thread block loads a tile of matrix A and matrix B into shared memory. - All threads in the block reuse that tile from shared memory (30-cycle latency). - Result: 10–20x fewer HBM reads, dramatically higher performance.
Typical optimization flow:
1. Identify data that is reused across threads in a block.
2. Load that data cooperatively from HBM into shared memory.
3. Synchronize threads (using __syncthreads()).
4. Perform computations using shared memory.
5. Write final results back to HBM.
📌 Summary — Key Takeaways for New Engineers¶
- Latency and bandwidth are not uniform — the GPU memory hierarchy has huge speed differences.
- L1/shared memory is the sweet spot — it's fast enough to avoid stalls, large enough to hold working sets, and engineer-controlled.
- Kernel optimization targets L1/shared because:
- It reduces latency by 10–20x compared to HBM.
- It amplifies effective bandwidth by enabling data reuse.
- It gives you explicit control over memory access patterns.
- Always ask yourself: "Can I load this data once into shared memory and reuse it multiple times?" If yes, you have an optimization opportunity.
🔗 Further Learning Path¶
- Understand the CUDA memory model (registers, shared, global, constant, texture).
- Practice tiling in matrix multiply or convolution kernels.
- Use NVIDIA Nsight Compute to profile memory stalls and identify L1/shared memory opportunities.
- Study bank conflicts in shared memory — they can reduce performance if not handled.
💡 Rule of thumb for new engineers: If your kernel spends more than 50% of its time waiting for memory (memory-bound), the first optimization to try is moving frequently accessed data from HBM to shared memory.