Background. CUDA's model is a hierarchy. Threads run in warps, warps in blocks, blocks in a grid, and every level has memory attached to it, from registers and shared memory out to L2 and global. This guide is where NVIDIA defines that model and the rules for using it.

Problem. The API does not tell you what anything costs. Whether a load coalesces, how much shared memory a block may claim, and how block size caps the warps resident on a streaming multiprocessor are rules that live outside the function signatures, and getting one wrong surfaces only as an unexplained slowdown.

Key idea. The guide is the normative statement of those rules, organized by level of the hierarchy. Treat it as reference, not reading. Know where the memory-hierarchy and occupancy sections are and go to them at the point of a decision.

Why it matters. Every CUDA kernel in PyTorch, CUTLASS, cuDNN, and FlashAttention is written against this document, and Triton and CUDA Python both inherit its memory model and launch configuration. The attention kernels and compilers later in the meeting hide the thread hierarchy; they do not remove it, and when a tiled or generated kernel underperforms the explanation is usually in here.

Background. By 2009 multicore processors had more arithmetic throughput than their memory systems could feed, and tuning was done per machine and per kernel with no shared way to state what the ceiling was.

Problem. A machine's peak FLOP/s says nothing about the ceiling for one kernel. Without knowing whether memory traffic or arithmetic is the binding constraint, you can spend weeks optimizing the side that was never the limit.

Key idea. Plot attainable performance against operational intensity, the FLOPs a kernel does per byte it moves from memory. The machine contributes two ceilings, a bandwidth slope and a flat compute limit, and which one a kernel sits under names what to fix.

Why it matters. It gives you the units for the rest of the course, and it is the right first step in Assignment 5. Most later results in this field are a claim about raising one of the two ceilings or about moving a kernel along the intensity axis.

Adoption. Nsight Compute and Intel Advisor both ship a roofline chart as a built-in profiler view, and NERSC's Empirical Roofline Toolkit measures the ceilings a real machine reaches.

Background. The roofline divides a kernel's time between arithmetic and memory traffic. A deep learning program is a Python process issuing thousands of small operators, so it spends time in a third place: getting work onto the GPU at all.

Problem. Such code often sits in neither roofline regime. Python dispatch and per-kernel launch cost can leave the GPU idle between operators, and no amount of tiling or fusion touches that time.

Key idea. Diagnose first, against three regimes rather than two: compute-bound, bandwidth-bound, and overhead-bound. Each has its own test and its own fix, so decide which one you are fighting before you touch anything.

Why it matters. It is the roofline in the form these workloads need, and it promotes overhead from measurement noise to a regime with a name. It also explains why fusion and graph capture are two different tools rather than two speedups.

Background. Attention over a sequence of length N forms an N×N matrix of scores, softmaxes it, and multiplies it by the values. A standard implementation writes that matrix to HBM and reads it back for each of those steps.

Problem. The matrix grows quadratically in sequence length, and the traffic to HBM, not the arithmetic, sets the runtime. Memory also caps the sequence length you can train on, so the context limit came from an implementation choice rather than from the model.

Key idea. Tile the computation so a block of queries meets a block of keys in SRAM, accumulate the softmax with a running normalizer, and never write the N×N score matrix to HBM at all. The output is exact attention, not an approximation.

Findings. Against the paper's baselines, FlashAttention delivers a 15% wall-clock speedup on BERT-large over the MLPerf 1.1 record, a 3× speedup on GPT-2 at length 1,000, and a 2.4× speedup on Long Range Arena tasks at 1,000-4,000. Its block-sparse extension reaches 63.1% accuracy on Path-256 at length 64,000.

Why it matters. It is the clearest demonstration in the course that IO, not FLOPs, is the budget. An algorithm with identical arithmetic and identical output runs faster because it moved less data, which is the lesson every later kernel in this section applies.

Adoption. PyTorch's scaled_dot_product_attention dispatches to a FlashAttention kernel, and vLLM, SGLang, TensorRT-LLM, and Hugging Face Transformers all call one; xformers and FlashInfer ship variants.

Background. FlashAttention already keeps the score matrix out of HBM, so what remains is internal to the kernel: how the work is split across thread blocks and warps, and how much of the arithmetic runs somewhere other than the tensor cores.

Problem. The original kernel parallelizes over batch and heads only, which leaves SMs idle when the batch is small and the sequence is long. It also does its rescaling on the general-purpose units, whose throughput per operation sits far below the tensor cores'.

Key idea. Better work partitioning and fewer non-matmul FLOPs. Parallelize over the sequence dimension as well, split each tile's work across warps so they no longer exchange partial results through shared memory, and defer the softmax division so fewer rescaling operations run per tile.

Findings. On A100, FlashAttention-2 runs about 2× faster than FlashAttention and reaches 50-73% of theoretical peak throughput. End-to-end GPT-style training reaches 225 TFLOP/s per A100, or 72% model FLOPs utilization.

Why it matters. It shows that removing the memory bottleneck is only the first step. Once IO is handled, the remaining wins come from where work is placed: across blocks, across warps, and on which functional unit. That is the vocabulary the rest of this section uses.

Adoption. PyTorch's scaled_dot_product_attention, Hugging Face Transformers' flash_attention_2 implementation, Megatron-LM, DeepSpeed, and vLLM's prefill implement FlashAttention-2.

Background. Hopper changes the execution model the earlier kernels assume. It adds a tensor memory accelerator for asynchronous copies, warpgroup-wide matmul instructions that issue asynchronously, and FP8 tensor cores.

Problem. A kernel written for synchronous loads and per-warp matmuls cannot keep those units busy, because the copy, the matmul, and the softmax run in sequence when the hardware could run them at once. FP8 adds a second difficulty, since attention's exponentials and outliers do not survive naive quantization.

Key idea. The same algorithm rewritten for a new memory and execution model. Warps specialize into producers that issue asynchronous copies and consumers that run the matmuls, the softmax of one tile overlaps the matmul of the next, and FP8 is applied with block quantization and incoherent processing so accuracy holds.

Findings. On H100, the FP16 kernel is 1.5-2.0× faster than FlashAttention-2 and reaches 740 TFLOP/s, or 75% utilization. The FP8 path approaches 1.2 PFLOP/s and has 2.6× lower numerical error than the paper's baseline FP8 attention.

Why it matters. It is the case study in what a hardware generation costs a kernel author. The mathematics did not change; the code did, because the units that mattered became asynchronous. Expect that cost at every new architecture, and see FlashAttention-4 for the Blackwell instance.

Adoption. vLLM builds and installs a dedicated FlashAttention-3 component for its Hopper attention backend.

Background. FlashAttention-3 is tuned for Hopper, where the balance among tensor-core throughput, SRAM bandwidth, and the unit that computes exponentials held a particular shape. Blackwell breaks that balance: tensor-core throughput grew and the other two did not.

Problem. When one unit gets faster and its neighbors do not, the bottleneck moves. Attention's softmax needs exponentials and needs tiles resident in SRAM, so a pipeline balanced for Hopper spends its Blackwell time waiting on the units that did not scale.

Key idea. Co-design the algorithm and the kernel pipeline for the new ratios instead of porting the Hopper schedule. The subtitle names both halves, and the point is that they move together: a pipeline balanced for one set of unit throughputs is wrong for another.

Findings. On B200 with BF16, FlashAttention-4 reaches 1,613 TFLOP/s, or 71% utilization, and is up to 1.3× faster than cuDNN 9.13 and 2.7× faster than Triton. Its CuTe-DSL implementation compiles 20-30× faster than the reported C++ template path.

Why it matters. A concrete lesson in redesigning a kernel when hardware scales asymmetrically. Peak FLOP/s is the figure vendors advertise and the least useful one here; what sets the rewrite is which units failed to keep up with it.

Background. The FlashAttention papers report what a finished kernel achieves, and the asymmetric scaling FlashAttention-4 names is the reason its Blackwell schedule differs from the Hopper one. On B200 the tensor cores roughly doubled while the unit that computes exponentials did not move, so softmax costs about as many cycles as the matrix multiplies despite far fewer floating-point operations.

Problem. A paper reporting one number for a co-designed kernel does not say which of its choices bought the speedup, and step-by-step walkthroughs on this hardware existed for matrix multiplication rather than for attention. A reader who wants to write such a kernel, or to judge which optimization to try on a new shape, has no ordered account of what each one is worth.

Key idea. Build the kernel one optimization at a time and measure every step. Fourteen versions of a dense non-causal BF16 kernel at head dimension 128, each a single change from the one before it, take the same tile decomposition from a naive baseline to near the reference implementation, with the Blackwell mechanisms introduced where they are first needed: tcgen05 tensor-core instructions and the tensor memory that holds their accumulators, swizzled shared-memory layouts, asynchronous copies driven by barriers, warp specialization, and load and compute pipelines that overlap the softmax with the matrix multiplies.

Findings. Every version is scored as a fraction of stock FlashAttention-4 on the same B200, over the 4,000, 8,000, and 16,000-token shapes the FlashAttention-4 paper uses, and the ladder runs from 14.2% at the baseline to 94.4% at the end. The large individual gains come from shared-memory swizzling (15% to 26%), a second query tile per block sharing the same key and value loads (27% to 41%), the load pipeline (49%), the compute pipeline (58%), and two steps that attack the exponential unit directly, calling the hardware instruction rather than a wrapper (64%) and emulating a quarter of the exponentials in three chained fused multiply-adds to relieve pressure on that unit (70%). Two steps, moving the probability tile into tensor memory and the first warp specialization, bought almost nothing on their own and were kept because later steps depend on them.

Why it matters. It is the roofline lesson at instruction granularity, and the binding resource is not the one the hardware is sold on: the specialized unit that failed to scale sets the runtime, and six of the fourteen steps exist to hide it. It also shows that an optimization ladder does not decompose into independent wins, since a step worth nothing measured alone is what makes the next one possible. For anyone writing a kernel in this course, it is the closest thing to a worked example of the method, including the discipline of scoring each attempt against a fixed reference on the same GPU.

Background. A serving engine never hands attention a clean batch. Requests arrive at different lengths, share prefixes, get preempted, and their KV cache lives in pages scattered across HBM. Each engine grew its own attention kernels to match its own layout.

Problem. Kernel work and scheduling work were being solved separately, and both suffered. A kernel written for one KV layout does not transfer, variants such as grouped-query attention, sliding windows, and custom masks each fork the code, and ragged batch shapes ruin the load balance a fixed tile assignment assumes.

Key idea. Give the kernel layer one interface. Represent every KV layout, paged or shared-prefix, as block-sparse; generate each attention variant from a template by just-in-time compilation rather than writing it by hand; and assign tiles with a load balancer that absorbs varying sequence lengths while staying compatible with CUDA graphs.

Findings. Against compiler backends in the paper's serving benchmark, FlashInfer reduces inter-token latency by 29-69%. It cuts latency by 28-30% for long-context inference and speeds up serving with parallel generation by 13-17%.

Why it matters. This is where the kernel layer and the scheduler layer meet. The lesson is that peak kernel throughput on a uniform batch is the wrong target for serving; what you get to keep is the throughput that survives a ragged batch and a paged cache.

Adoption. FlashInfer is the default attention backend in SGLang and a supported backend in vLLM and MLC-LLM.

Background. Paged attention keeps the KV cache in fixed-size blocks reached through a block table, so a batch of requests with different lengths reads a ragged set of blocks rather than one dense tensor. The kernels that do this are CUDA work.

Problem. A TPU has no warps and a different memory system, so the CUDA kernel cannot be ported line by line. That leaves the question of which parts of a paged-attention kernel are the algorithm and which parts are CUDA-specific folklore.

Key idea. Implement paged attention natively for TPU: the same block-table KV layout, a ragged batch so sequences of different lengths are served in one kernel rather than padded to a common length, and a tiling strategy chosen for the TPU memory system.

Findings. With Llama 3 8B on TPU7x, the kernel reaches up to 86% memory-bandwidth utilization during decode and 73% model FLOPs utilization during prefill.

Why it matters. It is a forcing function for separating an algorithm from its implementation. If paged attention survives a move to hardware without warps, the block table and the tile-at-a-time traversal are the portable content and the rest was tuning for one vendor.

Adoption. vLLM's TPU backend runs a ragged paged attention kernel of this shape, written in Pallas, against the same block-table KV layout the GPU backends use.

Background. Tensor parallelism splits a layer across devices to hold weights that do not fit on one, and every split layer ends in a collective that recombines the partial results. §7.20 counts 64 of them per decode iteration for the reference model.

Problem. Those collectives sit on the dependency path, since the next layer cannot start until the current one's results are combined. At decode the messages are small, so their cost is latency rather than bandwidth, and it does not fall as the degree rises even though the compute per device does.

Key idea. Decompose the communication and the computation into pieces much finer than a layer, then fuse those pieces into a single kernel, so one tile's data can be sent while other tiles are still being computed. Overlap becomes a property of the kernel rather than something a runtime arranges around it.

Findings. The paper reports up to 96% of communication hidden inside a fused kernel, with 1.66× prefill and 1.30× decoding speedups over vLLM on eight GPUs, and 1.24× over Megatron-LM for training on 128.

Why it matters. It is the concrete form of §7.20's third consequence, that an engine can overlap a layer's collective with arithmetic where the dependence allows, and it names the price. The overlap has to be built into the kernel, so it arrives per fused operation rather than as a setting you turn on.

Background. Frameworks in 2018 ran a model by dispatching its operators to vendor kernel libraries such as cuDNN. A model therefore ran fast only on hardware whose vendor had written and tuned kernels for the operators that model used.

Problem. Every new operator, fusion, or accelerator needs another hand-written kernel, and the number of operator and hardware pairs grows as a product. A graph-level compiler alone cannot fix this, because the performance is decided inside the operators.

Key idea. Separate what a tensor operator computes from how it executes. The compute expression is written once; a schedule then specifies tiling, loop order, vectorization, and memory placement, and code is generated per target from the pair. Search over schedules replaces hand tuning.

Findings. Across low-power CPUs, mobile GPUs, server GPUs, and an FPGA accelerator, generated implementations are competitive with the paper's state-of-the-art hand-tuned libraries. The result establishes performance portability across unlike targets rather than a win on one device.

Why it matters. The compute and schedule split is the organizing idea of the kernel languages that follow, Triton and TileLang included, and it names the layer at which kernel performance is decided. It also sets up the auto-tuning problem the next reading attacks.

Adoption. Apache TVM ships as a real compiler stack: MLC-LLM builds its on-device LLM runtime on TVM Unity, and AWS SageMaker Neo compiles customer models with it.

Background. A GPU tensor program exists at three levels at once: kernels launched on the device, thread blocks within a kernel, and threads within a block. Compilers usually optimize one level at a time and take the operator graph as given.

Problem. The best implementation sometimes requires changing all three levels together, for instance replacing a sequence of operators with an equivalent one that tiles differently. A pass that rewrites a single level cannot reach those programs, because the win only appears once every level changes.

Key idea. Search for the program rather than optimizing the given one. Candidate implementations are generated jointly across the kernel, thread-block, and thread levels, and each candidate is checked for equivalence with the original before it is accepted.

Findings. Across the paper's heavily optimized DNN workloads, Mirage improves performance by as much as 3.3× over existing compiler and library implementations.

Why it matters. It shows where performance still hides once per-operator kernels are individually tuned, which is in the boundaries between them. The unit of optimization becomes a region of the model instead of an operator, and that is the premise the next reading takes to its conclusion.

Background. A model normally runs as a sequence of kernel launches, at least one per operator, with a device-wide synchronization between them. At decode batch sizes each kernel is short, so launch and synchronization time is a real share of the step.

Problem. The launch boundary forbids the two things that would hide it. An operator cannot begin on the tiles that are ready while its predecessor finishes, and a collective cannot proceed while a compute kernel holds the device. Writing a megakernel by hand for one model is possible; doing it per model is not.

Key idea. Compile a whole multi-GPU model into one persistent megakernel whose in-kernel scheduler dispatches SM-level tasks from a task graph. Per-operator launch boundaries disappear, which makes cross-operator pipelining and compute and communication overlap expressible inside a single kernel.

Findings. Across the paper's multi-GPU inference workloads, MPK delivers an end-to-end inference speedup of up to 1.7× against kernel-per-operator serving systems.

Why it matters. It relocates scheduling from the CUDA runtime into the kernel, and it names the price: the runtime must carve SMs into workers and schedulers and size the task graph to the physical SM count, so the compiled artifact is tied to the GPU it was built for.

Background. Work reaches a GPU one launch at a time, and each launch costs host time whether or not the sequence repeats. A decode step issues nearly the same sequence of kernels on every token it produces.

Problem. Repeating a known sequence still pays to describe it. The ordinary launch API offers no way to tell the driver that this iteration's work has the same shape as the last, so the description is rebuilt and resubmitted every step. At a decode step measured in single-digit milliseconds, the host can become the thing the GPU waits for.

Key idea. Separate defining the work from running it. A graph records a sequence of operations and their dependencies once, either by capturing a stream or by building the nodes directly, and the instantiated graph is then launched as one submission.

Why it matters. Replay requires a recorded shape, so engines capture graphs for a fixed list of batch sizes and pad to the nearest one. Shape immutability is an API constraint; avoiding pointer updates is an implementation choice.

Adoption. vLLM, SGLang, and TensorRT-LLM all capture decode graphs over a fixed set of batch sizes, which is why their configuration exposes a capture list and a maximum captured batch.

Background. PyTorch runs a model one tensor operation at a time, dispatching each from Python as the program executes. That is why the framework is easy to debug, and it is also why every operation is a separate launch the host has to issue.

Problem. Capturing a whole-graph representation without giving up eager semantics had defeated earlier attempts. A tracer that records one execution specializes on whatever it happened to observe, and Python's data-dependent control flow, mutation, and side effects mean the recording can be silently wrong on the next input.

Key idea. Intercept CPython bytecode before it runs, extract the tensor operations into a graph, and attach guards stating exactly which properties of the observed state the graph depends on. A guard failure retraces or falls back to Python, so correctness never rests on one trace being universally valid. A backend then lowers the graph, emitting Triton for GPUs.

Findings. Across more than 180 real-world models on an NVIDIA A100, TorchInductor delivers geometric-mean speedups of 2.27× for inference and 1.41× for training over eager PyTorch. These are broad framework benchmarks rather than measurements of a continuously batched serving engine.

Why it matters. Guards decide how often a served model recompiles. A serving engine sees a new shape when its batch or sequence length changes, so the guard set determines whether compilation is paid once at startup or repeatedly under load. The evaluation uses open-source benchmark suites rather than a served workload.

Background. The roofline arguments this course makes rest on a claim about where time goes: prefill is compute-bound and decode is bound by memory bandwidth. That claim is usually asserted from first principles rather than measured on a named machine.

Problem. A phase label does not tell you what the phase is doing. Time in decode can go to kernel mix, issue stalls, or scattered memory access, and which of those dominates decides whether a given optimization helps at all.

Key idea. Split inference into prefill and decode and trace each phase’s cost to kernel mix, roofline bounds, issue stalls, and memory access patterns, measured on an A100 server and on a Jetson Orin edge device.

Findings. On the A100, both phases spend 70-80% of cycles stalled, but for different reasons. Tensor parallelism cuts prefill time from 1.24 seconds to 0.66 while increasing decode latency because of communication and synchronization.

Why it matters. It is the empirical grounding for the roofline claims the rest of the course makes, and the two very different devices show which conclusions are about transformers and which are about one particular memory system.