Efficient Training Techniques
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. |
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.
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.
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
| Level | Size (H100) | Bandwidth | Role |
|---|---|---|---|
| Registers | ~256 KB/SM | Fastest | Per-thread scratch |
| SRAM (shared) | ~228 KB/SM | ~20 TB/s | On-chip tiles |
| HBM (global) | 80 GB | ~3.3 TB/s | Model, activations |
| Inter-GPU (NVLink) | — | ~0.9 TB/s | Cross-device |
| Inter-node (IB) | — | ~0.1–0.4 TB/s | Cross-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.
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)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
| Strategy | Activation memory | Extra compute |
|---|---|---|
| Store all activations | O(L) — full | None (baseline) |
| Checkpoint every layer | O(√L) | ~33% (one recompute per segment) |
| Selective recompute | Tunable | Recompute only cheap ops |
| Full recompute | O(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.
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.
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.
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 operation | What it combines |
|---|---|
| Fused linear + activation | Matmul + bias + GELU/SwiGLU in one pass |
| Fused LayerNorm/RMSNorm | Mean, variance, normalize, scale in one kernel |
| Fused attention (Flash) | QKᵀ, scale, mask, softmax, ·V — all fused |
| Fused optimizer step | AdamW update across all params in one kernel |
| Fused dropout + residual | Dropout mask + add in one pass |
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.
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.
@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.
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.
# 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.
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
| Format | Bits | Use in training |
|---|---|---|
| fp32 | 32 | Master weights, optimizer state, loss accumulation |
| bf16 | 16 | Default compute precision (Chapter 15) |
| fp16 | 16 | Older default; needs loss scaling |
| fp8 (E4M3) | 8 | Forward-pass matmuls (more precision) |
| fp8 (E5M2) | 8 | Gradient matmuls (more range) |
| int8 / int4 | 8 / 4 | Inference 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.
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
| 1 | Capture | TorchDynamo traces the Python code into an FX graph |
| 2 | Lower | AOTAutograd builds forward AND backward graphs |
| 3 | Optimize | The graph is fused, simplified, and pattern-matched |
| 4 | Codegen | TorchInductor generates fused Triton/C++ kernels |
| 5 | Execute | Compiled kernels run, cached for reuse |
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.
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.
| Technique | What it gives | Cost |
|---|---|---|
| bf16 mixed precision | 2× throughput, half memory | Negligible (Ch. 15) |
| FlashAttention | Fast, memory-linear attention | None (exact) |
| torch.compile | Automatic fusion, 1.3–2× | Compile time |
| Activation recompute | Fit bigger models/batches | ~33% compute |
| Fused optimizer | Faster optimizer step | None |
| fp8 (where supported) | Further 2× on matmuls | Scaling complexity |
| FSDP / ZeRO | Fit large models (Ch. 18) | Communication |
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.
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.
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).Efficiency Quick-Reference
| Technique | Attacks | Mechanism |
|---|---|---|
| Kernel fusion | HBM round trips | Keep intermediates in SRAM |
| FlashAttention | Attention memory traffic | Tiling + online softmax |
| Activation recompute | Activation memory | Recompute in backward |
| fp8 / bf16 | Memory + throughput | Fewer bits per number |
| torch.compile | Un-fused op chains | Automatic graph fusion |
| Fused optimizer | Optimizer-step launches | One kernel for all params |
| LoRA / QLoRA | Fine-tuning cost | Low-rank, quantized adaptation |
Exercises
Exercises 1–10 are pen-and-paper or derivations; 11–20 require code.
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.