GPU Kernels Explained

An animated explainer of how a GPU actually executes the maths: the memory hierarchy, the roofline, warps and divergence, coalescing, shared-memory bank conflicts and occupancy, then the kernels built on them: GEMM step by step, reductions and warp shuffles, online softmax and FlashAttention, overlap and split-K, quantised kernels, and the warp-specialised tensor-core pipeline of Hopper and Blackwell (TMA, wgmma → tcgen05, Tensor Memory, FP8 and FP4 with block scaling). Every chapter is built around an animation, and every frame of every animation is computed by a small GPU execution model whose TypeScript port matches its Python reference exactly. Every hardware figure comes from a published source.

It is the fourth of a family of companion sites: the Transformer Decoder Explainer shows one forward pass, LLM Inference Explained shows how a model is served, LLM Architectures Explained shows how the models differ, and this site is the layer underneath: the kernels. They share one design system and link to each other from the header, in four tracks, and the header always shows where you are ("Track › Site"): "Foundations" (Decoder · Training: one small model, end to end), "Hardware" (Kernels · Systolic · Accelerators · Racks: the machines underneath, chip to rack), "Inference" (Serving · Architectures · Numerics · Trade-offs; Serving is LLM Inference Explained) and "Agents", which starts with Agent Harnesses Explained (the loop, tools, context and permissions that turn a model into an agent; then Agent Protocols Explained, MCP and A2A on the wire; then Agent Context Explained, retrieval, memory and context engineering; then Agent Orchestration Explained, graphs, multi-agent patterns and durable execution; then Agent Evals Explained, measuring agents and how sure to be; then Agent Security Explained, defences that keep agents safe). Numerics Explained covers the accuracy side of the quantised kernels in chapter 11, and Systolic Arrays Explained is the other way to build a matrix engine (chapters 7 and 10 link it). Inference Trade-offs Explained measures what the quantised kernels (chapter 11) and the overlap of communication (chapter 10) are worth to a whole serving system.

Live: gpu-kernels-explained.vercel.app

Part of

This project sits in the LLMs hub, next to the Transformer Decoder Explainer, LLM Inference Explained and LLM Architectures Explained. Every chapter links the matching slides of the NVIDIA GPU Architectures and CUDA Programming series.

Chapters

#	Chapter	The animation
01	The memory hierarchy	Data flows HBM → L2 → shared memory → registers → ALUs; pipe widths are bandwidths to scale, packets move at each level's real use. Vector add, two GEMMs, A100 or H100.
02	The roofline	A GEMM's tile size grows from 1 to 128 and its point slides from the memory-bound slope onto the roof; a slider places any intensity.
03	Warps, SIMT and divergence	A warp issues every instruction of both paths of a branch, with lanes masked off; four kinds of condition.
04	Coalescing	Each lane's address maps to 32-byte sectors; stride, alignment and element size are adjustable.
05	Shared-memory bank conflicts	A 32-bank heat map fills request by request, then drains pass by pass; one word of padding fixes the transpose live.
06	Occupancy	Thread blocks land on an SM until warp slots, registers or shared memory run out; sliders for each.
07	GEMM, step by step	Tiles of A and B march into a 16 × 16 × 16 multiply, four ways (naive, shared-memory tiles, register blocking, tensor cores), with byte counters and the intensity they give.
08	Reductions and warp shuffles	64 numbers summed by one block: divergent, strided (bank conflicts), sequential, then warp shuffles; arrows show every add.
09	Softmax and FlashAttention	FlashAttention-2's tile loop with HBM bytes against standard attention's; online softmax rescaling its running sum, checked against the ordinary softmax.
10	Split-K, streams and overlap	A load / compute / store timeline with one, two or three buffers; split-K filling the SMs wave by wave.
11	Quantised kernels	Eight INT4 weights unpacked and dequantised in one register; a layer's time against batch size in four formats.
12	Hopper and Blackwell GEMM	A producer warp's TMA loads fill a ring of shared-memory stages while single-thread MMAs consume them into Tensor Memory, with the epilogue overlapped; stages, accumulators, latency and epilogue as controls.

The hero animations

Recorded frame by frame from the model's states (pnpm animations); each is also a WebM video in docs/media/.

Screenshots

Regenerate them with pnpm build && pnpm start in one shell and pnpm screenshots in another.

The execution model

reference/gpu_model.py is an illustrative, parameterised model of a GPU, not a cycle-accurate simulator. A GPU is a set of figures (SM count, clock, FP32 and tensor peaks, warps, registers, shared memory with 32 banks, L2 and HBM bandwidths and sizes, tensor-core tile shapes, measured latencies). For a kernel configuration the model computes:

• the bytes moved at each memory level (registers, shared memory, L2, HBM) for element-wise kernels and tiled GEMMs, and the kernel time as a hierarchical roofline (the slowest level sets the pace);
• arithmetic intensity and the roofline position, for the FP32 units and the BF16 tensor cores;
• occupancy, by the rules of the CUDA Toolkit's cuda_occupancy.h (register and shared-memory allocation granules, the reserved kilobyte, the block limit);
• bank conflicts for an access pattern, with broadcast;
• coalescing: the 32-byte sectors a warp's request touches;
• SIMT divergence: the active mask of every instruction a warp issues;
• a load / compute / store timeline per tile, single- or multi-buffered, and its event-by-event states;
• GEMM, four ways (naive, shared-memory tiles, register blocking, BF16 tensor cores): bytes at each level, intensity and time on a 4096³ GEMM, and the tile-by-tile march of a 16³ one;
• reductions (Harris's divergent, strided and sequential versions, and warp shuffles): every value, shared-memory access, barrier, divergent warp and bank conflict;
• online softmax block by block, against the ordinary softmax, and FlashAttention-2's HBM traffic tile by tile (with and without L2 reuse of K and V) against standard attention's;
• split-K: blocks, waves, SM use and the cost of the partial sums;
• quantised layers (BF16, W8A16, W4A16, W8A8) against batch size, and dequantising eight INT4 weights in a register.

The functions that drive the animations return a list of states; a frame on the site is a pure function of one state.

Presets and sources

Two CUDA-core presets ship, A100 SXM4 40 GB and H100 SXM5 80 GB, plus three tensor-core presets for chapter 12 and the landing page (reference/blackwell.py): the same A100 and H100, and the B200 as in a GB200 NVL72 (BF16 2.5, FP8 5 and NVFP4 10 PFLOP/s dense, 8 TB/s, 186 GB, from NVIDIA's GB200 NVL72 page and the Blackwell Ultra blog's Table 2: the same figures as the AI Racks and AI Accelerators sites; 148 SMs from Jarmusch and Chandrasekaran, arXiv:2512.02189). NVIDIA does not publish the B200's clock, L2 or latencies, so the CUDA-core chapters stay on the A100 and H100. Every figure has a status (spec, rule, measured, derived or approximation) and a source with the table or page; the site lists them all at /gpus. The sources: NVIDIA's A100 and H100 architecture whitepapers, the H100 product page, the CUDA Programming Guide, the CUDA Toolkit's cuda_occupancy.h, the PTX ISA (tensor-core shapes), and Luo et al., Benchmarking and Dissecting the Nvidia Hopper GPU Architecture (measured latencies and the H100's L2 bandwidth, which NVIDIA does not publish). Derived figures and their formulas:

• shared memory: SMs × 32 banks × 4 bytes × clock;
• L2: bytes per clock × clock;
• registers: 6 bytes per flop × the FP32 peak (three 4-byte operands per two-flop FFMA);
• the H100's clock: its FP32 peak ÷ (2 × FP32 cores).

Illustrative, and labelled so on the site: the kernel-time model (every level at peak, perfectly overlapped; a GEMM's matrices either fit in L2 or every load misses it; latency ignored), the H100's L2 bandwidth (measured on an H800, the same die), the classic SIMT path order (Volta and later may order the paths differently, at the same cost), the occupancy carveout (set to its maximum), FlashAttention's traffic as two bounds (no L2 reuse of K and V, or every re-read an L2 hit), split-K's one block per SM at 1/SMs of the tensor peak, the quantised layers' roofline times (no cost for the dequantising instructions), and the overlap timeline's unit-free tile times.

Checked

• tests/python/test_gpu_model.py checks the reference against its sources' worked examples: the CUDA Programming Guide's 4-sector and 12.5% coalescing cases, its two-way and 32-way bank conflicts and the padding fix, its 75% and 50% occupancy examples, and the GEMM traffic closed forms.
• src/lib/gpu/model.ts repeats every function with the same operations in the same order. scripts/make_fixtures.py writes the reference's results over grids of every parameter (3,744 occupancy configurations, 180 coalescing patterns, 44 bank patterns, 72 branch traces, 9 overlap schedules, 8 FlashAttention loops and 20 traffic totals, 12 online softmaxes, 9 dequantisations, every animation's states), and tests/unit/model.test.ts requires the port to reproduce every value exactly, no tolerance, except the online softmax, which uses exp: V8's Math.exp and the C library's differ in the last bit for 3 of the 12 cases, so it is compared to a relative 1e-14. CI fails if the fixtures are out of date.
• Frame tests: tests/unit/frames.test.ts and tests/e2e/frames.spec.ts set key frames of every animation and require the state, and the caption on the page, to match the ones built from the Python reference's state.
• Numbers in the chapters are printed from the model at build time (<V of="a100.bw.smem" />), and every code block shown is cut from the model's source (content.test.ts); the CUDA snippets are labelled as not compiled here (CI has no GPU). Numbers the prose states in words are recomputed from the model in values.test.ts.

The animations

Every animation follows the family's visual standard: play and pause, step back and forward, a scrub bar, speed from 0.25× to 4× and reset; Space and the arrow keys when it has focus; a live caption per step, also in an aria-live region; no auto-play with prefers-reduced-motion; paused when scrolled out of view (IntersectionObserver); the equation beside the picture, with the term the animation is on highlighted (and hovering a term highlights the picture). Colours are Okabe and Ito's colour-blind-safe palette, one colour per memory level (registers orange, shared memory green, L2 sky blue, HBM purple), the same in light and dark mode; a stall is a warning hue and a hatch.

The animations are plain SVG and HTML drawn by React from the model's states, with a requestAnimationFrame clock (src/lib/anim/clock.ts, pure and unit-tested). No animation library and no D3: the pictures have at most a few thousand elements, the scales are linear or logarithmic, and the companion sites draw their charts the same way.

On phones the pictures keep their story with fewer, larger labels: every SVG label renders at 11 px or more at a 390 px width (useSvgFont measures the drawn scale; an e2e test checks every label of every animation), and labels that would then collide are dropped or shortened.

Stack

The same stack as the companion sites, minus the backend:

• Framework: Next.js 14 (App Router) + TypeScript (strict)
• Styling: Tailwind CSS, Tailwind plugin for ESLint + Prettier
• Content: MDX via next-mdx-remote, KaTeX rendered on the server
• Model: Python reference, TypeScript port, JSON fixtures
• Testing: pytest, Vitest (exact parity, frames, content; 100% line coverage on src/lib/gpu/), Playwright (every page at 1280 and 390 px, light and dark; every animation's controls; frame tests; reduced motion; axe-core scans)
• CI / deploy: GitHub Actions (model and fixtures, lint, typecheck, unit, e2e, Lighthouse), Vercel

No database and no sign-in: every page is statically rendered, and each animation is a code-split client component.

Design system: where each piece came from

Copied from llm-architectures-explained at commit 7b02ec6, which copied it from the inference site and the explainer:

Here	From
tailwind.config.ts, src/app/globals.css, src/app/layout.tsx	identical apart from titles; globals.css adds the equation-highlight rules
src/components/ui/SiteHeader.tsx	the same header; new navigation links
src/components/ui/SiteSwitch.tsx	redesigned for six sites: a full row from sm up, a <details> dropdown on phones
src/app/learn/, src/lib/mdx/, Layer.tsx, LayerToggle.tsx, MdxTable.tsx	copied; MdxTable.tsx adds a focusable pre, the chapter page makes display maths focusable
src/components/ui/Controls.tsx, Callout.tsx	unchanged (the animation panel takes the place of WidgetFrame.tsx)
.eslintrc.json, .prettierrc.json, tsconfig.json, vitest.config.ts, playwright.config.ts, lighthouserc.json, .github/workflows/ci.yml	adapted (model job, new pages)
scripts/smoke-check.ts, scripts/capture-screenshots.ts, RUNBOOK.md	adapted

New here: the animation clock, hook and panel (src/lib/anim/, src/components/anim/), the family palette (src/lib/viz/palette.ts) and the server-rendered equation component (src/components/mdx/Eq.tsx). A shared npm package for the design system would be cleaner in principle; for a handful of small sites, copying and recording the origin stays simpler.

Local development

• Node ≥ 20.11 and pnpm ≥ 9 (pinned via packageManager); Python ≥ 3.10.
• No environment variables, no database.

git clone https://github.com/BrendanJamesLynskey/gpu-kernels-explained
cd gpu-kernels-explained
pnpm install
python3 -m venv .venv && .venv/bin/pip install -r reference/requirements.txt
pnpm dev                              # http://localhost:3000

Changing the model

# edit reference/gpu_model.py, then the same change in src/lib/gpu/model.ts
.venv/bin/python -m pytest tests/python      # the reference's closed forms
.venv/bin/python scripts/make_fixtures.py    # regenerate src/data and the fixtures
pnpm test                                    # the port must match exactly

Testing

.venv/bin/python -m pytest tests/python   # reference model
pnpm lint && pnpm typecheck && pnpm format:check
pnpm test:coverage                        # Vitest with thresholds (exact parity included)
pnpm test:e2e                             # Playwright on a production build (builds first)
pnpm lighthouse                           # Lighthouse CI on a `pnpm build`
pnpm smoke <url>                          # post-deploy check of every page

Deploying

See RUNBOOK.md: a CLI deploy from a clean git archive export, then pnpm smoke.

Project layout

content/chapters/     The MDX chapters, each opening with its animation
reference/            The Python GPU execution model
scripts/              make_fixtures, smoke-check, capture-screenshots, capture-animations
src/app/              Routes: /, /learn, /learn/[slug], /gpus, /about
src/lib/gpu/          The TypeScript model, the captions, the values the prose quotes
src/lib/anim/         The animation clock
src/data/             Presets and sources (generated by make_fixtures.py)
src/components/anim/  The animation hook and panel (controls, caption, equation)
src/components/interactive/  One widget per chapter
tests/python/         pytest
tests/unit/           Vitest (fixtures in tests/fixtures/)
tests/e2e/            Playwright + axe-core

References

• NVIDIA, A100 Tensor Core GPU Architecture and H100 Tensor Core GPU Architecture (whitepapers).
• NVIDIA, CUDA Programming Guide and Nsight Compute Profiling Guide.
• Williams, Waterman and Patterson, 2009 — Roofline: an insightful visual performance model for multicore architectures, doi:10.1145/1498765.1498785.
• Luo et al., 2024 — Benchmarking and Dissecting the Nvidia Hopper GPU Architecture.
• Okabe and Ito, 2008 — Color Universal Design (the palette).
• Mark Harris — Optimizing Parallel Reduction in CUDA (NVIDIA).
• Milakov and Gimelshein, 2018 — Online normalizer calculation for softmax.
• Dao et al., 2022 — FlashAttention; Dao, 2023 — FlashAttention-2; Shah et al., 2024 — FlashAttention-3.
• Osama et al., 2023 — Stream-K.
• Frantar et al., 2022 — GPTQ; Lin et al., 2023 — AWQ; Dettmers et al., 2022 — LLM.int8().

Contributing

PRs welcome. CI runs the model job (fixtures up to date, pytest), format:check, lint, typecheck, unit tests with coverage thresholds, e2e on a production build, and Lighthouse CI (performance, accessibility and best practices must each score at least 90 on /, /gpus and five chapters).

Sharing

Every page carries Open Graph and Twitter tags, titled and addressed as that page. The link-preview cards (1,200 × 630 PNG) are drawn at build time by src/app/_og/card.tsx, the same file on every companion site, from this site's src/lib/share.ts: one for the site and one per chapter. The favicon (src/app/icon.svg) is the Hardware track's chip with this site's letter, K; /sitemap.xml lists every page and /robots.txt allows everything. tests/e2e/sharing.spec.ts checks the tags, the cards, the icons, the sitemap and the landing page's start button, strip and hero. The cards are set in Noto Sans and Noto Sans Mono (subsets, SIL Open Font License 1.1: src/app/_og/OFL.txt).

Licence

MIT — see LICENSE.