《The Stack》 EP.42 2026 · APR · 19 38 MIN IMAGINARY

Why did your GPT
end up using my GPU this way

Replay in reverse the 7-week path of an engineer who walked from vector_add to Flash Attention — and you start to see why the modern LLM evolved into exactly this shape, and how hardware and model co-evolved.

Host

Sam

imaginary · asking from the model side

Guest

Jensen

imaginary · answering from the chip side

CHAPTERS · 6 ACTS38:12

02:1401Distance is the costliest thingHBM vs FMA — 100× 08:0302Parallelism isn't freeatomic 133 ms vs tree 1 ms 13:4103Matmul feeds the LLM0.4 → 7.9 TFLOPS 20:0204Fusion, the dark artHBM trips 4 → 2 25:3805The moment everything convergesFlash Attention · 65× less 33:0706Why PyTorchtorch.ops · single entry point

00 · 02:14Cold Open

Let's go in reverse — then everything explains itself.

SAM

Slightly unusual concept today. I watched an engineer learn CUDA from the sidelines. Seven weeks, from vector_add to Flash Attention. And one question kept circling. Why did GPGPU evolve into exactly this shape?

JENSEN

(laughs) Then let's walk it backward. The kernel people call most these days is Flash Attention, right? Go backward from there and almost everything explains itself.

01 · 08:03Distance is the costliest thing

Compute is nearly free, and moving data is 100× more expensive.

Inside a GPU, adding one number is 1 ns. Fetching it from HBM is 100 ns. This single ratio governs every decision in GPU programming.

FIG 1 · access cost by tierlower = faster

FMA (register)~1 ns

L1 / smem~5 ns

L2 cache~20 ns

HBM~100 ns

PCIe H↔D~10 µs

InfiniBand~2 µs /byte²

What lessons 1 & 2 confirmed in practice

The vector_add kernel: 3.5 ms. The PCIe copy of the same data alone: 60 ms. The kernel is 17× shorter. Without pinned memory, end-to-end runs 5.6× slower.

SAM

So… how long does adding one number inside the GPU take?

JENSEN

About 1 nanosecond. Fetching it from HBM: 400–500 cycles, nearly 100 nanoseconds. Compute is nearly free, and data movement is 100× more expensive.

"The moment you leave the silicon, 100×.
That one fact set half the shape of LLM architecture."— Jensen, 09:41

SAM

So the game isn't compute faster, it's avoid moving data.

JENSEN

Exactly. That's why we built NVLink, pushed HBM density, and grew L2 cache to 48 MB. We keep shrinking the "distance." Physics only lets us go so far.

SAM

So that's why people scream to keep GPUs in the same node when they build clusters.

↳ Lesson 01 · vector_add Lesson 02 · pinned vs pageable

02 · 13:41Parallelism isn't free

Feed a serial algorithm to parallel hardware and the hardware turns serial.

If you have tens of thousands of threads, doing atomic adds together should be fast, right? No. 100× slower. That's why GPUs ship weird instructions like __shfl_down_sync.

SAM

The GPU has tens of thousands of threads. Wouldn't adding everything in parallel be fast?

JENSEN

That's the trap in Lesson 3. Version 1: every thread does atomicAdd on the same address. Guess how many ms?

SAM

… fast, I'd think?

133 ms. Other versions: 1 ms.
100× difference.— Jensen, 14:58

SAM

Wait — so the GPU hardware's weird features — shared memory, warp shuffle — are all "the hardware went in because software needed these patterns"?

JENSEN

That's co-design. __shfl_down_sync didn't exist in 1.0. People used tree reductions so much we added it six years later. Now it's standard.

FIG 2A · atomicAdd — a million in a single lineSERIAL

FIG 2B · tree reduction — log N levelsPARALLEL

__syncthreads() is not free either

Remove five of them and you get 29% back. In small kernels, a single sync is a third of runtime. That's why "last warp only shuffles, no sync" became an idiom.

↳ Lesson 03 · Four versions of Reduction

03 · 20:02Matmul feeds the LLM

Four jumps — and the last one is a new piece of hardware.

Why Lesson 4 was the longest. Tiling and register blocking are software techniques that raise arithmetic intensity; the final breakthrough past the 8-TFLOPS ceiling came from a new unit called the Tensor Core.

FIG 3 · matmul v1 → v4 — T4, 4096³TFLOPS

SAM

The axis just jumped at v4.

JENSEN

The first three are "read less HBM" — lifting arithmetic intensity (AI), flops per byte. No matter what you do, FP32 peak 8 TFLOPS is the ceiling. So Volta baked a 4×4 matmul block into the hardware.

SAM

The FP16 / BF16 / FP8 stuff I kept hearing while training GPT-4 in 2023 — all of it was to feed Tensor Cores.

100% — Tensor Cores only work on reduced-precision inputs. The descent to FP4 is about the chip as much as it is about the model.— Jensen, 22:48

occupancy trap

"Bigger tiles raise AI, so they're faster" is only half true. You also need enough blocks to feed all the SMs. That's why decode kernels and prefill kernels look different.

↳ Lesson 04 · Four versions of Matmul

04 · 25:38Fusion, the dark art

Fuse three kernels into one and HBM trips are halved — that's 2× speedup.

FIG 4A · BEFORE · 3 kernels, 4 HBM tripsunfused

k1max(x) → mR → W

k2exp(x − m) → num, sumR → W

k3num / sum → yR → W

total4 trips · 1.0×

FIG 4B · AFTER · 1 kernel, 2 HBM tripsfused online

k1pass 1: running (m, s) in registersR

k1pass 2: y = exp(x−m)/sR → W

total2 trips · ~2×

SAM

Softmax is mathematically three steps. Fine to just call three kernels, right?

JENSEN

Then HBM trips 4. Merge the three into one kernel: HBM trips 2. 2× faster. Measured: 2.02×, 1.86×, 1.92×. Almost exactly theoretical. That's operator fusion.

Why not fuse everything into one giant kernel?
— shared memory is finite, code combinatorics explode, precision issues.
"Not always good" — "picked only for hot paths."— Jensen, 27:10

SAM

And online softmax — Lesson 5's v3 — felt like it mattered more than it looked.

JENSEN

That's the mathematical half of Flash Attention. The formula that exactly merges (max, sum) when you can't fit the whole row. It was in an NVIDIA employee's 1985 paper; Tri Dao applied it to attention in 2022.

ONLINE MERGE · formulaexact, O(N) memory

# merging two partial softmax statistics
new_max = max(m1, m2)
new_sum = s1 * exp(m1 − new_max)
        + s2 * exp(m2 − new_max)

↳ Lesson 05 · Softmax & Online

05 · 28:16The moment everything converges

Flash Attention uses no fundamentally new technique — the combination was just sharp.

It never physically creates the N×N intermediate. All five earlier techniques — coalesced loads, HBM intuition, warp reduce, tiled matmul, online softmax — fit inside one kernel.

FIG 5 · HBM traffic at N=409665× less

65× less · time 4.79×

SAM

65× less should mean 65× faster, shouldn't it?

JENSEN

(laughs) In practice it's 4.79×. Why? naive wasn't really HBM-bound. L2 cached most of S. FA's real win shows up beyond the size L2 can hold — N=8k, 16k, 32k. At those sizes naive won't even run.

The reason GPT-4o's 128k context works is FA-2 + FA-3. Without them, it's impossible.— Jensen, 31:44

FIG 5B · five lessons inside one FA kernelsynthesis

L1Q / K / V coalesced load→ reg

L2HBM ↔ L2 intuition · tile sizesBr, Bc

L3warp-reduce for row max / sum__shfl_xor

L4tiled matmul in registersQ@Kᵀ, P@V

L5online (m, s) mergeAct 4 formula

↳ Lesson 06 · Flash Attention (capstone)

06 · 33:07Why PyTorch

The single entry point of the modern LLM-serving stack — `torch.ops.*`.

SAM

Until 2018 a CUDA kernel was mostly a standalone binary. Nobody does that now.

JENSEN

Because of the ecosystem. GPT-5 training runs under PyTorch — checkpointing, data loading, FSDP, autograd, optimizers, model zoo, everything. Rebuilding that is insane. So to ship a new kernel you have to plug into PyTorch.

vLLM, FlashAttention-3, Mamba — all CUDA kernels registered under torch.ops.*.
This became the single entry point of production LLM serving.— Jensen, 35:20

FIG 6 · a kernel plugging into the ecosystemtorch.compile fullgraph

Measurement · Lesson 9

Placed the whole AttentionBlock inside torch.compile(..., fullgraph=True). Graph breaks: 0. eager vs compiled err = 0.00e+00 — bit-for-bit identical.

↳ Lesson 07 · PyTorch custom op Lesson 09 · MHA causal

So why did it end up looking like this

Running the engineer's 7-week path backward, you can see that Flash Attention was inevitable.

FIG 7 · 6-step inevitability chainforward → reverse

SAM

One last thing. Five years from now, what do you think you'll be optimizing besides matmul?

JENSEN

(a brief silence) … let's save that for the next episode.

All numbers come from cudatraining lesson 1–9 handoff docs, measured on T4 sm_75 / L4 sm_89.

← Back Index · 11-post log Ep.43 → What Triton hides, and what it exposes

Why did your GPTend up using my GPU this way

Let's go in reverse — then everything explains itself.

Compute is nearly free, and moving data is 100× more expensive.

Feed a serial algorithm to parallel hardware and the hardware turns serial.

Four jumps — and the last one is a new piece of hardware.

Fuse three kernels into one and HBM trips are halved — that's 2× speedup.

Flash Attention uses no fundamentally new technique — the combination was just sharp.

The single entry point of the modern LLM-serving stack — torch.ops.*.

So why did it end up looking like this

Why did your GPT
end up using my GPU this way

The single entry point of the modern LLM-serving stack — `torch.ops.*`.