GPU Cache Hierarchy Explained: What L1, L2, and L3 Actually Do Inside Your GPU

Uday Sharma

12 min read

·

Aug 5, 2026

--

Listen

Share

From registers to HBM — a visual, practical breakdown of how GPUs keep thousands of cores fed, and why NVIDIA and AMD do it differently

If you’ve ever waited on a slow CUDA kernel, watched a training run get “memory-bound” for no obvious reason, or read a paper like FlashAttention and wondered why it’s such a big deal that it “avoids reading and writing to HBM” — you’ve already run into the topic of this article, even if nobody called it by name.

GPUs have a memory hierarchy, just like CPUs do. But it’s built on a completely different philosophy, and understanding it is one of those things that quietly separates engineers who can write GPU code from engineers who can write fast GPU code.

This article breaks down the three (well — two-and-a-half) levels of cache inside a modern GPU: what they are, where they physically sit, how big they are on real chips like the H100 and MI300X, and why AMD and NVIDIA don’t even agree on whether “L3” exists.

Table of Contents

1. Why GPU Caches Matter
2. The Kitchen Analogy
3. The Full Memory Hierarchy at a Glance
4. L1 Cache & Shared Memory: The SM’s Private Workspace
5. L2 Cache: The GPU’s Shared Warehouse
6. L3 / Infinity Cache: The Tier Most GPUs Don’t Have
7. NVIDIA vs AMD: Two Different Philosophies
8. GPU Caches vs CPU Caches
9. Real Numbers: Cache Sizes Across Modern GPUs
10. Seeing It in Code: Why Cache-Aware Kernels Matter
11. Why This Matters for AI Engineers
12. Production Checklist
13. Key Takeaways
14. References

1. Why GPU Caches Matter

A modern GPU can have over 16,000 CUDA cores. Every one of those cores needs a constant stream of data to stay busy. The problem is that the GPU’s main memory — HBM3 or GDDR6X — while extremely fast by RAM standards, is still hundreds of times slower to reach than the compute cores themselves.

If every single core had to fetch every single value directly from VRAM, the cores would spend most of their time waiting, not computing. This is called being memory-bound, and it’s one of the most common performance killers in GPU programming — including in the transformer kernels that power every LLM you’ve used.

Caches exist to solve exactly this problem: keep the data that’s about to be reused as physically close to the compute units as possible, so the cores almost never have to wait on the slow, far-away memory.

2. The Kitchen Analogy

Think of a GPU like an enormous restaurant kitchen with thousands of cooks working at once.

• Registers are what’s literally in a cook’s hands right now — the one ingredient they’re chopping. Nobody else can touch it, and grabbing it takes zero time.
• L1 Cache / Shared Memory is that cook’s personal prep station, shared only with the small team working the same section of the kitchen. It’s fast to reach and the team can pass ingredients to each other directly.
• L2 Cache is the kitchen’s walk-in fridge. Every station in the kitchen can pull from it, and it holds a lot more than any one prep station — but walking there and back takes noticeably longer.
• L3 / Infinity Cache (only in some kitchens) is a big shared pantry connecting multiple kitchen branches in the same building. Not every restaurant has this — some just go straight from the walk-in fridge to the supplier.
• VRAM is the supermarket down the street. It has everything, but every trip costs real time.
• System Memory is the wholesale distributor across town — huge capacity, but you really don’t want to make that trip often.

Keep this picture in mind. Every section below is really just this same idea, in hardware.

3. The Full Memory Hierarchy at a Glance

Notice the dashed border around the L3 / Infinity Cache tier — that’s intentional. Not every GPU has an L3 cache. This is the single most misunderstood part of GPU cache hierarchy, and we’ll spend a full section on it below

4. L1 Cache & Shared Memory: The SM’s Private Workspace

Every NVIDIA GPU is built from repeating units called Streaming Multiprocessors (SMs) — AMD’s equivalent is the Compute Unit (CU), grouped into Workgroup Processors (WGPs). Each SM has its own private, on-chip memory that never has to be shared with any other SM.

On NVIDIA hardware since Volta, this on-chip memory is unusual in a specific way: the L1 cache, the texture cache, and the programmer-controlled “shared memory” are physically the same block of SRAM. CUDA lets you choose, per kernel, how much of that block acts as an automatic hardware cache versus a scratchpad you manage yourself with explicit __shared__ variables.

Key facts about L1 / shared memory:

• Scope: private to one SM (not visible to other SMs)
• Size: as small as 24 KB per SM on Pascal (2016), up to 256 KB per SM on Hopper H100 (2022)
• Latency: roughly 20–30 clock cycles
• Managed by: partly automatic hardware caching, partly explicit programmer control
• Typical use: staging a “tile” of data (like a block of a matrix) that many threads in the same block will read repeatedly

This is the only cache level in the entire hierarchy that CUDA and HIP kernels can address directly. When you write __shared__ float tile[32][32]; in a kernel, you're reaching straight into this layer. That direct control is exactly why shared memory is central to almost every high-performance GPU kernel, including matrix multiplication and attention kernels — more on that in Section 10.

5. L2 Cache: The GPU’s Shared Warehouse

Where L1/shared memory is private to a single SM, L2 cache is shared by every SM on the chip. It’s the GPU’s central meeting point — when one SM writes data that another SM will need, L2 is usually where that handoff happens without having to round-trip through VRAM.

Key facts about L2 cache:

• Scope: shared across the entire GPU die
• Size: 40 MB on A100, 50 MB on H100, and a surprisingly large 72 MB on the consumer RTX 4090
• Latency: roughly 200–300 clock cycles — much slower than L1, but far faster than VRAM
• Managed by: hardware, automatically, though modern architectures expose some manual control

That last point is worth expanding on. Starting with Ampere, NVIDIA added an API — cudaFuncSetAttribute with cudaFuncAttributePreferredSharedMemoryCarveout, along with L2 persistence controls — that lets you hint to the GPU that specific data (like frequently-reused weights) should be kept "pinned" in L2 rather than evicted. For inference workloads that repeatedly touch the same tensors, this is a meaningful lever.

Every modern GPU — NVIDIA or AMD — has an L2 cache. It’s the one universal middle layer of the hierarchy. The disagreement starts at the next level up.

6. L3 / Infinity Cache: The Tier Most GPUs Don’t Have

Here’s the part that trips up a lot of people coming from CPU architecture, where L1/L2/L3 is a settled, universal naming convention.

Most NVIDIA GPUs stop at L2. There is no L3. The H100, A100, and every RTX 40-series card go straight from a chip-wide L2 cache to off-chip HBM or GDDR memory. NVIDIA’s own documentation and tuning guides never mention an “L3” tier for these chips.

AMD, on the other hand, introduced a genuine last-level cache starting with RDNA2 in 2020, which it branded Infinity Cache. Functionally, this is an L3 cache — a large pool of SRAM shared across the whole chip (or across multiple chiplets, in AMD’s newer designs), sitting between L2 and HBM/GDDR.

The clearest example is AMD’s MI300X, a chiplet-based accelerator built for AI workloads. Each of its eight compute dies (XCDs) has its own small 4 MB L2 slice, but all of them share a single 256 MB Infinity Cache — genuinely enormous compared to anything on an NVIDIA chip, and explicitly documented by AMD as “the last level cache of the cache hierarchy.”

So when someone asks “what’s the L3 cache on an NVIDIA GPU,” the honest answer is: there usually isn’t one. The question itself carries a CPU assumption that doesn’t map cleanly onto NVIDIA’s architecture — though it maps very well onto AMD’s.

7. NVIDIA vs AMD: Two Different Philosophies

Put the two diagrams side by side and a clear pattern emerges:

• NVIDIA leans on a large, fast, per-SM L1/shared memory and gives programmers direct manual control over it. The philosophy is: give kernel authors a fast scratchpad they can hand-tune, and back it with a solid but comparatively modest chip-wide L2.
• AMD leans on a much larger shared last-level cache (Infinity Cache) to reduce how often the GPU needs to touch HBM at all — especially valuable on chiplet designs like MI300X, where trips to memory have to cross die boundaries.

Neither approach is strictly “better” — they reflect different bets about where data reuse tends to happen in real workloads, and different constraints from monolithic-die versus chiplet manufacturing.

8. GPU Caches vs CPU Caches

If you’re coming from CPU architecture, a few differences are worth calling out explicitly:

Press enter or click to view image in full size

The biggest conceptual shift: a CPU cache exists mainly to make a single thread run fast. A GPU cache exists to make it possible to keep thousands of threads fed simultaneously — latency to any individual piece of data matters less than sustained throughput to the whole SM.

9. Real Numbers: Cache Sizes Across Modern GPUs

Press enter or click to view image in full size

A few things jump out from this table:

• L1/shared memory capacity per SM roughly 10x’d from Pascal to Hopper — a sign of how much NVIDIA has leaned into giving kernels more on-chip scratchpad.
• The RTX 4090’s 72 MB L2 is actually larger than the H100’s 50 MB L2 — a reminder that “datacenter GPU” doesn’t automatically mean “bigger cache,” since H100 prioritizes raw HBM bandwidth over cache capacity for its target workloads.
• AMD’s Infinity Cache figures (128 MB, then 256 MB) dwarf anything on the NVIDIA side — because it’s doing a different job, standing in for the L3 tier NVIDIA’s design doesn’t have.

10. Seeing It in Code: Why Cache-Aware Kernels Matter

Registers and L2 are managed automatically. But L1/shared memory is the one tier you control directly — and using it well is often the difference between a kernel that’s memory-bound and one that isn’t.

Here’s the classic example: tiled matrix multiplication. Instead of every thread reading directly from slow global memory (VRAM) for every multiplication, threads cooperate to load a small tile into fast shared memory once, then reuse it many times.

#define TILE_SIZE 32

__global__ void tiledMatMul(const float* A, const float* B, float* C, int N) {
    // Shared memory tiles — this lives in L1/shared memory, not VRAM
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < N / TILE_SIZE; ++t) {
        // Each thread loads ONE element into shared memory
        tileA[threadIdx.y][threadIdx.x] = A[row * N + (t * TILE_SIZE + threadIdx.x)];
        tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];

        __syncthreads(); // wait until the whole tile is loaded

        // Now reuse the tile TILE_SIZE times from fast shared memory
        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }

        __syncthreads(); // wait before loading the next tile
    }

    C[row * N + col] = sum;
}

Without tiling, every multiply-add would issue a fresh read from VRAM. With tiling, each value loaded from VRAM gets reused TILE_SIZE times (32 times, here) straight out of L1/shared memory — cutting global memory traffic by roughly the same factor. This single pattern — load once into fast on-chip memory, reuse many times — is the foundation of nearly every high-performance CUDA and HIP kernel, and it's exactly the same idea behind FlashAttention, which tiles attention computation into SRAM specifically to avoid repeated round-trips to HBM.

11. Why This Matters for AI Engineers

If you write CUDA kernels for a living, this is old news. But if you’re an AI engineer who mostly works a layer or two above raw kernels — PyTorch, vLLM, TensorRT-LLM — this hierarchy still shapes performance in ways you’ll run into directly:

• Attention kernels. FlashAttention and its successors are fast specifically because they’re cache-hierarchy-aware — they tile Q/K/V into shared memory instead of materializing the full attention matrix in VRAM.
• KV-cache reuse. During autoregressive decoding, the KV-cache gets re-read on every generated token. Whether that data stays warm in L2 (or Infinity Cache, on AMD) versus falling back to HBM has a real, measurable effect on decode latency.
• Batch size and occupancy tuning. Cranking batch size up too far can push a kernel’s working set out of L1/shared memory and into slower memory tiers — sometimes making a “bigger batch” run slower per-token, not faster.
• Choosing hardware. An MI300X’s 256 MB Infinity Cache versus an H100’s 50 MB L2 isn’t just a spec-sheet number — it reflects two different bets about how much data reuse a workload has, which matters when you’re picking (or renting) hardware for a specific model and batch size.

You don’t need to hand-write CUDA to benefit from this. Just understanding why a fused, tiled kernel outperforms a naive one — and why that gap is bigger on some GPUs than others — will make you a sharper reader of benchmark numbers and a better-informed decision-maker when picking inference hardware.

12. Production Checklist: Writing Cache-Aware GPU Code

• Tile your data. If threads in the same block reuse the same values, load them into __shared__ memory once instead of re-reading global memory repeatedly.
• Watch your shared memory budget. Requesting too much shared memory per block reduces how many blocks can run concurrently on an SM — profile occupancy, don’t just guess.
• Use coalesced memory access patterns. Even before data reaches L1, how threads access global memory determines how efficiently it gets cached at all.
• Consider L2 persistence for hot data (Ampere and newer). If a kernel repeatedly reads the same tensor (e.g., cached weights), cudaFuncAttributePreferredSharedMemoryCarveout and L2 persistence hints can keep it resident longer.
• Prefer fused, tiled kernels over naive ones for anything running at scale. This is the entire premise behind FlashAttention-style kernels — and it applies to any repeated-read workload, not just attention.
• Don’t assume “L3” exists. If you’re reasoning about cache behavior on NVIDIA hardware, remember the hierarchy stops at L2. On AMD, factor in the (much larger) Infinity Cache.
• Profile, don’t guess. Tools like NVIDIA Nsight Compute report actual L1/L2 hit rates per kernel — use them before optimizing blind.

13. Key Takeaways

• GPU caches exist to keep thousands of cores fed without constantly waiting on slow VRAM.
• L1 cache / shared memory is private per-SM, small (up to 256 KB on H100), fast, and the only tier programmers can address directly.
• L2 cache is shared across the whole GPU, larger (40–72 MB on modern chips), and fully automatic.
• L3 / Infinity Cache is not universal. Most NVIDIA GPUs don’t have one; AMD’s RDNA2+ and CDNA3+ chips do, and it can be enormous (up to 256 MB on MI300X).
• NVIDIA and AMD optimize for different things: NVIDIA leans on programmer-controlled shared memory, AMD leans on a very large automatic last-level cache.
• Understanding this hierarchy explains why techniques like FlashAttention and tiled matrix multiplication are fast — and gives you a better mental model for reading GPU benchmarks and choosing inference hardware.

14. References

• NVIDIA Hopper Tuning Guide — docs.nvidia.com/cuda/hopper-tuning-guide
• NVIDIA CUDA C++ Programming Guide — docs.nvidia.com/cuda/cuda-c-programming-guide
• Chips and Cheese, “Nvidia’s H100: Funny L2, and Tons of Bandwidth” — chipsandcheese.com
• Chips and Cheese, “Testing AMD’s Giant MI300X” — chipsandcheese.com
• Chips and Cheese, “AMD’s CDNA 3 Compute Architecture” — chipsandcheese.com
• AMD ROCm Documentation, Device Hardware Glossary — rocm.docs.amd.com
• Wikipedia, “Ampere (microarchitecture)” and “Hopper (microarchitecture)” — en.wikipedia.org
• Dao et al., “FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness” (arXiv)

If you found this helpful, follow for more deep dives into the systems that power modern AI. Claps and comments are always appreciated! 👏