Memory hierarchy, roofline, arithmetic intensity

Neither roof, which means the binding constraint is not in this roofline at all. Roofline gives an upper bound. Sitting under both roofs tells you only that something outside the model is limiting you. Ranked suspects:

1. You are bound at a level this roofline cannot see. Shared-memory or L2 bandwidth, or register-file bandwidth. Draw the hierarchical roofline before drawing conclusions.
2. Latency-bound, not bandwidth-bound. Insufficient memory-level parallelism to satisfy Little’s Law: low occupancy, long dependent chains, address arithmetic on the critical path. The pipes are wide open and you are not putting enough in flight.
3. Issue-bound. The limiter is instruction throughput on a pipe that is neither FMA nor memory: address math, predication, shared-memory instructions, or the epilogue.
4. Wave quantization. A partly-empty final wave dilutes the whole-kernel average. 132 SMs and 140 blocks means a second wave running at 6% utilization.
5. Synchronization. __syncthreads() in an inner loop serializes what the roofline assumes is concurrent.

The discipline worth internalizing: roofline names what you are not limited by. It never names the limiter.

Count bytes at the binding boundary, not FLOPs. Two passes over $N$ bytes, each reading and writing: $4N$ total. Fused: $2N$. Ceiling is 2×, and that is the whole answer at the HBM level.

Reality lands below it for three reasons and above it for one. Below: if $N$ is small enough that the intermediate stayed resident in L2, the second pass never touched HBM, so fusion bought you nothing at that boundary. On H100, with 50 MB of L2, anything under ~25 MB of working set is already in this regime. Fusion also raises register pressure, which lowers occupancy, which lowers memory-level parallelism, which lowers achieved bandwidth, so the fused kernel can run at a worse point on the same roof. Above: you also eliminated a kernel launch, roughly 3–5 µs, which dominates everything else for small $N$.

Generalized: the ceiling on any fusion is bytes-before over bytes-after, evaluated at whichever boundary is actually binding. The same arithmetic sizes recomputation-versus-storage tradeoffs, FlashAttention included.

Because in decode, arithmetic intensity is tokens-verified-per-weight-load. Reading the weights is the entire cost. The FLOPs are free, since you are sitting 300× below the ridge with all that compute idle. Verifying $k$ drafted tokens in a single forward pass loads the weights exactly once and does $k\times$ the arithmetic: $I$ goes from ~1 to ~$k$, and you move right along the memory roof at constant byte cost. You are spending headroom you already owned.

Four ceilings, in the order you hit them:

1. Acceptance. Expected accepted tokens per round is set by draft-target agreement, typically 2–4 for a well-matched draft. Everything past that is discarded work.
2. The draft model’s own serial decode, which is itself batch-1 memory-bound and sits on the critical path.
3. KV cache traffic. The “bytes are free” premise holds only while weights dominate. At long context the KV cache becomes the larger read and it scales with the tokens you process, so the $I \approx k$ argument decays.
4. The ridge itself. Once effective batch × $k$ approaches ~300 you become compute-bound and further $k$ buys nothing.

The same one-line test explains GQA, quantization, and continuous batching: they are all byte-count or intensity moves. None of them is a FLOP move. In a memory-bound regime, FLOP moves are not interesting.

Close to nothing, and the direction it points is wrong. Three distinct failures:

• It is asymptotic. $\min(\pi, \beta I)$ presumes enough concurrent work to reach steady state and hide latency. At batch 1 the machine is overwhelmingly idle and you are under both roofs by a wide margin, in a regime the model does not describe.
• Its only lever is unavailable. “Raise $I$” means batch or reuse. Batch is pinned at 1 by the problem statement, and there is no reuse to find inside a single sample.
• It assumes fixed costs amortize to zero. At 5 µs nothing amortizes. Kernel launch alone is 3–5 µs. A PCIe round trip is 1–2 µs. Either one consumes the budget before a single FLOP retires.

The replacement model is $\text{latency} = \sum(\text{dependent stage latencies}) + \text{fixed overheads}$, with the objective being depth of the critical path measured in cycles or gate delays, not bytes moved. That reframing is the content of Part B, and it is why this workload class ends up on an FPGA or ASIC with weights held in on-chip SRAM and registers. Flattening the memory hierarchy to a single level does not make “bytes over bandwidth” smaller. It removes the term from the equation.

Measure $I$ at the binding boundary and compare to $I^{*}$. Below the ridge, performance scales roughly linearly in $\beta$ and not at all in $\pi$. Above it, the reverse. That is the model’s answer, and it is worth exactly as much as your byte accounting.

Two empirical tests that need no model at all. Clock sweep: pin the memory clock and sweep the SM clock, then swap. The slope of each curve tells you which roof you are on. Byte ablation: halve the bytes without touching the math, by quantizing weights fp16 → fp8. If the kernel gets ~2× faster you were memory-bound. If it barely moves, you were not, and you just learned something more reliable than any profiler counter.