LLM serving basics (Sep 28) closed by moving the binding constraint off the allocator and onto the SLO — B = 164 rather than the 199 the memory permitted — and left two debts for today. Paging removed no bytes, only reservations, so decode still streams 13.5 GB per token; and it handed the attention kernel a block table to gather through, which is a hardware question wearing a data structure's clothes. Lecture 2 treated the machine those debts are owed to as two numbers — 989 TFLOP/s and 3,350 GB/s — and got the semester's most-used fact from their ratio, a ridge point of 295 FLOP/byte. Today we open the box, and then we write against it. First the machine: a GPU gave up everything a CPU uses to make one thread finish sooner — large caches, branch prediction, out-of-order execution — and spent the die on arithmetic units and enough resident parallelism to hide a roughly 500-cycle memory latency behind other threads' work. Then the one idea kernel writing consists of: a kernel's FLOPs are fixed by the mathematics it implements, but its bytes are a design decision, so count them, then restructure the computation until each byte fetched is used many times before it is discarded. Tiling does that inside an operator, fusion does it across operator boundaries, and Triton lets you express both in an afternoon instead of a month of CUDA. By the end you should be able to read a kernel's resource usage and predict its occupancy, say whether an access pattern wastes bandwidth and by what factor, derive a tiled matmul's HBM traffic from its tile size and place it on the roofline before running it, and say when fusion buys nothing.

Where this sits

Lecture 2 derived the roofline, the ridge point, and the decode floor from two datasheet numbers; Monday (LLM serving basics, Sep 28) built the serving loop that lives inside those constraints. Both treated the GPU as a black box with a bandwidth and a peak. Today opens it — the execution model, the memory hierarchy, the arithmetic of hiding latency — and then changes the numbers instead of computing them: same FLOPs, same GPU, and the bytes moved drop by two orders of magnitude because you chose a better kernel.

Monday (GPU kernels, Oct 5) is the sequel and owns the payoff: it takes today's two moves, tiling and fusion, composes them into FlashAttention, and reads the roofline paper properly. Today supplies the machine and the method; Oct 5 supplies the canonical application and the attention arithmetic. Everything after that — batching (Oct 7), disaggregation (Oct 14), prefix reuse (Nov 4) — dispatches the kernels this lecture explains.

Instructor notes — Timing plan

75-minute class (Mon/Wed 11:15am–12:30pm, SEC LL2.221). This is the merged GPU meeting: two former lectures' centerpieces in one room, so the budget is unforgiving and the reading-only sections are how it fits.

Time	Segment	Notes
0–4	Recap and framing	Put 989 / 3,350 / 295 back on the board; announce that today those numbers get physical addresses, then get moved.
4–8	§7.1 The bet	Die-area argument; end on 33.8 MB of registers versus a CPU's caches.
8–12	§7.2 SIMT	Thread/warp/block/grid; the divergence worked case only.
12–23	§7.3 Memory hierarchy	Centerpiece 1. Protected — never cut. Build the table row by row; end on the 118×.
23–31	§7.4 Little's law	Centerpiece 2. Protected — never cut. Derive 55-versus-64 live, digit by digit.
31–36	§7.5 Occupancy	Register table live; "cover your latency" is the line to land.
36–42	§7.6 Coalescing	The 1/32 worked case, then the 128 ms thought experiment.
42–47	§7.7 Tensor cores	989-versus-67; the n/3 climb; hand off to §7.8.
47–61	§7.8 Bytes are the budget	Centerpiece 3. Protected — never cut. Naive 0.5 → 82 ms, the traffic formula, the table, the register-file cap.
61–66	§7.9 Triton	The ownership table only; the kernel listing is reading.
66–75	§7.10 Fusion	The five-traversal table struck down to two, then the three-line decision procedure.

Reading-only, not scheduled: §7.9's code listing, §7.11, §7.12, §7.13. Each composes earlier arithmetic without adding a mechanism, so it costs nothing to leave on the page.

If running long: compress §7.6 to the 1/32 case, drop §7.5's shared-memory example, and cut §7.9 to the two-question contrast. Never cut §7.3's 118×, §7.4's 55-versus-64, §7.8's traffic formula, or §7.10's 5-passes-to-2 — those four must be on the board when the bell goes.

Learning objectives

By the end of this class you should be able to:

1. Explain, in die-area terms, why a GPU has no branch predictor, no out-of-order window, and a 33.8 MB register file — and what the design demands from your code in exchange.
2. Map a kernel's threads, warps, blocks, and grid onto lanes, schedulers, SMs, and the chip, and predict the throughput cost of a divergent branch.
3. Reproduce the H100 memory hierarchy — capacities, rough latencies, sharing scopes — and compute why no on-chip level can hold the reference 7B's weights.
4. Derive, via Little's law, the number of in-flight loads a stated bandwidth and latency demand, and compare it against the hardware's warp budget.
5. Compute a kernel's occupancy from its per-thread registers and per-block shared memory, and the coalescing efficiency of a stated access pattern.
6. Derive a matmul kernel's HBM traffic as a function of its tile sizes, compute its arithmetic intensity, place it on the roofline before running it, and name which on-chip budget caps the tile.
7. Count the bytes an unfused operator chain moves, decide which fusions pay, and quantify the saving — and say why no tile size rescues batch-1 decode.

7.1 The bet: throughput over latency

A CPU and a GPU face the same physical fact — DRAM answers in hundreds of cycles — and spend their transistors on opposite responses to it. A CPU spends them making a single instruction stream finish sooner: megabytes of private cache so most accesses never reach DRAM, a branch predictor so fetch never waits for a decision, an out-of-order window and speculation so independent work fills the stall. All of that machinery exists to make latency smaller, and it is most of the die.

A GPU deletes nearly all of it: no branch prediction, no speculation, no out-of-order execution, and caches that are small per thread and shared by thousands. The recovered area buys arithmetic units and register files large enough that thousands of threads stay resident — register state allocated, ready to issue. Latency is then not reduced but tolerated: when a warp issues a load that will take 500 cycles, the scheduler issues from a different warp on the next cycle, and the switch costs zero because every resident warp's registers are allocated simultaneously. Nothing is ever saved or restored; a "context switch" is a multiplexer picking a different warp id.

Question	CPU answer	GPU answer
A load misses to DRAM — now what?	speculate past it in one thread	issue from a different warp
Where do the transistors go?	caches, predictors, reorder logic	ALUs and register files
What is kept close to the ALUs?	recently used data	ready-to-run thread state
Cost of a wrong guess / a stall	pipeline flush	zero, if another warp is eligible
What the design demands of software	locality	parallelism

Neither column is smarter; they optimize different objectives. A CPU minimizes the time to finish this instruction stream; a GPU maximizes the number of streams finished per second and lets any individual one wait. Serving systems care about the second objective almost everywhere — which is why the last row is the one to remember: the machine only works if your kernel shows up with tens of thousands of independent things to do.

Sit with the register number. A server CPU keeps a few hundred hardware threads; an H100 keeps a quarter of a million resident, and its register file — 33.8 MB — is within sight of its own 50 MB L2 and larger than the last-level cache of most CPUs of its era. The largest memory-like structure near the arithmetic units is not a cache. It is thread state. That inversion is the design, and §7.4 and §7.5 will price exactly how much parallelism it demands of you in exchange.

Instructor notes

Minutes: 4. Board: Two columns, "CPU spends area on" / "H100 spends area on". Fill the left with cache/predictor/OoO, the right with ALUs and "33.8 MB of registers". Circle the register number. Ask the room: "What does a GPU do on a branch mispredict?" Nothing — it never predicted. Let the silence land; it reframes the whole machine. Expect confusion: Students assume the GPU is a CPU with more cores. Say: "It is not a thousand small CPUs. It is one machine that refuses to wait, because someone else is always ready."

7.2 SIMT: threads, warps, blocks, grids

The programming model has four levels, and each maps to a physical home:

Software unit	Size	Hardware home	What it shares
thread	1	one lane	its own registers (private)
warp	32 threads	one scheduler slot	one program counter, issued together
block	up to 1,024 threads	one SM	shared memory, barrier sync
grid	all blocks	the chip	L2 and HBM only

The warp is the real unit of execution: the hardware never issues fewer than 32 lanes, whatever your code says. Each SM has four warp schedulers, and each cycle every scheduler picks one eligible warp — operands ready, not stalled — and issues its next instruction. A stalled warp costs nothing while any other warp is eligible; this free switch is the mechanism §7.4 quantifies. This model is called SIMT — single instruction, multiple threads — and it is a contract: you write scalar-looking code per thread, the hardware executes it 32 lanes at a time, and the price appears only when the 32 disagree.

A CUDA kernel reads threadIdx and blockIdx to compute which slice it owns — i = blockIdx.x * blockDim.x + threadIdx.x is the idiom — while Triton (§7.9) drops the thread level entirely and hands you the block, precisely because warps and lanes are better left to the compiler. Either way the grid names hundreds of thousands of independent work items, which §7.4 will show is not generosity but the minimum the memory system demands.

Divergence is the price. When lanes of one warp take different sides of a branch, the warp runs both paths serially under an active mask — lanes on the not-taken side execute the instructions and discard the results; in the worst case a 32-way divergent branch runs at 1/32 throughput. (Since Volta each thread carries its own program counter, which changes what deadlocks are possible but not the cost: the lanes still share issue slots.)

The LLM-relevant instance is per-token control flow: a batch of ragged sequence lengths handled with branches diverges in nearly every warp, which is why serving kernels prefer masking and padding — do the work for every lane, multiply the unwanted results by zero — over branching around it. Wasting lanes on discarded arithmetic looks inefficient and is usually the fast path; the arithmetic was going to be paid for either way. (§7.9's mask= arguments are this idea in source form.)

Blocks are the other half of the contract: the independence guarantee. The hardware promises nothing about the order blocks run in or which SM gets which; in exchange, the same kernel scales from 1 SM to 132 without recompiling. The cost of that freedom is that blocks cannot communicate within a kernel (with narrow exceptions) — synchronization at grid scope means ending the kernel and launching another, at Lecture 2's ≈5 µs a launch, a figure §7.10 will spend.

Instructor notes

Minutes: 4. Board: The four-row table, then the divergence example with 2⁻³¹ ≈ 5e-10 next to it. Ask the room: "One thread of a warp loads from memory; the other 31 don't need data. What does the hardware do?" All 32 lanes participate; you cannot buy less than a warp. Expect confusion: "Threads" imported from POSIX — students expect independent progress and locks. Say: "A GPU thread is a lane of a vector that is pretending, politely, to be a thread."

7.3 The memory hierarchy, with numbers

Here is the machine's memory system, from the arithmetic units outward. Build it row by row; the whole rest of Part II is a walk down this table.

Level	Capacity	Bandwidth	Latency	Managed by
Registers	256 KB/SM → 33.8 MB chip	operands every cycle	~0 (allocated, not fetched)	compiler
Shared memory (SMEM)	228 KB/SM → 30.1 MB chip	~an order of magnitude above HBM, aggregate	~30 cycles	you (§7.8)
L2	50 MB, chip-shared	a few × HBM	a few hundred cycles	hardware
HBM3	80 GB	3,350 GB/s	~500+ cycles	hardware / allocator

Three notes on the entries. Shared memory is the same physical SRAM array as the L1 cache, carved up by configuration; it is the one level you address explicitly, and programming it is most of §7.8. Treat its aggregate bandwidth as "~10× HBM": at roughly 128 bytes per cycle per SM — an approximate figure — 128 × 132 × 1.8 GHz ≈ 30 TB/s against HBM's 3.35. The L2 row is deliberately vague because NVIDIA publishes neither its latency nor its bandwidth; such figures come from the pointer-chasing kernels and stride sweeps Dissecting Volta and Dissecting Hopper build, and "a few hundred cycles, a few times HBM bandwidth" is the defensible summary. Finally, the sharing scope is the hierarchy — registers per-thread, SMEM per-block, L2 and HBM chip-global — because wider sharing costs more cycles, which is what distance on a die means.

Now derive the two cliffs the table implies.

No caching strategy, however clever, makes decode's weight reads cheap: every generated token must stream essentially all 13.5 GB from HBM, because there is nowhere on the chip to keep them. This is the physical location of Lecture 2's 4.0 ms decode floor — 13.5 GB ÷ 3,350 GB/s — and the reason the floor is a statement about HBM and nothing else.

The latency cliff: HBM at ~500 cycles versus SMEM at ~30 is a ~17× gap, and at 1.8 GHz, 500 cycles is ≈ 278 ns. Every trip your kernel's data takes down one level of the table and back is a trip a better kernel does not take — that sentence, quantified, is §7.8.

Now divide the headline bandwidth by the machine that wants to consume it, because "3,350 GB/s" is a chip-wide figure and no kernel experiences the chip; it experiences an SM.

One SM's fair share of the monster is fourteen bytes a cycle — a 36× shortfall against what its own issue width can ask for. The quoted bandwidth sounds enormous until you divide it by the arithmetic sitting next to it, and the entire discipline of kernel writing is arranging for most operands to come from the upper rows of the table, where the supply actually matches the demand.

The hierarchy is not an implementation detail to abstract over; it is the resource you program. §7.8 is, in its entirety, about moving the same bytes fewer times down this table.

Instructor notes

Minutes: 11. Centerpiece 1. Board: The table, row by row, capacities first. Then 33.8 + 30.1 + 50 ≈ 114 MB, then "13.5 GB / 114 MB ≈ 118×" — box it. Then 3,350 ÷ 132 ≈ 25.4 GB/s → 14.1 B/cycle → the 36× shortfall. Ask the room: Before revealing the total: "Can the 7B's weights fit anywhere on this chip?" Let them add the column themselves. Expect confusion: Students expect L2 to behave like a CPU's LLC and "cache the hot weights". Say: "During decode every weight is touched exactly once per token. There is no hot set. A cache needs reuse, and batch-1 decode has none." Common wrong answer: "Then add more SRAM." 114 MB is already a large fraction of the die; ×118 more is not a manufacturing option — it is a different machine, and §7.12's TPU shows what choosing more SRAM and less DRAM bandwidth looks like.

7.4 Latency hiding, by Little's law

The question this section answers: what does it take for HBM to actually deliver 3,350 GB/s when each access takes ~500 cycles? The answer is a one-line theorem from queueing:

If the pipe is to stay full, the amount of data in transit at any instant must equal the rate times the trip time. Apply it to the H100:

Now put that next to the hardware's residency limit. If each warp keeps roughly one load in flight — the natural state of a simple streaming kernel — then an SM needs ~55 warps resident and stalled on memory just to keep HBM busy, and the hardware's ceiling is 64 warps per SM. The occupancy limit is not an arbitrary constant: it is sized to be barely enough to cover HBM latency. The 64-warp budget of §7.1 and the ~500-cycle latency of §7.3 are two halves of one equation.

One honest caveat, which is also a forward reference: a warp is not limited to one outstanding load. Unrolled loops and independent loads give instruction-level parallelism within a warp — several loads in flight before the first is consumed — so fewer warps can suffice if each carries more. This is the loophole that pipelined, low-occupancy-high-ILP kernels exploit deliberately (§7.9's num_stages): fewer warps, each fatter with in-flight work, leaving more registers and SMEM per warp.

The corollary cuts the other way and matters more day to day: a kernel that cannot put thousands of independent loads in flight cannot run at quoted bandwidth, whatever its arithmetic looks like. Any serialization — pointer chasing, a dependent reduction done naively, tiny grids that leave SMs empty — caps effective bandwidth far below 3,350 GB/s.

The pointer-chase corollary, reading-only. One thread walking a linked list, where each load's address is the previous load's result, keeps exactly one load in flight. At 1.8 GHz and 500 cycles that is 3.6 million dependent loads/s; at an 8-byte pointer each, ≈29 MB/s of useful traffic — about 116,000× below the machine's useful ceiling, on the same DRAM and the same pins. So "bandwidth-bound" is never a property of hardware alone; it is a property of an algorithm that earned the bandwidth by exposing concurrency. It is also why every serving-system structure we meet later — paged KV blocks, radix trees over prefixes — chases its pointers on the CPU and hands the GPU flat, pre-resolved addresses.

Batch-1 decode passes the concurrency test easily: streaming 13.5 GB of weights is embarrassingly parallel in the load dimension, every address known in advance. So the floor of 4.0 ms is achievable — just useless, at 0.34% of the arithmetic peak (Lecture 2). The GPU's problem with decode was never getting the bytes; it is having nothing to do with them.

Sensitivity, in one line: the 55 scales linearly with the assumed latency and inversely with clock — at 400 cycles it is 44, at 600 it is 66. The conclusion "you need most of the warp budget" survives any defensible choice of constants.

Instructor notes

Minutes: 8. Centerpiece 2. Protect this budget absolutely. Board: bytes in flight = BW × latency, then the four lines of the worked block, digit by digit, ending with "55 needed, 64 available" boxed next to §7.3's 118×. Ask the room: Before dividing by 132: "The chip needs 7,270 loads in flight. Who supplies them?" You do. The hardware provides slots; the kernel provides the parallelism. Expect confusion: Latency and bandwidth treated as independent virtues. Say: "Bandwidth is a rate you only get if enough requests are in the pipe. An empty pipe has infinite bandwidth and delivers nothing." Common wrong answer: "Caches fix this." Caches shorten the trip for reused data; a streaming read of 13.5 GB has no reuse, so the only fix is concurrency.

7.5 Occupancy: what fits on an SM

Occupancy is resident warps divided by the 64-warp maximum. It is set at launch time by three budgets — take the minimum:

1. Registers: 65,536 32-bit registers per SM, allocated per thread by the compiler (any single thread caps at 255).
2. Shared memory: 228 KB per SM, allocated per block by you.
3. Hard caps: 64 warps, 2,048 threads, and a resident-block limit (32 on Hopper) per SM.

The register budget alone produces a table worth memorizing the shape of:

Registers/thread	Resident threads (65,536 ÷ regs)	Warps	Occupancy
32	2,048	64	100%
64	1,024	32	50%
128	512	16	25%

So the compiler's register allocation is a throughput decision, not a code-generation detail: spill a value to memory and you buy slow load/store instructions; keep it in a register and you may halve the number of warps that fit. Neither choice is free, and -maxrregcount-style knobs exist precisely because the compiler cannot know which side of the trade your kernel is on.

Shared memory imposes the same arithmetic per block, and the two budgets bind independently:

The binding resource is whichever budget divides worst; it is the same min() logic as every capacity computation in this course, from Lecture 2's KV-cache ceiling onward. And notice the tension it sets up with §7.8: bigger tiles in SMEM mean more reuse per byte fetched (good for the roofline) but fewer resident blocks (bad for latency hiding). Tile size is not a free parameter — it trades one section of this lecture against another, which is why real kernels sweep it (§7.11).

Now connect to §7.4, where "occupancy" stops being folklore. Occupancy is not a virtue; it is the supply of stall-tolerant work. A memory-streaming kernel with one outstanding load per warp cannot saturate HBM below ~55 resident warps per SM — §7.4's demand side — and above that point extra occupancy buys nothing. "Maximize occupancy" is the folk advice; "cover your latency" is the rule, and enough instruction-level parallelism per warp covers it at 25% occupancy while a naive kernel starves at 50%.

Instructor notes

Minutes: 5. Board: The three budgets as min(regs, smem, caps), then the register table — make them compute the 64-regs row. Then write "§7.4 said 55" underneath the occupancy column. Ask the room: "Your kernel is at 40% occupancy. Is that bad?" Wrong question — ask whether 25 warps' worth of in-flight loads covers its stalls. Occupancy is supply; the kernel's stall profile is demand. Expect confusion: Occupancy read as utilization. Say: "100% occupancy with every warp stalled on the same dependent load is 0% utilization. Occupancy counts bodies, not work." If short on time: Cut the SMEM example; the register table plus the §7.4 connection is the section.

7.6 Coalescing: how a warp touches memory

A warp does not issue 32 loads; it issues one memory request with 32 addresses, and the memory system services it in 128-byte aligned segments. The hardware moves segments, never bytes. Everything about access-pattern efficiency follows from one rule of thumb: efficiency = useful bytes ÷ segment bytes touched.

Best case: 32 threads read 32 consecutive fp32 values — 128 contiguous bytes, one aligned segment, 100% of moved bytes useful. This is a coalesced access, and it is the pattern the 3,350 GB/s figure is quoted for.

One indexing choice demotes an H100 below a laptop's memory system. Cash it out on the course's central number as a thought experiment: if the reference 7B's weight reads during decode were fully strided like this, the 4.0 ms batch-1 floor would become 32 × 4.0 = 128 ms per token — about 8 tokens/s. Same GPU, same model, same FLOPs; the entire 32× is address arithmetic. Real weight layouts are contiguous, of course — which is precisely why the floor is reachable. The thought experiment tells you what the layout is worth.

The milder cases follow from the same rule, and they tabulate cleanly — this is a table to be able to regenerate, not to memorize:

Pattern (one warp)	Segments touched	Moved / useful bytes	Efficiency	Effective bandwidth
32 consecutive fp32, aligned	1	128 / 128	1	3,350 GB/s
consecutive fp32, misaligned	2	256 / 128	1/2 for that warp; ~1 amortized over a stream	≈3,350 GB/s
stride-2 fp32	2	256 / 128	1/2	~1,675 GB/s
stride-32 fp32 (column read)	32	4,096 / 128	1/32	~105 GB/s
stride-64 fp16 (column read)	32	4,096 / 64	1/64	~52 GB/s

Two things to notice. Misalignment is nearly free for streams because adjacent warps reuse each other's boundary segments — the segments all get consumed by someone — while strides are expensive at any scale because nobody ever wants the rest of the segment. And smaller elements make strided access worse, not better: an fp16 read whose address stride still spans a full segment (≥ 64 elements) uses 2 bytes of each 128-byte segment → 1/64 (exercise 3). Halving your precision halves your bytes only if the layout keeps them adjacent.

Two fixes, both used in §7.8. Either change the data layout so the fast-moving index is contiguous, or stage through shared memory: load coalesced row-wise into SMEM, then read column-wise out of it — SMEM has banks, not segments, and does not care about this pattern. (Its own hazard is the bank conflict: 32 banks, and two lanes hitting the same bank at different addresses serialize. Same shape of problem, one level up the table — and in Triton it is the compiler's problem, not yours.)

Bandwidth is quoted for the machine's favorite access pattern; every other pattern gets a fraction, and the fraction is computable in your head before you write a line of code.

Instructor notes

Minutes: 6. Board: Draw one 128-B segment as a box; 32 arrows into one box (coalesced), then 32 arrows into 32 boxes (strided). Write "4,096 moved / 128 used" and then "4.0 ms → 128 ms". Ask the room: "Whose fault is a strided access — the hardware's or the program's?" The layout's. Then: who owns the layout? You do — that is §7.8. Expect confusion: Coalescing confused with caching. Say: "The segment is the unit of transfer, not of retention. You pay for the whole segment on this access whether or not L2 keeps it." If short on time: The 1/32 case and the 128 ms line; drop stride-2 and misalignment.

7.7 Tensor cores: matmul as a hardware primitive

Everything so far describes lanes doing scalar arithmetic. Since Volta, that is no longer where the FLOPs are. A tensor core executes a small matrix-multiply-accumulate — a warp-wide MMA over tile fragments held in registers — as one instruction, rather than as hundreds of scalar fused multiply-adds spread across 32 lanes. The H100 SXM feeds them at 989 TFLOP/s dense BF16; the ordinary fp32 vector lanes manage about 67 TFLOP/s (datasheet figure). The matmul path is ~15× the everything-else path, and per SM it is 989 ÷ 132 ≈ 7.5 TFLOP/s.

Why did matmul, of all operations, get its own silicon? Because it is the one common workload whose arithmetic intensity grows with problem size: an n×n matmul does 2n³ FLOPs on 3n²·b bytes of operands, so intensity ≈ 2n/(3b) — at bf16, n/3 FLOP/byte, assuming each operand is moved once.

Read the first row carefully, because it is §7.8's motivation in one line: a matmul the size of one 128-wide tile, with operands fetched from HBM, is bandwidth-bound — tensor cores idle at 14%. Large GEMMs escape only because a tile's operands, once staged in SMEM and registers, are reused against many other tiles; the n/3 climb is real only if the reuse actually happens on chip. Tensor cores are the hardware half of that bargain; tiling is the software half, keeping the MMA units fed from the top rows of §7.3's table instead of from HBM.

The asymmetry this creates is the one Lecture 2 flagged and this lecture can explain. Elementwise and normalization ops run at ~1 FLOP/byte, so their ceiling is 1 × 3,350 GB/s = 3.35 TFLOP/s — 0.34% of peak — and no tensor core helps an operation starved of bytes. At a 32,768-token microbatch those ops are 0.19% of the FLOPs but ~15× the bytes of the layer weights (Lecture 2). GEMMs get 989 TFLOP/s; everything between the GEMMs gets bandwidth, so left as separate kernels the cheap ops dominate wall-clock — the entire case for fusion, §7.10.

Precision is a throughput dial on the same units: FP8 doubles the feed rate to ≈1,979 TFLOP/s at unchanged bandwidth, so the ridge doubles to 1,979 ÷ 3.35 ≈ 591 FLOP/byte — every precision drop makes the compute ceiling easier to buy and the intensity bar higher to clear (Lecture 2's fp8 row, now with a physical mechanism).

The modern GPU is no longer a general parallel processor with a matmul habit; it is a matmul engine with a general parallel processor attached, and your job as a kernel writer is to keep everything that is not matmul off the critical path.

Instructor notes

Minutes: 5. Board: "989 vs 67 → 15×", then intensity(n×n) ≈ n/3, then n ≈ 885 crosses 295. Then re-write Lecture 2's "0.19% of FLOPs, 15× the bytes" and draw the arrow to "§7.10, fusion". Ask the room: "You add tensor cores to a chip. What happens to every kernel that isn't a GEMM?" Relatively speaking, it gets 15× worse. Peak rose; their diagonal didn't move. Expect confusion: "Use tensor cores for everything." They compute matmuls; a LayerNorm has no n³ to offer them. The bound is bytes, and §7.3 set the bytes.

7.8 Bytes are the budget: naive matmul, and tiling

Everything so far measured the machine. From here we change the numbers instead of computing them. The contract §7.3 and §7.7 impose: a kernel's FLOPs are fixed by the mathematics it implements, but its bytes are a design decision, and the ratio of the two decides everything.

The method, stated once and then reused for the rest of Part II:

1. Count the algorithmic FLOPs. These do not change.
2. Count the HBM bytes moved, under a stated reuse assumption.
3. Intensity = FLOPs ÷ bytes.
4. Attainable throughput = min(peak, intensity × bandwidth) — Lecture 2's roofline.
5. Time = max(FLOPs ÷ peak, bytes ÷ bandwidth). Whichever term wins names the bound.

Only step 2 is yours to change. Recall Lecture 2's two anchors on the intensity axis: batch-1 decode sits at exactly 1 FLOP/byte, a 4K prefill at ≈4,100 — the workload spans four orders of magnitude. Today's subject is how a kernel moves along that axis for a fixed problem: the same matmul can sit at 0.5 FLOP/byte or at 128, depending on choices you make.

Naive matmul. C = A·B with A of shape M×K, B of shape K×N, all bf16 (b = 2 bytes). FLOPs = 2·M·N·K, fixed. One thread per output element reads its full row of A and its full column of B from HBM: under zero-reuse accounting the reads are 2·M·N·K elements = 4·M·N·K bytes, plus a negligible 2·M·N of writes. Intensity = 2MNK ÷ 4MNK = 0.5 FLOP/byte, independent of problem size — both terms scale with MNK, so the ratio is pinned. Attainable: 0.5 × 3,350 GB/s = 1.68 TFLOP/s, 0.17% of peak, which is 0.5/295.

One honest footnote before we fix it: a real "naive" CUDA kernel measures better than this, because the 50 MB L2 catches some row and column re-reads by accident. The 0.5 is the zero-reuse accounting. Tiling's contribution is not that reuse exists — it is that reuse happens by design, in a memory you control, at a rate you can derive, rather than by luck in a cache you share with 131 other SMs.

Tiling, the mechanism. Partition C into tiles of BLOCK_M × BLOCK_N and assign each tile to one program (CUDA: thread block). The program loops over K in chunks of BLOCK_K, loading one A-tile (BLOCK_M × BLOCK_K) and one B-tile (BLOCK_K × BLOCK_N) into shared memory each iteration and multiplying them into an accumulator held in registers. Inside that step every element of the A-tile participates in BLOCK_N multiply-adds and every element of the B-tile in BLOCK_M: each byte fetched from HBM is used BLOCK times before it is discarded, instead of once.

The derivation. Each program reads its full A row-strip (BLOCK_M × K) and full B column-strip (K × BLOCK_N) over its K-loop, and there are (M/BLOCK_M) · (N/BLOCK_N) programs, so A is re-read once per column of tiles and B once per row:

(The per-program route gives the same total: each of the MN/(BLOCK_M·BLOCK_N) programs loads K·(BLOCK_M + BLOCK_N) elements. Two routes, one answer — a useful check whenever you count bytes.)

Read the formula before the table. Traffic is proportional to 1/BLOCK_M + 1/BLOCK_N, so doubling the tile edge halves the bytes; and for a fixed tile area that sum is minimized when the tile is square — perimeter versus area, the same reason cells are round. Now the 4096³ example, where b·MNK = 137.4 GB and bytes = 137.4 GB × (2/BLOCK) for square tiles:

BLOCK	HBM bytes	Intensity (FLOP/byte)	Attainable	Memory time	vs 82 ms naive
32	8.59 GB	16	53.6 TFLOP/s	2.56 ms	32×
64	4.29 GB	32	107 TFLOP/s	1.28 ms	64×
128	2.15 GB	64	214 TFLOP/s	0.64 ms	128×
256	1.07 GB	128	429 TFLOP/s	0.32 ms	256×

Traffic falls as 1/BLOCK; the improvement over naive is exactly BLOCK, because naive is the degenerate tiling BLOCK = 1 at intensity 0.5 — the formula covers it too (2·1·1/(2·2) = 0.5).

The feasibility check, against §7.3's and §7.5's per-SM budgets. Take BLOCK = 128 with BLOCK_K = 64. Shared memory per program: (128·64 + 64·128) elements × 2 bytes = 32 KB; double buffering (loading step k+1 while computing step k) doubles it to 64 KB of the 228 KB — comfortable. The accumulator must be fp32 (Lecture 2's precision discipline: bf16 inputs, fp32 accumulation) and lives in registers: 128·128·4 B = 64 KB of the 256 KB register file — a quarter of the file. Now try BLOCK = 256: the SMEM tiles still fit (128 KB double-buffered), but the accumulator is 256²·4 = 256 KB — the entire register file of an SM, leaving nothing for addresses, indices, or operands. The register file, not shared memory, is what caps the tile.

The honest ceiling. Set BLOCK/2 = 295 to reach the ridge: BLOCK ≈ 590, whose accumulator would need 590²·4 B = 1.39 MB — 5.3× the register file — and whose double-buffered SMEM tiles ≈302 KB against 228 KB. Impossible on both budgets. Square shared-memory tiling therefore tops out around BLOCK = 128–256, intensity 64–128, an attainable ceiling of 128 × 3,350 = 429 TFLOP/s ≈ 43% of peak — from this level of the hierarchy alone. Yet well-tuned matmuls exceed 90% of peak on large shapes. The gap is real and this derivation cannot close it; what closes it is the same formula applied at the other levels of §7.3's table:

• Coalescing changes no term in the byte count, yet it is the single largest jump in Boehm's CUDA ladder. HBM serves segments (§7.6), so an uncoalesced kernel delivers several times its algorithmic bytes. Step 2 above silently assumed coalesced access — every derivation today inherits that assumption.
• Register tiling is the formula one level down, with SMEM as the slow memory: each thread owns a TM × TN micro-tile of the accumulator (256 threads, a 128×128 tile → 64 accumulators per thread, an 8×8 micro-tile), reading TM + TN = 16 SMEM elements per 2·TM·TN = 128 FLOPs — 8 FLOPs per SMEM element where the naive within-block scheme managed 1. SMEM has bandwidth too, and without this rung it becomes the bottleneck the moment HBM stops being one.
• L2 reuse across programs is what closes the 43% gap. Neighboring programs share A row-strips and B column-strips; launch them in an order that keeps a "supergroup" of adjacent tiles in flight together (grouped or swizzled program ordering) and their shared strips stay resident in the 50 MB L2, served at L2 bandwidth. The effective BLOCK in the formula becomes the supergroup's footprint, which 50 MB can hold where a 256 KB register file could not. This, plus rectangular tiles tuned per shape, is how real kernels travel from 43% toward 90%+.

Reading-only aside — wave quantization. The 4096² output at BLOCK = 128 makes (4096/128)² = 1,024 programs on 132 SMs: 7 full waves of 132, then a tail wave of 100 while 32 SMs idle. Tile choice changes how ragged that tail is, which is one reason §7.11's autotuner sometimes prefers a tile that is worse on the formula and better on the calendar.

Tiling is not a trick. It is the formula bytes ∝ (1/BLOCK_M + 1/BLOCK_N), and §7.3's hierarchy gives you three places to apply it.

Instructor notes

Minutes: 14. Centerpiece 3. Never cut the formula or the feasibility check. Board: Naive first: 274.9 GB before 82 ms, then 0.139 ms next to it and the 590× bracket, then "590 = 295 ÷ 0.5". Then the strip picture (one tile of C, its A row-strip and B column-strip), the total-traffic count, and box the formula. Table next. Feasibility digit by digit: 32 KB, 64 KB, then 256 KB = "the whole register file". End with "43%" and "90%+" side by side and the three bullets as the bridge. Ask the room: After the formula, before the table: "Predict the improvement over naive at BLOCK = 64." They should say 64× from the formula alone — that moment is why the derivation precedes the table. Expect confusion: Students think shared memory is the binding budget because it has "memory" in the name. The accumulator arithmetic settles it: registers bind first. Common wrong answer: "So use BLOCK = 590 across two SMs." Programs do not span SMs; that is what the L2 rung is for.

7.9 Triton: the tile is the primitive

Everything in §7.8 is expressible in CUDA, and in Boehm's post it is — one hand-written kernel per rung, each a real afternoon (or week) of engineering. Triton's bet is that the decisions worth a human are the ones §7.8 derived, and the rest should be a compiler's problem.

The model. You write what one program — one output tile — does, at block granularity. There is no threadIdx anywhere in your source: tl.program_id(0) says which tile you are, tl.arange(0, BLOCK) builds index vectors so pointer arithmetic operates on whole blocks, tl.load/tl.store take boolean mask= arguments that handle the ragged edges where a tile overhangs the matrix, tl.dot is the tile-times-tile product lowered to tensor cores, and tl.max/tl.sum/tl.exp are block-level reductions and elementwise math for kernels like §7.10's.

What the compiler owns — each row of this table is a hand-written kernel or rewrite in the CUDA ladder:

Decision	In CUDA	In Triton
Thread-to-element mapping, coalescing	you, per kernel	compiler
SMEM allocation, layout, bank-conflict avoidance	you	compiler
Double buffering / software pipelining of the K-loop	you	compiler (num_stages)
Vectorized loads/stores	you	compiler
Register allocation, micro-tiling	you	compiler (num_warps)

What you still own: the tile sizes, the grid, the masks, and the algorithm's byte count. Triton automates the ladder's labor, not §7.8's arithmetic. A Triton kernel with bad BLOCK sizes is a fast implementation of a slow policy, and the compiler will faithfully pipeline your way to the wrong roofline point.

The cost, stated fairly. You give up warp-level control: shuffles, hand-placed tensor-core fragments, exotic layouts. A well-tuned Triton matmul is competitive with vendor libraries on standard shapes — the tutorials benchmark theirs against cuBLAS at rough parity — but the last few percent on the shapes NVIDIA cares about belongs to CUTLASS specialists, and some tricks are simply not expressible. For custom fusions on LLM shapes, where the alternative is an unfused chain rather than a hand-tuned library, the trade is heavily in Triton's favour.

CUDA asks "what does thread 17 do?"; Triton asks "what does tile (i, j) do?" — and tile (i, j) is exactly the unit §7.8's derivation was written in. The language's primitive matches the analysis's primitive, which is why the analysis transcribes into code almost mechanically.

Reading-only: the annotated matmul. Below is the canonical kernel, in the shape of Triton's own tutorial 03, with each region labelled by the §7.8 quantity it implements. Nothing in it is new — the derivation already happened; this is the formula wearing syntax. Work through it alongside tutorials 01–03 before Oct 5.

@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K,
                  stride_am, stride_ak, stride_bk, stride_bn,
                  stride_cm, stride_cn,
                  BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr,
                  BLOCK_K: tl.constexpr):
    # -- which BLOCK_M x BLOCK_N tile of C is mine? (one program = one tile)
    pid = tl.program_id(axis=0)
    num_pid_n = tl.cdiv(N, BLOCK_N)
    pid_m, pid_n = pid // num_pid_n, pid % num_pid_n

    # -- index vectors for my rows, columns, and the K-chunk
    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    offs_k = tl.arange(0, BLOCK_K)

    # -- block pointers into A's row-strip and B's column-strip
    #    (these strips are the M*K*(N/BLOCK_N) and K*N*(M/BLOCK_M) terms
    #     of the traffic formula: each program walks one of each)
    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak
    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn

    # -- the accumulator tile: fp32, lives in registers (the 64 KB of Sec 7.8)
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

    for k in range(0, tl.cdiv(K, BLOCK_K)):
        # -- masked block loads: each loaded element will be reused
        #    BLOCK_N (for a) or BLOCK_M (for b) times -- the reuse itself
        a = tl.load(a_ptrs, mask=(offs_m[:, None] < M) &
                                 (offs_k[None, :] < K - k * BLOCK_K), other=0.0)
        b = tl.load(b_ptrs, mask=(offs_k[:, None] < K - k * BLOCK_K) &
                                 (offs_n[None, :] < N), other=0.0)
        acc += tl.dot(a, b)          # tensor cores; 2*BM*BN*BK FLOPs per trip
        a_ptrs += BLOCK_K * stride_ak
        b_ptrs += BLOCK_K * stride_bk

    # -- epilogue: cast once, store once (Sec 7.10 will graft more work here)
    c = acc.to(tl.bfloat16)
    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn
    tl.store(c_ptrs, c, mask=(offs_m[:, None] < M) & (offs_n[None, :] < N))

Grid: (tl.cdiv(M, BLOCK_M) * tl.cdiv(N, BLOCK_N),) — one program per output tile, the (M/BM)·(N/BN) of the derivation. Three things to notice. The masks handle boundary tiles in one comparison per load where CUDA needs cloned edge-case kernels — but masked-off lanes still occupy their slots, so a ragged edge wastes the padding's compute (one reason §7.11's autotuner cares about shape, not just hardware). The two invisible optimizations: nothing mentions shared memory, yet the A and B tiles pass through it in a bank-conflict-free layout the compiler chose; nothing mentions pipelining, yet at num_stages ≥ 2 the compiler issues iteration k+1's loads while tl.dot chews on iteration k — the double buffering §7.8 budgeted 2× SMEM for. Two rungs of the ladder, absorbed. The epilogue casts and stores exactly once, and that single tl.store is the seam §7.10 exploits.

Instructor notes

Minutes: 5. The code listing is reading-only — do not walk it in class. Board: The two-question contrast ("what does thread 17 do?" / "what does tile (i, j) do?"), then the ownership table's left column only — ask the room to sort each row into "you" or "compiler" before filling it in. Ask the room: "You wrote a Triton kernel with BLOCK_M = BLOCK_N = 16. The compiler pipelines it perfectly. How fast is it?" Intensity 8, so ≈27 TFLOP/s max — the compiler cannot fix step 2. Expect confusion: "Compiler handles memory" gets heard as "memory is solved." The ownership split is the correction: the count is yours, the choreography is the compiler's.

7.10 Fusion: never write the intermediate

Tiling fixed the GEMM. Now look at the GEMM's neighbors, with §7.7's ledger open: elementwise and normalization ops are 0.19% of a transformer's FLOPs but ~15× the bytes of the layer weights at a 32,768-token microbatch, and normalization sits at 1 FLOP/byte — hopelessly bandwidth-bound. Around every well-tiled GEMM is a parade of small kernels, each reading the full activation tensor from HBM and writing it back: norm, activation, residual-add, each a round-trip.

The move. Fuse. Either graft elementwise work onto the producing GEMM's epilogue — apply bias, activation, residual to the accumulator tile while it is still in registers, before §7.9's single tl.store — or collapse a chain of elementwise kernels into one pass that reads once, does all the math, and writes once. Every fused op deletes one full write and one full read of an intermediate from HBM.

The set piece: fused softmax. Take one S×S attention score matrix at S = 4,096, bf16. One traversal of the matrix (a read of every element, or a write) is S²·2 = 33.55 MB. Numerically safe softmax needs the row max, so the straightforward three-kernel version traverses five times:

Kernel	Reads	Writes	Traversals
1. row-max	S²	S maxima (≈0)	1
2. exp(x − max) + row-sum	S²	S²	2
3. divide by sum	S²	S²	2
Unfused total			5 × 33.55 = 167.8 MB → 50.1 µs

Fused: one program per row. A 4,096-element bf16 row is 8 KB — it fits in registers, let alone the 228 KB of SMEM. Load the row once, compute tl.max, tl.exp, tl.sum, and the divide entirely on-chip, store once: 2 traversals, 67.1 MB, 20.0 µs. A 2.5× traffic cut — 5 passes to 2 — with zero FLOPs changed and the ratio independent of S. Check the FLOP side to see why fusion, not tiling, was the right tool: softmax does ~5 ops per element against 4 bytes moved even when fused — intensity ≈ 1.25 FLOP/byte, still two orders of magnitude under the ridge. No restructuring makes softmax compute-bound; its only lever is fewer bytes, and 2 passes is the floor if the S×S matrix must exist in HBM at all. Refusing that premise is Monday's lecture (GPU kernels, Oct 5): FlashAttention fuses the softmax into the attention GEMMs so the S×S matrix is never materialized, and the 167.8 MB above is the before-picture it is measured against. We stop at the seam; Oct 5 crosses it.

The second payoff: launches. Lecture 2 counted ~1,500 kernel launches per training step at ≈5 µs each and showed the step goes launch-bound below ~75 tokens — precisely the decode regime, where a forward-only pass still issues about 490 launches for 2.45 ms of dispatch against a 4.0 ms memory floor. Every fusion is one fewer launch, so in decode fusion pays twice: once in bytes and once in fixed overhead — and exercise 5 shows the second payment is the one actually available, because the weight bytes are irreducible and the dispatch is not.

When fusion does not help. Fusion deletes bytes, so it pays only where bytes bind. A big prefill GEMM sits at intensity in the hundreds to thousands (Lecture 2 put the 4K prefill at ≈4,100 against a ridge of 295), its runtime set by arithmetic, so deleting boundary traffic that is already a rounding error changes nothing. GEMM-into-GEMM fails mechanically: one output tile of GEMM2 needs a full row-strip of GEMM1's output, not a tile, and that strip does not fit in 228 KB at LLM sizes — and the two GEMMs want different tilings anyway. Hence frameworks fuse epilogues onto GEMMs but not GEMM chains. The famous exception proves the rule: attention's softmax has row structure admitting an online reformulation — a running max and running sum, corrected as new columns arrive — which is the one licence to fuse GEMM–softmax–GEMM, and Oct 5's business, not ours. Fusion also has costs when it works: a fatter epilogue raises register pressure, which can lower occupancy (§7.5) and slow the GEMM it was grafted onto; a fused op matches no library kernel, so you own its performance and its bugs; and the autotuning space multiplies.

The decision procedure, three lines:

1. Compute the intensity of the fused region (step 3 of §7.8's method).
2. If the unfused chain was bandwidth-bound and the fused working set fits on-chip → fuse.
3. If any member is compute-bound → fuse only its boundaries (epilogue/prologue), never its interior.

Fusion is tiling across operator boundaries: the "tile" is the intermediate tensor, and the reuse is "use it while you still have it." Administered to a compute-bound kernel it does nothing but complicate the code.

Instructor notes

Minutes: 9. Board: The five-traversal table live, then strike through three of the five. Write "167.8 → 67.1 MB" large; it is the section. The 0.43 ms epilogue number goes in a corner with "× 32 layers". Then the three-line procedure, verbatim. Ask the room: "Fused softmax is 2.5× less traffic. Why not 5×?" You still must read the input and write the output — 2 passes is the floor while the matrix lives in HBM. Then: "who refuses to let it live in HBM?" — Oct 5. Expect confusion: Students expect fusion to raise intensity into the compute-bound regime. Show the 1.25: fusion shrank the denominator, but softmax has no FLOPs to be bound by. If short on time: The softmax table and the three lines; the register-pressure caveat is on the page.

7.11 Autotuning: tile sizes are facts about hardware × shape

Reading-only; not scheduled. Composes §7.8–§7.10 with no new mechanism.

There is no universal BLOCK. The optimum trades §7.8's traffic (wants big tiles) against register pressure and occupancy (want small ones), the SMEM budget, and wave quantization (a tile size that divides the shape into 7.05 waves wastes most of a wave). Every term is hardware-specific: an A100 — 2,039 GB/s, ridge 153, less shared memory per SM than the H100's 228 KB — balances them differently, so a config tuned on one card is a guess on the other.

Shape matters as much as hardware, and this is the LLM-specific point. Decode GEMMs are tall-skinny: M = the batch size B — anywhere from 1 to ~116 in Lecture 2's accounting — against K = N = 4,096. When M < BLOCK_M, the mask (§7.9) pads the tile, and the reuse of each loaded weight element is capped at M, not BLOCK_M. Rerun the intensity formula with the cap:

At M = 1 that is ≈1 FLOP/byte whatever the tile sizes — we have re-derived Lecture 2's "decode intensity ≈ B" from inside the kernel. No tile size can rescue batch-1 decode: the weight bytes have exactly one token to be reused against, and the fix is B, not BLOCK — batching (Oct 7), not kernel heroics. Kernels and batching are complements, not substitutes.

The mechanism, in Triton: decorate the kernel with @triton.autotune(configs=[...], key=['M', 'N', 'K']), listing candidate BLOCK_M/BLOCK_N/BLOCK_K triples with their num_warps and num_stages. Triton benchmarks every config the first time it sees a new key and caches the winner. The cost is honest: first-call latency per new shape, which is a real serving concern when shapes vary step to step — a decode server sees a new M whenever the batch composition changes, and bucketing shapes is part of the fix.

A kernel is finished when its tile sizes are a lookup table keyed on (GPU, shape) — and someone measured every entry.

Instructor notes

Minutes: 0 — reading-only. If five minutes appear, do only the M = 1 re-derivation: it ties the whole meeting back to Lecture 2's serving arithmetic in four lines. Ask the room (if run): "Assignment 4 gives you a serving system. Where does today's lecture show up in it?" Every custom kernel's config table, and every fusion the engine chose.

7.12 The other answer: the TPU's systolic array

Reading-only; not scheduled for class.

Google's TPU makes the same bet as the GPU — spend the die on arithmetic, hide latency with structure rather than caches and speculation — with the opposite control strategy. TPU v1 (the 2017 paper) is built around a systolic array: a 256×256 grid of 8-bit multiply-accumulate units through which operands are pumped, neighbour to neighbour, one step per cycle. Weights stay put, activations flow in from one edge, partial sums out the other. No per-access address, no scheduler, no warp — the dataflow itself is the latency hiding, because every value's next hop is wired.

The paper's numbers, internally consistent and worth checking: 256 × 256 = 65,536 MACs at 700 MHz → 65,536 × 2 × 0.7e9 ≈ 92 TOPS (int8). The array is fed from a 24 MiB unified buffer of software-managed SRAM (plus 4 MiB of accumulators) — the TPU's answer to shared memory, at nearly the H100's entire 30.1 MB — and backed by only 34 GB/s of DDR3. Divide: 92e12 ÷ 34e9 ≈ ~2,700 ops/byte of ridge point against the H100's 295. The TPU does not repudiate the roofline bet; it is the same wager with the intensity requirement an order of magnitude higher, and the paper's own roofline figure shows most of its 2015-era workloads pinned to the bandwidth diagonal as a result. (Later generations replaced DDR3 with HBM for exactly this reason; we do not quote their figures.)

The trade, stated symmetrically. The GPU pays area and energy for generality: any 32-lane computation runs, dynamic control flow is legal, and the taxes are scheduling overhead and divergence. The TPU pays rigidity for efficiency: matmul-shaped work runs at remarkable ops-per-joule, the compiler must produce a fully static schedule, and divergence is not expensive — it is inexpressible. Irregular workloads, the TPU's tax, simply cannot run well; divergent code, the GPU's tax, at least runs.

For this course the punchline is that both architectures answer Lecture 2's roofline identically — raise arithmetic intensity via on-chip reuse — which is why §7.8's tiling ideas transfer across accelerators even when not a line of the code does.

7.13 Reading the datasheet: where the roofline lives in silicon

Reading-only; not scheduled for class.

Reassemble Lecture 2's roofline from today's parts. The flat ceiling, 989 TFLOP/s, is §7.7's tensor cores — 132 SMs × 7.5 TFLOP/s of MMA hardware. The diagonal, 3,350 GB/s, is §7.3's HBM stacks and the pins that connect them. The ridge point, 295 FLOP/byte, is nothing but the ratio of two independently engineered subsystems — which is why it is a choice: NVIDIA could have shipped fewer tensor cores or more HBM stacks and moved it. And reaching either ceiling is conditional: the diagonal requires §7.4's thousands of in-flight loads, and the flat roof requires §7.8's on-chip operand reuse.

The generational trend now has a physical reading. A100 to H100: compute 3.2×, bandwidth 1.6×, ridge 1.9× (Lecture 2). Compute grows by widening tensor cores, cheap in die area; bandwidth grows by adding HBM stacks and pins, expensive in packaging. The bandwidth wall is a packaging-economics fact, which is why the ridge keeps rising and why this course's memory-side lectures age better than any chip.

The checklist to leave with, applicable to any accelerator datasheet: extract (a) peak ÷ bandwidth → the ridge point, hence how much batching or fusion the chip demands before it earns its price; (b) SRAM per compute unit → the tile budget, hence how much reuse a kernel can stage; (c) the resident-parallelism limit against bandwidth × latency → whether quoted bandwidth is even reachable, and with how much headroom. Applied to the three machines this lecture touched:

Checklist item	H100 SXM	A100 80GB SXM	TPU v1
(a) ridge = peak ÷ bandwidth	989 ÷ 3.35 → 295	312 ÷ 2.039 → 153	92 ÷ 0.034 → ~2,700
(b) staging SRAM per unit	228 KB/SM	164 KB/SM	24 MiB unified buffer
(c) latency-hiding supply vs demand	64 warps vs ~55 (§7.4)	64 warps vs ~53 (ex. 2)	static schedule — hidden by construction

Those three rows predict most of what the rest of Part II measures, and filling them in for a chip you have never programmed is the fastest sanity check on any accelerator pitch you will ever hear.

Lecture 2 gave you the roofline as a graph. Today it became a machine, and then it became your code.

Discussion seeds

1. Occupancy, argued honestly. You inherit a kernel at 25% occupancy that already saturates HBM. A colleague "fixes" it to 75% by halving BLOCK. Using §7.4's demand of ~55 outstanding loads per SM and §7.8's BLOCK/2 intensity, say what that trade cost — and name the one measurement that settles the argument without a slide.
2. Where the missing 47% lives. Square SMEM tiling tops out near 43% of peak; vendor GEMMs exceed 90%. Given one week, which of §7.8's three rungs — coalescing, register tiling, L2 reuse across programs — do you implement first, and what do you measure to know you bought the right one?
3. Fusion against the SLO. A fat epilogue raises register pressure and can lower occupancy (§7.5), while every fusion deletes a ≈5 µs launch. On the decode path — Sep 28 put the SLO-capped batch at B = 164 — which effect dominates? Design the experiment that answers it in one afternoon.
4. The layout you did not choose. §7.6's 1/32 is pure address arithmetic, and Sep 28 turned attention into a block-table gather. Which of the two fixes — change the layout, or stage through shared memory — applies to a paged KV cache, and what does Lecture 10 §10.3's block-size table have to do with your answer?
5. Moving the ridge. 295 is a ratio of two independently engineered subsystems, so NVIDIA could have shipped fewer tensor cores or more HBM stacks. Running the agent workload of Lectures 4–5 (≈60:1 prefill:decode), which way would you want the ridge moved, what would you pay for it, and which of today's sections would change most?
6. What you would write by hand. Given §7.9's ownership table, name the one kernel in your Assignment 4 serving stack you would write yourself, the one you would take from a library, and the one you would delete by fusing it into a neighbour.

Key takeaways

Numbers worth memorizing

Quantity	Value	Source
SMs / warps per SM / threads per SM	132 / 64 / 2,048	H100; §7.1
Register file / shared memory	256 KB/SM → 33.8 MB; 228 KB/SM → 30.1 MB, ~30 cycles	132 × each
L2 / HBM3	50 MB / 80 GB at 3,350 GB/s, ~500 cycles	datasheet
On-chip total vs 7B weights	~114 MB vs 13.5 GB → 118×	33.8 + 30.1 + 50
In-flight demand to saturate HBM	~0.93 MB ≈ 7,270 × 128 B ≈ 55 per SM	1,861 B/cycle × 500
Coalescing worst case (fp32, stride ≥ 32)	1/32 → ~105 GB/s effective	128 of 4,096 B
Tensor vs vector peak	989 vs ~67 TFLOP/s ≈ 15×	datasheet
Ridge point	295 FLOP/byte bf16; 591 fp8	989/3.35; 1,979/3.35
Naive matmul intensity (zero-reuse)	0.5 FLOP/byte → 0.17% of peak	2MNK ÷ 4MNK; 0.5/295
Tiled HBM traffic / square-tile intensity	b·MNK·(1/BM + 1/BN); BLOCK/2 FLOP/byte	§7.8 derivation
4096³ on H100	137.4 GFLOP; floor 0.139 ms; naive 82 ms; BLOCK=128 → 2.15 GB, 0.64 ms	§7.8
SMEM-tiling ceiling	BLOCK ≈ 128–256 → ≈429 TFLOP/s ≈ 43% of peak	register file: 256²·4 B = 256 KB
Fused softmax	5 passes → 2 passes; 2.5×, size-independent	§7.10
One unfused activation round-trip (32,768 tok, 7B up-proj)	1.443 GB → 0.43 ms, ×32 layers	32,768·11,008·2·2 B
Decode GEMM intensity cap	≈ M (= B) FLOP/byte at bf16, any tiles	§7.11

Self-check

Exercises

1. Occupancy is a min(). A kernel uses 64 registers/thread, 256-thread blocks, and 32 KB of SMEM per block. Compute the occupancy, then decide whether it can saturate HBM under §7.4's model. Solution sketch: Registers: 65,536 ÷ 64 = 1,024 threads = 4 blocks. SMEM: ⌊228 ÷ 32⌋ = 7 blocks = 1,792 threads. Caps: 2,048 threads / 32 blocks — not binding. Minimum is registers → 1,024 threads = 32 warps = 50%. Against §7.4's demand of ~55 outstanding loads per SM, 32 warps at one load each falls short; the kernel needs ~55/32 ≈ 2 in-flight loads per warp (unroll by two) to saturate — exactly what num_stages buys in §7.9.
2. Little's law on the A100. The A100 has 2,039 GB/s and 108 SMs; assume ~500 cycles of latency at a ~1.4 GHz working clock (stated assumption, same status as our 1.8 GHz). How many outstanding 128-B transactions per SM does saturation demand? Solution sketch: 500 ÷ 1.4e9 = 357 ns; 2.039e12 × 357e-9 ≈ 728 KB in flight; ÷ 128 ≈ 5,690 transactions; ÷ 108 SMs ≈ 53 per SM — essentially the H100's 55. The per-SM latency-hiding burden is a constant of the DRAM era, not of one chip: bandwidth per SM and latency have moved together.
3. Coalescing accounting at fp16. A warp reads one column of a 4,096×4,096 row-major fp16 matrix — one fp16 per lane, address stride 4,096 × 2 = 8,192 B. Compute the efficiency, the effective bandwidth, and the cost of the SMEM fix. Solution sketch: Every lane is in its own segment: moved = 32 × 128 = 4,096 B; useful = 32 × 2 = 64 B → 1/64 — worse than the fp32 case, because smaller elements waste more of each segment. Effective bandwidth ≈ 3,350 ÷ 64 ≈ 52 GB/s. Fix: stage a 32×32 fp16 tile through SMEM (load coalesced row-wise, read column-wise from banks): 32 × 32 × 2 = 2 KB, about 0.9% of the 228 KB budget — a 64× bandwidth recovery for under 1% of one resource.
4. 8192³ matmul, full accounting. For M = N = K = 8,192 bf16 on one H100, compute the FLOPs and compute floor, the naive traffic and time, and the tiled traffic and time at BLOCK = 128 and 256. Is the BLOCK = 256 kernel compute-bound? What closes the remaining gap? Solution sketch: FLOPs = 2·8192³ = 1.10 TFLOP → floor 1.10e12 ÷ 989e12 = 1.11 ms. Naive: 4·8192³ = 2.20 TB → 656 ms (8× the 4096³ case in every column). BLOCK = 128: bytes = 2·8192³·2/128 = 17.2 GB → 5.13 ms; BLOCK = 256: 8.59 GB → 2.56 ms. At intensity 128 < 295 it is still memory-bound (2.56 ms > 1.11 ms). The rest comes from §7.8's L2 rung — grouped launch order keeping shared strips resident in the 50 MB L2 — which raises effective reuse beyond what any single program's registers can hold.
5. Launches in the decode regime. Using Lecture 2 §2.9's figures — a forward-only eager pass over 32 layers issues on the order of 490 launches at ≈5 µs each — compare launch overhead with the 4.0 ms batch-1 decode floor. Then say what cutting the launch count 3× saves, and what it cannot save. Solution sketch: 490 × 5 µs = 2.45 ms of dispatch against a 4.0 ms memory floor, so a naive eager decode step costs about 6.45 ms and dispatch is 38% of it — this is Lecture 2's "launch overhead alone can be half the budget of a decode step". Fusing to ~163 launches saves 327 × 5 µs = 1.64 ms/step, a 25% cut in step time, landing at 4.65 ms, within 16% of the floor. What fusion cannot save is the floor: the 13.5 GB of weights is read once per token regardless, so nothing goes below 4.0 ms. Moral: in decode, fusion pays through two meters, and the dispatch meter is the one you can drive to zero — which is why CUDA graphs exist, and why the residual gap is a bytes problem rather than a dispatch problem.

Reading guide

Optional — CUDA C Programming Guide. A reference to keep open through Part II, not a text to finish. Read the Programming Model chapter for threads/blocks/grids and the Hardware Implementation chapter for SIMT, warps, and scheduling; then the performance-guidelines material on memory coalescing and occupancy. Skip the API reference and everything driver-API. The compute-capability appendix is the authoritative source for the per-architecture limits we used today (64 warps, 65,536 registers, 228 KB SMEM). Question to hold: for each rule the guide states — "prefer coalesced access," "avoid divergence," "balance occupancy against registers" — which line of today's arithmetic is it the compressed version of?

Optional — Dissecting Volta and Dissecting Hopper. Read for method: NVIDIA publishes capacities but not latencies, so the authors reconstruct the hierarchy — line sizes, per-level latencies, register-bank structure, tensor-core behavior — from hand-built microbenchmarks: pointer-chasing for latency, stride sweeps for geometry. The specific cycle counts belong to their generations; carry away the technique and the ordering of levels. Question to hold: which of today's board numbers could you verify yourself with a 50-line pointer-chase kernel — and which (peak FLOP/s, say) need a different experiment entirely?

Optional — Triton docs. Not to be read — to be run, before Oct 5. Tutorial 01 (vector add) is program ids and masks in ten lines; tutorial 02 (fused softmax) is §7.10 executable; tutorial 03 (matmul) is §7.9 with the grouped launch ordering §7.8's L2 rung described — find the swizzle and say which memory it is tiling for. Hold one question through all three: for each line, which term of §7.8's byte formula does it implement? Skip the compiler-internals and backend pages.

Optional — Triton (MAPL 2019) (Tillet, Kung, Cox). The introduction and language section, for the thesis: make the block the primitive and a compiler can own the layout, shared-memory, and scheduling decisions CUDA delegates to the programmer. Skim the compiler-pass and evaluation sections for shape only — the result to retain is competitive-with-vendor-library performance on the GEMMs tested, not any specific figure. The paper's surface syntax (Triton-C) predates today's Python tl.* API; the block-programming model is unchanged.

Optional — How to optimize a CUDA matmul kernel (Boehm). The empirical companion to §7.8's ladder: one hand-written kernel per rung, each measured. His numbers are fp32 on his GPU, not our bf16/H100 — do not transplant the GFLOP figures; take the ordering and the relative magnitudes: coalescing and shared-memory tiling are the giant steps, and the last 2× is many small ones. Question to hold: which rungs did Triton's compiler absorb (§7.9's table), and which — tile-size choice — survive as your problem?

Optional — TPU. Read §2 for the architecture — the 256×256 MAC array, the unified buffer, and the framing that design decisions were driven by 99th-percentile latency — and the roofline figures with their analysis. Skim the workload table (2015 Google ran mostly MLPs, a useful corrective) and skip the CPU/GPU comparison controversies and die-photo detail. This is v1: inference-only, int8, DDR3-era; treat its constants as historical and its argument as current. Question to hold: the paper's own roofline shows most of its workloads bandwidth-bound under a ~2,700 ops/byte ridge — is the TPU a refutation of the GPU's design, or the same bet with the intensity requirement turned up?

Looking ahead

Monday (Efficient LLM computing: GPU kernels, Oct 5) is the direct sequel and composes today's two moves into one kernel: FlashAttention is fused softmax plus tiling, arranged so the S×S score matrix never touches HBM at all. §7.10's 167.8 MB is the before-picture, the online max-and-sum is the licence, and the roofline paper — read properly there rather than borrowed from Lecture 2 as we did today — supplies the units. Bring the Triton tutorials; run 01–03 first.

Then the serving layer, where the kernels this lecture explains are what the scheduler dispatches: Oct 7 (batching and scheduling) and Oct 14 (disaggregation and chunked prefill), where §7.11's "the lever is B, not BLOCK" becomes a scheduling policy, then Oct 26 on routing. Nov 9 takes the other lever, fewer bytes per weight. The longer arc: paged KV allocation (Sep 23), KV-cache optimization (Oct 28), prefix reuse (Nov 4), and quantization (Nov 9) all manage the HBM row of §7.3's table, and the final project (announced Oct 26, report due Dec 8) is that row taken personally.

Assignment 2 (design an agent) is due Sunday, 11:59pm, and Assignment 3 (optimize the agent) goes out the Monday after, due Sun Oct 25 — A3 still sits outside the API, so none of today's levers are yours yet. They arrive with Assignment 4 (serve your own agent, out Oct 19, due Tues Nov 10), and Assignment 5 (optimize the full stack, out Nov 11, due Wed Dec 2) is where they pay.

One sentence to carry out of the room: the FLOPs in your kernel were decided by whoever wrote the model, and the bytes were decided by you — so every performance argument you will have for the rest of this course is an argument about step 2.