hardNewNVIDIA

What can you rely on from the L2 cache on a GPU? When does it save a kernel, and when does it mislead you?

An H100 has 50 MB of L2, split into two partitions, with a bandwidth several times HBM's. That is enough to change the roofline for a kernel whose working set fits and to do nothing at all for the ones that stream. The arithmetic of what fits, the residency controls, and the two ways the L2 lies to a profile.

Updated Sep 2026 · Grounded in real AI infrastructure interview loops and written to a senior-engineer editorial bar, with every number worked and every diagram hand-built.

20 answers per topic instead of 10, plus saved progress and bookmarks · no cardor unlock all 283 remaining answers · ₹2,000 / $25

The concepts behind this question

Ranked by how closely each one overlaps this question's topic, so the first card is the thing to read if the answer above moved too fast.

Foundational

🧩 GPU & Accelerator Architecture

GPU Memory HierarchyA GPU has four places a byte can live, and they differ by a thousandfold in bandwidth: registers, shared memory on the SM, a chip-wide L2, and HBM off-chip. Almost every kernel optimization is a decision about which level a value is read from and how many times. Knowing the sizes and bandwidths for an H100 cold is what lets you say why a kernel is slow before you profile it.

Advanced

⚡ Kernels & Compilers🔒 Premium

Tiled Matrix MultiplicationA matrix multiply has enough reuse to be compute-bound, but only if the kernel captures that reuse in shared memory and registers instead of re-reading HBM. Tiling is how: a block owns an output tile, streams K-slices of A and B through shared memory, and each thread accumulates a small register tile. It is the live-coding exercise that separates people who know the roofline from people who have climbed it.

Advanced

⚡ Kernels & Compilers🔒 Premium

Shared Memory and Bank ConflictsShared memory is the programmer-managed SRAM inside each SM, split into 32 four-byte banks that serve one word each per cycle. When several lanes of a warp hit the same bank at different addresses the access serializes, and a 32-way conflict makes a shared-memory-bound loop run over ten times slower. Padding, XOR swizzles, cp.async and TMA are the tools that decide whether a tiled kernel gets the bandwidth it staged data for.

Foundational

🧩 GPU & Accelerator Architecture

GPU Execution ModelA GPU hides memory latency with parallelism instead of caches: thousands of threads in flight, scheduled in warps of 32, pinned to streaming multiprocessors that switch between warps for free whenever one stalls. Every performance conversation in an AI infra loop, from occupancy to why decode is slow, rests on this one mechanism.

Explore the full curriculum →

UP NEXT ON YOUR JOURNEY

Next in this trackWhat does ECC on a GPU cost you, and why do you keep it on across a fleet?Next in this trackExplain NUMA and CPU affinity on a GPU node. Why can the wrong socket halve host-to-device bandwidth?Next in this trackCompare A100, H100 and B200 for a 70B serving fleet. Which gives the most tokens per dollar, and where do fp8 and fp4 change the ranking?

FEDITOR'S NOTE

Scored on treating the L2 quantitatively (what fits per wave, what bandwidth it adds), knowing the residency controls exist and what they cost, and recognizing the cases where a good hit rate does not mean a fast kernel.

DISCUSSION · 0

No comments yet — be the first to share your approach.