Part IV: Pretraining at Scale
Chapter 20

Efficient Training Techniques

FlashAttention, activation recomputation, kernel fusion with Triton, low-precision formats, and the compiler-level optimizations that raise hardware utilization and squeeze more model out of every GPU-hour.
20 Exercises

Learning Objectives

1.Understand the GPU memory hierarchy and why most kernels are memory-bound.
2.Explain Model FLOPs Utilization and what limits it in practice.
3.Apply activation recomputation (gradient checkpointing) and reason about its trade-off.
4.Understand kernel fusion and why it reduces memory traffic.
5.Write a simple fused kernel in Triton and understand the programming model.
6.Explain how FlashAttention fuses the attention computation.
7.Apply low-precision formats (fp8) and the torch.compile graph compiler.
8.Combine the techniques into an efficient training configuration.
20.1

Chapter 16's 6ND rule counts the FLOPs a training run requires; Chapter 18 spread those FLOPs across thousands of GPUs. This chapter asks a different question: of the compute a GPU is theoretically capable of, how much is actually used for useful work? The answer is often surprisingly low — and closing that gap is worth millions of dollars on a frontier run.

Model FLOPs Utilization

Model FLOPs Utilization (MFU) is the fraction of a GPU's peak theoretical throughput that a training run actually achieves on useful model computation. A well-optimized large-model run reaches perhaps 40–55% MFU; a naive one can be far lower. The gap between achieved and peak is the target of every technique in this chapter.

textModel FLOPs Utilization
MFU = (model FLOPs per second) / (GPU peak FLOPs per second)

model FLOPs/s = 6 · N · (tokens processed per second)   # the 6ND rule

Example: 7B model, 40k tokens/s, on an H100 (~990 TFLOP/s bf16):
    MFU = (6 × 7e9 × 40000) / 990e12 ≈ 1.7e15 / 990e12 ≈ 0.48 (48%)

Why isn't MFU 100%? Because GPUs spend much of their time NOT doing matrix multiplications: waiting on memory, launching kernels, communicating between devices, and running non-matmul operations (softmax, normalization, activation) that the hardware's tensor cores cannot accelerate. The techniques in this chapter attack each of these inefficiencies.

A Few Percent of MFU Is a Fortune
On a nine-figure training run, raising MFU from 40% to 50% cuts wall-clock time and cost by 20%. At frontier scale that is tens of millions of dollars and weeks of time. This is why labs invest enormous engineering effort in the unglamorous work of kernels, memory layout, and compiler optimization.
Efficiency is not a nice-to-have at scale — it is a primary competitive lever. The same compute budget, used more efficiently, trains a meaningfully better model.
20.2

To understand efficiency, you must understand where time actually goes on a GPU. The central fact: arithmetic is cheap, but moving data between levels of the memory hierarchy is expensive. Most deep-learning operations are bottlenecked not by compute but by memory bandwidth — they are 'memory-bound', spending their time waiting for data rather than calculating.

The Hierarchy

LevelSize (H100)BandwidthRole
Registers~256 KB/SMFastestPer-thread scratch
SRAM (shared)~228 KB/SM~20 TB/sOn-chip tiles
HBM (global)80 GB~3.3 TB/sModel, activations
Inter-GPU (NVLink)~0.9 TB/sCross-device
Inter-node (IB)~0.1–0.4 TB/sCross-node

Each level down is far larger but far slower. SRAM is ~6× the bandwidth of HBM but thousands of times smaller. The art of efficient GPU code is keeping data in the fast levels as long as possible — doing as much work as you can on each byte before it has to travel back to slow HBM.

Compute-Bound vs Memory-Bound

An operation's 'arithmetic intensity' — FLOPs performed per byte of memory moved — determines whether it is compute-bound or memory-bound. Large matrix multiplications have high intensity (lots of compute per byte) and are compute-bound: they actually use the tensor cores. Elementwise operations (add, activation, normalization) have low intensity and are memory-bound: they spend their time shuttling data, leaving the tensor cores idle.

textArithmetic intensity and the roofline
Arithmetic intensity = FLOPs / bytes moved

Matmul (M×K · K×N):  high intensity  → compute-bound (good)
Elementwise (add, GELU): ~1 FLOP/element → memory-bound (wasteful)

Roofline: performance = min(peak compute, intensity × bandwidth)
The Memory Wall Drives Every Optimization
Almost every technique in this chapter — kernel fusion, FlashAttention, recomputation — is fundamentally about the memory wall: reducing how much data crosses the slow HBM boundary. Once you see the hierarchy, the optimizations become obvious: keep data in SRAM, fuse operations so intermediates never hit HBM, and recompute cheap things rather than storing and reloading them.
This reframes 'efficiency' away from 'fewer FLOPs' and toward 'less memory traffic.' On modern GPUs, the latter usually matters more.
20.3

Chapter 11 introduced gradient checkpointing; here we examine it as a core efficiency technique with quantitative trade-offs. The backward pass needs the forward activations to compute gradients, and storing them all costs memory proportional to depth × batch × sequence. Activation recomputation discards most of these and recomputes them on demand during the backward pass.

The Trade-off, Quantified

StrategyActivation memoryExtra compute
Store all activationsO(L) — fullNone (baseline)
Checkpoint every layerO(√L)~33% (one recompute per segment)
Selective recomputeTunableRecompute only cheap ops
Full recomputeO(1) minimal~2× forward (recompute everything)

The standard choice — checkpointing at layer boundaries — stores only the input to each Transformer block and recomputes the block's internals during backprop. This typically costs about 33% extra compute for a large reduction in activation memory, enabling bigger models or longer sequences on the same GPU.

Selective Recomputation

A refinement (Korthikanti et al., 2022) recognizes that not all activations are equally costly to recompute. The attention softmax is cheap to recompute but expensive to store; the big matmuls are the opposite. Selective recomputation stores the expensive-to-recompute activations and recomputes only the cheap-to-recompute ones, getting most of the memory savings for far less than 33% overhead.

PythonActivation recomputation in PyTorch
import torch; from torch.utils.checkpoint import checkpoint
 
class Block(torch.nn.Module):
    def forward(self, x):
        x = x + self.attn(self.norm1(x))
        x = x + self.ffn(self.norm2(x))
        return x
 
class Model(torch.nn.Module):
    def forward(self, x, recompute=True):
        for block in self.blocks:
            if recompute:
                # Don't save block internals; recompute them in backward
                x = checkpoint(block, x, use_reentrant=False)
            else:
                x = block(x)
        return x
 
# Rule of thumb: enable recomputation when activation memory > weight memory,
# which is common for long sequences and large batches. The ~33% compute
# cost is worth it when it lets you fit a model that otherwise wouldn't.
Recomputation Interacts with Everything
Activation recomputation is rarely used alone. It combines with FSDP (Chapter 18) to fit large models, with FlashAttention (which already recomputes attention internally), and with the parallelism strategies. The decision of WHAT to recompute is increasingly automated by frameworks and compilers that estimate the memory-vs-compute trade-off per operation.
The key intuition: recomputation converts a memory problem into a compute problem. Since large training is often memory-constrained but has spare compute (during memory-bound phases), this trade is frequently favorable.
20.4

Kernel fusion is the single most important idea in GPU efficiency. A 'kernel' is one GPU function launch. Running a chain of separate kernels — say, a matmul, then a bias add, then a GELU — forces each intermediate result to be written to slow HBM and read back by the next kernel. Fusion combines them into ONE kernel that keeps intermediates in fast registers/SRAM, eliminating the round trips.

Why Fusion Helps

Consider computing GELU(x @ W + b). Unfused, this is three kernels: the matmul writes its output to HBM, the bias-add reads and rewrites it, and GELU reads and rewrites it again — the data crosses the slow HBM boundary multiple times. A fused kernel computes all three in one pass, with the intermediate values living only in fast on-chip memory. For memory-bound elementwise chains, fusion can give large speedups.

textThe cost of unfused vs fused
Unfused (3 kernels):  matmul → HBM → bias → HBM → GELU → HBM
    HBM round trips: 3× (each op reads and writes the full tensor)

Fused (1 kernel):     matmul + bias + GELU, intermediates in SRAM
    HBM round trips: 1× (read inputs once, write output once)

Common Fusion Patterns

Fused operationWhat it combines
Fused linear + activationMatmul + bias + GELU/SwiGLU in one pass
Fused LayerNorm/RMSNormMean, variance, normalize, scale in one kernel
Fused attention (Flash)QKᵀ, scale, mask, softmax, ·V — all fused
Fused optimizer stepAdamW update across all params in one kernel
Fused dropout + residualDropout mask + add in one pass
Frameworks Fuse Automatically — Up to a Point
PyTorch ships hand-written fused kernels for common patterns (fused AdamW, fused LayerNorm), and torch.compile (Section 20.8) can fuse many elementwise chains automatically. But the most impactful fusions — like FlashAttention — are hand-crafted because they require algorithmic restructuring (the online softmax) that a general compiler cannot discover.
This is why understanding fusion matters even if you never write a kernel: it tells you which operations are cheap (already fused) and which patterns to avoid (long chains of un-fusable custom ops that each round-trip to HBM).
20.5

Historically, writing fused GPU kernels meant CUDA C++ — difficult, low-level, and slow to iterate. Triton (Tillet et al., 2019), an open-source language from OpenAI, lets you write high-performance fused kernels in Python-like syntax. It handles the hardest parts — memory coalescing, SRAM management, and parallelization — automatically, while you express the computation at the level of blocks of data.

The Triton Programming Model

In Triton, you write a kernel that operates on BLOCKS of data rather than individual elements. The kernel is launched over a grid of program instances, each handling one block. You explicitly load blocks from HBM into SRAM, compute on them, and write results back — but Triton manages the low-level details. This block-level model maps naturally onto the GPU's architecture.

PythonA fused operation in Triton (vector add)
import triton; import triton.language as tl; import torch
 
@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    """Each program instance handles one BLOCK-sized chunk.""" 
    pid = tl.program_id(0)              # which block am I?
    offsets = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offsets < n                     # guard the last partial block
    # Load from HBM into SRAM, compute, store back -- one pass
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    tl.store(out_ptr + offsets, x + y, mask=mask)
 
def add(x, y):
    out = torch.empty_like(x); n = x.numel()
    grid = lambda meta: (triton.cdiv(n, meta['BLOCK']),)
    add_kernel[grid](x, y, out, n, BLOCK=1024)
    return out
 
# This trivial example shows the model. The power comes from FUSING:
# load once, do matmul + bias + activation + dropout in SRAM, store once.

A Fused Kernel That Matters

The real payoff is fusing a chain that would otherwise round-trip to HBM. A fused 'linear + bias + activation' kernel loads the inputs once, performs the matmul, adds the bias, applies the activation — all while the data stays in SRAM — and writes only the final result. The same pattern, scaled up with the online-softmax trick, is exactly how FlashAttention is implemented in Triton.

PythonFused bias + activation in Triton (sketch)
@triton.jit
def fused_bias_gelu(x_ptr, b_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0); offs = pid*BLOCK + tl.arange(0, BLOCK)
    mask = offs < n
    x = tl.load(x_ptr + offs, mask=mask)
    b = tl.load(b_ptr + offs, mask=mask)
    # bias add + GELU, fused -- intermediate never leaves SRAM
    z = x + b
    gelu = 0.5 * z * (1.0 + tl.math.tanh(0.7978845608 * (z + 0.044715 * z*z*z)))
    tl.store(out_ptr + offs, gelu, mask=mask)
 
# Unfused: 2 kernels, 2 extra HBM round trips.
# Fused: 1 kernel, the (x+b) intermediate stays in registers.
# For memory-bound ops, this is a direct ~2x speedup.
Triton Democratized Custom Kernels
Before Triton, writing a competitive fused kernel required deep CUDA expertise. Triton lets ML researchers — not just systems engineers — write kernels that approach hand-tuned CUDA performance. FlashAttention, many fused LayerNorm and attention variants, and much of the custom-kernel ecosystem are written in Triton.
torch.compile itself uses Triton as a backend: it generates Triton kernels for the fused operations it discovers. So even if you never write Triton by hand, the kernels running under your model often are Triton.
20.6

Chapter 19 introduced FlashAttention as an architecture-level efficiency win. Here we revisit it as the canonical example of the efficiency principles of this chapter: it is a fused kernel that exploits the memory hierarchy through tiling and an online softmax, avoiding the slow-memory round trips that make naive attention inefficient.

The Online Softmax Trick

The obstacle to fusing attention is the softmax: it seems to need the entire row of scores at once (to compute the max and the sum for normalization). FlashAttention's key algorithmic insight is the ONLINE softmax — a way to compute the softmax incrementally, block by block, by maintaining a running maximum and running sum and rescaling as new blocks arrive. This lets attention be tiled and fused without ever materializing the full score row.

textFlashAttention (tiled, online softmax)
# Process K,V in blocks; maintain running output and softmax stats
for each query block Q_i:
    m ← -∞   (running max)
    l ← 0    (running sum of exp)
    O ← 0    (running output)
    for each key/value block (K_j, V_j):
        S = Q_i K_jᵀ / √d         # block scores, in SRAM
        m_new = max(m, rowmax(S))
        P = exp(S - m_new)        # rescaled exponentials
        l = l·exp(m - m_new) + rowsum(P)
        O = O·exp(m - m_new) + P V_j   # rescale and accumulate
        m = m_new
    O ← O / l                     # final normalization
    # full T×T score matrix NEVER stored in HBM

Each block of scores lives only in SRAM, is consumed immediately, and is discarded. The running statistics (m, l, O) are tiny. The result is bit-for-bit the same as standard attention, but the HBM traffic drops from O(T²) to O(T). This is the union of every idea in this chapter: fusion, tiling, the memory hierarchy, and a clever algorithm to make it all possible.

FlashAttention-2 and -3
Successive versions refined the original: FlashAttention-2 (Dao, 2023) improved the work partitioning across GPU threads for better occupancy, and FlashAttention-3 (2024) exploited the asynchrony and fp8 capabilities of Hopper GPUs. Each version pushed attention closer to the hardware's peak throughput.
The progression illustrates a general truth of GPU efficiency: the algorithm and the hardware co-evolve. A kernel must be re-tuned — sometimes substantially rewritten — for each new GPU generation to extract peak performance.
20.7

Chapter 15 covered bf16 mixed precision. The frontier of low-precision training is fp8 — eight-bit floating point — supported natively on Hopper (H100) and newer GPUs. fp8 roughly doubles throughput and halves memory versus bf16, but its tiny dynamic range demands careful scaling to avoid overflow and underflow.

The Precision Ladder

FormatBitsUse in training
fp3232Master weights, optimizer state, loss accumulation
bf1616Default compute precision (Chapter 15)
fp1616Older default; needs loss scaling
fp8 (E4M3)8Forward-pass matmuls (more precision)
fp8 (E5M2)8Gradient matmuls (more range)
int8 / int48 / 4Inference quantization (Chapter 27)

fp8 has two variants: E4M3 (4 exponent bits, 3 mantissa) offers more precision for forward activations, while E5M2 (5 exponent, 2 mantissa) offers more range for gradients. Training in fp8 keeps a higher-precision master copy and uses per-tensor scaling factors to map values into fp8's narrow range — the same mixed-precision philosophy as bf16, pushed further.

⚠️
fp8 Requires Careful Scaling
fp8's dynamic range is so narrow that naive use overflows or underflows immediately. Successful fp8 training (as in NVIDIA's Transformer Engine) uses per-tensor scaling factors, updated dynamically, to keep each tensor's values within fp8's representable range — a more elaborate version of fp16's loss scaling.
Not every operation tolerates fp8: normalization, softmax, and accumulation typically stay in higher precision. fp8 is applied selectively to the large matmuls where it gives the most throughput and the precision loss is least harmful. Getting this right is delicate, which is why fp8 training is still maturing even as the hardware supports it.
20.8

Writing kernels by hand is powerful but laborious. Graph compilers automate much of it: they capture the model's computation as a graph, then apply optimizations — fusing operations, eliminating redundant work, generating efficient kernels — automatically. PyTorch's torch.compile is the most widely used; it can deliver substantial speedups with a single line of code.

How torch.compile Works

The torch.compile pipeline

1CaptureTorchDynamo traces the Python code into an FX graph
2LowerAOTAutograd builds forward AND backward graphs
3OptimizeThe graph is fused, simplified, and pattern-matched
4CodegenTorchInductor generates fused Triton/C++ kernels
5ExecuteCompiled kernels run, cached for reuse
PythonUsing torch.compile
import torch
 
model = GPT(...).cuda()
 
# One line: compile the model. First call traces & compiles; then it's fast.
model = torch.compile(model)
 
# Compilation modes trade compile time for runtime speed:
model = torch.compile(model, mode='max-autotune')  # slow compile, fastest run
 
# The training loop is UNCHANGED -- compilation is transparent
for batch in loader:
    with torch.autocast('cuda', dtype=torch.bfloat16):
        loss = lm_loss(model, batch.cuda())
    loss.backward(); opt.step(); opt.zero_grad()
 
# Typical speedups: 1.3-2x on training, more on inference. The compiler
# fuses elementwise chains, the norm, and activation into Triton kernels,
# reducing the HBM round trips described in Section 20.2.
⚠️
Graph Breaks Kill Compiler Gains
torch.compile works by capturing a contiguous graph of operations. Data-dependent Python control flow (an if that branches on a tensor's value, printing tensors, calling un-traceable libraries) causes a 'graph break' — the compiler splits the graph, losing fusion opportunities across the break. A model riddled with graph breaks sees little speedup.
Writing 'compiler-friendly' code — avoiding data-dependent control flow in the hot path, keeping operations on tensors — lets the compiler capture large fusable graphs. The torch._dynamo.explain tool reports where graph breaks occur so you can eliminate them.
20.9

The techniques in this chapter are complementary, and a well-tuned training run uses most of them together. Here is how they stack into a coherent, efficient configuration — the practical synthesis of the chapter.

TechniqueWhat it givesCost
bf16 mixed precision2× throughput, half memoryNegligible (Ch. 15)
FlashAttentionFast, memory-linear attentionNone (exact)
torch.compileAutomatic fusion, 1.3–2×Compile time
Activation recomputeFit bigger models/batches~33% compute
Fused optimizerFaster optimizer stepNone
fp8 (where supported)Further 2× on matmulsScaling complexity
FSDP / ZeROFit large models (Ch. 18)Communication
PythonCode Lab: an efficient training setup
import torch; from torch.distributed.fsdp import FullyShardedDataParallel as FSDP
 
model = GPT(...).cuda()
model.gradient_checkpointing_enable()        # activation recompute (20.3)
model = FSDP(model, auto_wrap_policy=block_policy)  # shard params (Ch.18)
model = torch.compile(model, mode='max-autotune')  # fuse + codegen (20.8)
 
# Fused AdamW kernel (one launch updates all params)
opt = torch.optim.AdamW(model.parameters(), lr=3e-4, fused=True)
 
for batch in loader:
    with torch.autocast('cuda', dtype=torch.bfloat16):  # bf16 (Ch.15)
        loss = lm_loss(model, batch.cuda())      # uses FlashAttention internally
    loss.backward()
    torch.nn.utils.clip_grad_norm_(model.parameters(), 1.0)
    opt.step(); opt.zero_grad()
 
# Stacked, these techniques can take MFU from ~25% (naive) to ~50%+,
# halving the time and cost of the same training run.
Measure, Don't Guess
Efficiency work must be guided by profiling, not intuition. Tools like the PyTorch profiler, NVIDIA Nsight, and the simple MFU calculation of Section 20.1 reveal where time actually goes. Often the bottleneck is surprising — a small un-fused operation, a data-loading stall, or a communication imbalance — and fixing the real bottleneck beats optimizing what you assumed was slow.
The discipline: profile, find the actual bottleneck, apply the relevant technique, measure again. Efficiency is empirical, and the MFU number is the scoreboard.
20.10

This chapter has focused on efficient PRETRAINING. A related family of techniques makes FINE-TUNING efficient, and they share the same spirit — doing more with less compute and memory. They become central in Part V (alignment), but it is worth previewing them here as efficiency techniques.

LoRA: Low-Rank Adaptation

Full fine-tuning updates all of a model's billions of parameters — expensive in compute and memory, and producing a full-size copy per task. LoRA (Hu et al., 2021) freezes the pretrained weights and learns only small low-rank update matrices added alongside them. A weight update ΔW is factored as the product of two thin matrices BA, so instead of training d×d parameters you train only 2×d×r for a small rank r.

textLoRA: low-rank weight update
Frozen weight W (d×d).  Learn ΔW = B A  where B is d×r, A is r×d, r ≪ d.

Forward: y = (W + BA) x = Wx + B(Ax)

Trainable params: 2dr instead of d².  For d=4096, r=8: 65k vs 16.8M (256× fewer).
LoRA and QLoRA Democratized Fine-Tuning
LoRA reduces the trainable parameters by orders of magnitude, so fine-tuning a large model fits on a single consumer GPU. QLoRA (Dettmers et al., 2023) goes further: it quantizes the frozen base model to 4-bit while training LoRA adapters in higher precision, enabling fine-tuning of 65B models on a single 48GB GPU.
These techniques are why a hobbyist can fine-tune a large open model at home. We return to them in depth in Chapter 22 (supervised fine-tuning), where they are the standard tool for adapting pretrained models to specific tasks and behaviours.
20.11

Efficiency Quick-Reference

TechniqueAttacksMechanism
Kernel fusionHBM round tripsKeep intermediates in SRAM
FlashAttentionAttention memory trafficTiling + online softmax
Activation recomputeActivation memoryRecompute in backward
fp8 / bf16Memory + throughputFewer bits per number
torch.compileUn-fused op chainsAutomatic graph fusion
Fused optimizerOptimizer-step launchesOne kernel for all params
LoRA / QLoRAFine-tuning costLow-rank, quantized adaptation

Exercises

Exercises 1–10 are pen-and-paper or derivations; 11–20 require code.

Exercise 1: Pen & Paper
Define Model FLOPs Utilization. For a 13B model processing 25k tokens/s on an H100 (~990 TFLOP/s bf16), compute the MFU using the 6N rule.
Exercise 2: Pen & Paper
Explain the difference between compute-bound and memory-bound operations using arithmetic intensity. Classify: matmul, GELU, LayerNorm, softmax.
Exercise 3: Pen & Paper
Sketch the GPU memory hierarchy (registers, SRAM, HBM, NVLink, IB) with sizes and bandwidths. Why does keeping data in SRAM matter so much?
Exercise 4: Pen & Paper
Quantify the HBM traffic for unfused vs fused GELU(x@W+b). Count read/write passes for each and explain the speedup for a memory-bound chain.
Exercise 5: Derive
For activation recomputation with checkpoints every √L layers, derive the O(√L) memory and the ~33% extra compute. State the assumptions.
Exercise 6: Pen & Paper
Explain the online-softmax trick. Why does maintaining a running max and sum let attention be computed block-by-block without the full score row?
Exercise 7: Pen & Paper
Compare fp8 E4M3 and E5M2. Why is E4M3 used for forward activations and E5M2 for gradients? What does per-tensor scaling accomplish?
Exercise 8: Pen & Paper
Explain what a graph break is in torch.compile and why it reduces speedup. Give two code patterns that cause graph breaks and how to avoid them.
Exercise 9: Derive
For LoRA with rank r on a d×d weight, derive the trainable-parameter count 2dr. For d=4096, r=16, compute the reduction factor versus full fine-tuning.
Exercise 10: Pen & Paper
List five efficiency techniques and, for each, state whether it primarily saves memory, compute, or memory traffic. Explain how they compose.
Exercise 11: Code
Compute MFU for a real training step: time it, count tokens, apply the 6N rule, and divide by your GPU's peak. Report the MFU and identify the likely bottleneck.
Exercise 12: Code
Profile a training step with the PyTorch profiler. Identify the top three time-consuming operations and classify each as compute- or memory-bound.
Exercise 13: Code Lab
Measure activation recomputation: train a deep model with and without gradient_checkpointing. Report peak memory, throughput, and confirm the ~33% compute overhead.
Exercise 14: Code
Demonstrate fusion's benefit: implement GELU(x@W+b) as separate ops vs a single fused expression, and measure the wall-clock and memory difference on a large tensor.
Exercise 15: Code Lab
Write a Triton vector-add kernel and verify it matches torch. Then extend it to a fused bias+GELU kernel and benchmark against the unfused PyTorch version.
Exercise 16: Code
Apply torch.compile to a small model. Measure the speedup vs eager mode for training and inference. Use torch._dynamo.explain to find any graph breaks.
Exercise 17: Code
Compare your from-scratch attention against torch's scaled_dot_product_attention (FlashAttention) at T = 1k, 4k, 8k. Plot speed and peak memory vs T.
Exercise 18: Code
Enable the fused AdamW kernel (fused=True) and measure the optimizer-step time vs the default. Report the speedup as a function of parameter count.
Exercise 19: Code
Implement a minimal LoRA layer: freeze a linear layer's weight and add a trainable BA update. Verify it trains and that only 2dr parameters receive gradients.
Exercise 20: Code (Challenge)
Build an efficiency ablation: train the same small model under (a) naive eager bf16, (b) + FlashAttention, (c) + torch.compile, (d) + activation recompute. Report MFU, throughput, and peak memory for each configuration, and write up which technique gave the biggest win and why.

Further reading: “FlashAttention” and “FlashAttention-2/3” (Dao et al., 2022–2024). “Reducing Activation Recomputation in Large Transformer Models” (Korthikanti et al., 2022) for selective recomputation. “Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations” (Tillet et al., 2019) and the Triton tutorials. The torch.compile and TorchInductor documentation. “LoRA” (Hu et al., 2021) and “QLoRA” (Dettmers et al., 2023). NVIDIA's Transformer Engine documentation for fp8 training.

Next → Chapter 21: Evaluation During Pretraining

You now have an efficient, distributed training run consuming clean, curated data — but how do you know if it is working? Chapter 21, the close of Part IV, covers evaluation during pretraining: tracking loss and perplexity, the benchmark suites (MMLU, HellaSwag, and others) used to measure capability, the perils of benchmark contamination, and how to read the signals of a months-long run to decide whether to continue, adjust, or stop. We close Part IV by turning the training process from a black box into something you can measure and steer.

20 Exercises in this chapter
Attempt each exercise before checking the worked solutions.
View Solutions →