Chapter 6: Large-Scale Model Training

Advanced JAX: Performance, Optimization, and Scale · Chapter 6 study notes

Once models get big, the hard question in training is no longer "how do I compute the gradient". It becomes does it fit, can it finish, and can communication between devices keep up. This chapter first looks at the three kinds of bottleneck in large-model training (memory, compute, communication), then covers the tools for each one: using Flax NNX / Optax to organize parameters and optimizer state, and scaling training up with data parallelism, gradient accumulation, gradient checkpointing (rematerialization), mixed precision, and tensor parallelism / FSDP / pipeline parallelism. It ends by comparing optimizers suited to large-scale training.

Learning goals

  • Estimate how much memory parameters, gradients, optimizer state and activations each take during training, and tell where the bottleneck is
  • Build a training loop with Flax NNX (nnx.Module, nnx.Rngs, nnx.Param / nnx.BatchStat, nnx.Optimizer, nnx.jit, nnx.split/merge) and Optax (chain, clip_by_global_norm, adamw, learning-rate schedules, MultiSteps), and be able to read older Linen / Haiku code
  • Migrate pmap code to jit + NamedSharding, and annotate parameters with sharding using nnx.with_partitioning
  • Implement gradient accumulation with lax.scan, and explain the differences that come from mean vs. sum and from BatchNorm
  • Trade memory for compute with jax.checkpoint and its policies, and verify the effect quantitatively with memory_analysis() and print_saved_residuals
  • Understand how bf16 differs from fp16, implement "fp32 parameters + bf16 compute", and implement dynamic loss scaling for fp16
  • Implement Megatron-style tensor parallelism and FSDP on a 2D data × model mesh, and verify they match single-device training
  • Compare the optimizer-state memory and use cases of AdamW, Adafactor, LAMB, Lion and LARS

Sections

SectionContent
6.1Challenges in training large models
6.2Introduction to JAX ecosystem libraries (Flax, Haiku)
6.3Managing model parameters and state
6.4Combining pmap with training frameworks
6.5Gradient accumulation
6.6Gradient checkpointing (rematerialization)
6.7Mixed precision training
6.8Model parallelism strategies
6.9Optimization algorithms at scale
6.10Practice: implementing gradient checkpointing
—Hands-on exercises · Quiz · Version notes · Summary · Cheat sheet · Where to go next

Original course link: https://apxml.com/courses/advanced-jax/chapter-6-large-scale-model-training-jax

How to read this: The material follows the course's main points, and every piece of code has actually been run. Outdated patterns go in ❌ cells, each followed right away by a ✅ recommended version.

Environment: Python 3.14, JAX / jaxlib 0.12.0.1, Flax 0.12.12 (including NNX), Optax 0.2.9, Orbax-checkpoint 0.12.10, CPU only. XLA_FLAGS=--xla_force_host_platform_device_count=8 simulates 8 devices on one CPU so we can demonstrate parallelism. All conclusions here are based on what this environment actually produced. Timings only reflect this CPU and are not asserted.

Links to the Basics course (this chapter doesn't repeat this material)

Basics notesFoundations used in this chapter
Basics 02 JIT compilationTracing and compilation, lower() / compile(), static arguments, recompilation
Basics 03 Autodiff with gradgrad / value_and_grad / has_aux, backprop needs to save residuals
Basics 04 Vectorization with vmapvmap and the batch dimension
Basics 05 Multi-device parallelism with pmappmap, axis_name, lax.pmean / psum, simulating multiple devices
Basics 06 State managementPyTree state, the lax.scan carry, hand-written SGD / Adam optimizers, donate_argnums
Basics 07 Practical topicsTrainState, a complete training loop

Building on that, this chapter switches to ecosystem libraries (NNX + Optax) and focuses on the new material that comes with "scale".

# Must be set before importing jax: simulate 8 devices on the CPU for the data-parallel / model-parallel / FSDP demos
import os
os.environ["XLA_FLAGS"] = "--xla_force_host_platform_device_count=8"

import sys, time, re, warnings, tempfile
from functools import partial

import numpy as np
import jax
import jax.numpy as jnp
from jax import lax
from jax.sharding import NamedSharding, PartitionSpec as P, AxisType, reshard
import optax
import flax
from flax import nnx

print("Python", sys.version.split()[0])
print("JAX   ", jax.__version__, "| Flax", flax.__version__, "| Optax", optax.__version__)
print("Devices", jax.devices())
assert jax.device_count() == 8

# ---------- Small helpers used throughout this chapter ----------
def tree_bytes(tree):
    """Total bytes of all array leaves in a PyTree."""
    return sum(x.nbytes for x in jax.tree.leaves(tree) if hasattr(x, "nbytes"))

def max_diff(a, b):
    """Largest elementwise absolute difference between two PyTrees of the same structure (compared in float64)."""
    la, lb = jax.tree.leaves(a), jax.tree.leaves(b)
    assert len(la) == len(lb)
    return max(float(np.max(np.abs(np.asarray(x, np.float64) - np.asarray(y, np.float64)))) for x, y in zip(la, lb))

def bench(f, *args, n=10):
    """Warm up first (including compilation), then time n runs and return the average in milliseconds."""
    jax.block_until_ready(f(*args))
    t0 = time.perf_counter()
    for _ in range(n):
        out = f(*args)
    jax.block_until_ready(out)
    return (time.perf_counter() - t0) / n * 1e3

def mem_of(f, *args):
    """Compile jit(f) and return XLA's memory analysis (without running it)."""
    return jax.jit(f).lower(*args).compile().memory_analysis()

def collectives(compiled_text):
    """Count how often each kind of collective appears in the compiled HLO."""
    ops = ["all-reduce", "all-gather", "reduce-scatter", "all-to-all", "collective-permute"]
    return {op: len(re.findall(op + r"(?:-start)?\(", compiled_text)) for op in ops}

MiB = 2 ** 20
Python 3.14.7
JAX    0.12.0.1 | Flax 0.12.12 | Optax 0.2.9
Devices [CpuDevice(id=0), CpuDevice(id=1), CpuDevice(id=2), CpuDevice(id=3), CpuDevice(id=4), CpuDevice(id=5), CpuDevice(id=6), CpuDevice(id=7)]

6.1 Challenges in training large models

Model sizes (especially in NLP and CV) keep growing. JAX gives you tools for acceleration and automatic differentiation, but past a certain size a model runs into the physical limits of the hardware. The course groups the challenges into three kinds: memory, compute and communication.

Memory limits: the main bottleneck

A single accelerator (one GPU or one TPU core) usually has only 16 GB to 80 GB of high-bandwidth memory (HBM). (Those are typical numbers for earlier generations. Newer accelerators such as H200, B200 and MI300X have reached 141–192 GB, but that's still not enough for models with tens of billions of parameters, so the conclusions below don't change.) Five things take up most of the memory during training:

  1. Model parameters. Models with tens of billions of parameters are common now. Storing 10 billion parameters in FP32 takes, for the parameters alone, $$10\times10^{9}\ \text{params}\times 4\ \text{bytes/param} = 40\ \text{GB}.$$
  2. Optimizer state. Adam / AdamW keeps a first moment $m$ and a second moment $v$ for every parameter, which needs an extra $$10\times10^{9}\times 2\times 4\ \text{bytes} = 80\ \text{GB},$$ so the "static memory" of parameters plus optimizer state is already 120 GB, more than almost any single accelerator holds.
  3. Activations. Intermediate results from the forward pass have to be kept for the backward pass. Their size grows with batch size, sequence length, model depth and width. For deep networks and long sequences they can be far larger than the parameter memory.
  4. Gradients. Same shape as the parameters, so another 40 GB in FP32.
  5. Workspace memory. Temporary buffers that the XLA compiler, cuDNN and other libraries need while running operations.

If you go over the device memory you get an OOM and training stops. The following sections each tackle one part: gradient checkpointing (6.6) mainly cuts activation memory, mixed precision (6.7) reduces the bytes used by activations and parameters, and FSDP (6.8) spreads parameters, gradients and optimizer state across several devices.

Compute cost and training time

  • FLOP growth: compute often grows nonlinearly with model dimensions. For example, Transformer self-attention compute grows quadratically with sequence length, and doubling depth or width can much more than double the compute.
  • Training time: it can take days, weeks or even months, so hardware cost, energy use and iteration cycles are all high.
  • Hardware requirements: you need high-performance GPU clusters or large TPU Pods.
  • Reducing total FLOPs usually means changing the algorithm or architecture. Mixed precision can speed things up by using faster low-precision compute units.

Communication overhead in distributed training

  • Data parallelism: every device holds a full copy of the model and processes a different shard of the data. After the local backward pass, an All-Reduce collective synchronizes the gradients, and then the optimizer update runs.
  • Communication cost depends on model size and on the bandwidth and latency of the interconnect between devices. For very large models, communication can eat a big chunk of each step and hurt scaling efficiency. A slow interconnect (Ethernet, compared with NVLink / TPU ICI) makes this worse.
  • Model parallelism / pipeline parallelism bring more complex communication patterns (6.8).
  • Ways to reduce the impact of communication: better collective algorithms, overlapping communication with compute, and choosing a parallel strategy that better balances compute and communication.

References from the course: Attention Is All You Need (Vaswani et al., 2017); Mixed Precision Training (Micikevicius et al., arXiv:1710.03740); Megatron-LM (Shoeybi et al., arXiv:1909.08053).

Next we turn the memory formulas above into code, then use Optax to check that "Adam state = 2× parameters".

GB = 1e9
n_params = 10e9                      # example: 10 billion parameters

mem = {
    "params (fp32)":      n_params * 4,
    "grads (fp32)":       n_params * 4,
    "Adam m+v (fp32)":    n_params * 2 * 4,
}
for k, v in mem.items():
    print(f"{k:18s} {v / GB:6.0f} GB")
static = mem["params (fp32)"] + mem["Adam m+v (fp32)"]
print(f"{'static (params+optimizer)':14s} {static / GB:6.0f} GB")
print(f"{'total (no activations)':15s} {sum(mem.values()) / GB:6.0f} GB")

assert mem["params (fp32)"] / GB == 40 and mem["Adam m+v (fp32)"] / GB == 80
assert static / GB == 120 and sum(mem.values()) / GB == 160

# Measure with Optax: Adam state is exactly 2x the parameter bytes (plus one int32 step count)
toy = {"w": jnp.zeros((1024, 512)), "b": jnp.zeros(512)}
adam_state = optax.adam(1e-3).init(toy)
ratio = tree_bytes(adam_state) / tree_bytes(toy)
print(f"\nOptax Adam state / params byte ratio = {ratio:.6f}")
assert abs(ratio - 2.0) < 1e-4
params (fp32)          40 GB
grads (fp32)           40 GB
Adam m+v (fp32)        80 GB
static (params+optimizer)    120 GB
total (no activations)    160 GB

Optax Adam state / params byte ratio = 2.000002

⚠️ Note: You'll often hear that Adam/AdamW optimizer state "needs 2–3× the parameter storage". In Optax 0.2.9, Adam/AdamW state is exactly 2× the parameters: one copy each of $m$ and $v$, plus a scalar step count. It only gets to 3× if you also count the fp32 master copy of the weights used in mixed precision training as optimizer state (see the Extra below).

🔍 Extra: how many bytes does each parameter really need? And how big are activations?

🟢 Must-read

1) The "16 bytes/param" rule of thumb for mixed precision + Adam (the accounting from the ZeRO paper): bf16 params 2 B + bf16 grads 2 B + fp32 master weights 4 B + fp32 $m$ 4 B + fp32 $v$ 4 B = 16 B/param. For 10 billion parameters that's 160 GB. So mixed precision does not reduce static memory. What it mainly saves is activations and bandwidth. To bring static memory down you need the FSDP/ZeRO sharding from section 6.8: split across $N$ devices, each device holds about $16\Phi/N$ bytes.

2) Estimating training compute: a dense Transformer needs about $6\Phi$ FLOPs per training token (forward $2\Phi$, backward $4\Phi$, where $\Phi$ is the parameter count). So total compute is $C\approx 6\Phi T$ ($T$ is the number of training tokens). The "backward ≈ 2× forward" ratio comes up again in section 6.6 when we discuss the cost of checkpointing.

3) Communication volume in data parallelism: with a ring All-Reduce, each device sends and receives about $2\frac{N-1}{N}S$ bytes ($S$ is the gradient size in bytes), which is almost independent of the number of devices. That's why data parallelism scales, but it also means every step moves about 2× the model size in data.

def bytes_per_param(mixed_precision=True):
    """Bytes needed per parameter for Adam training (excluding activations)."""
    if mixed_precision:
        return dict(bf16_params=2, bf16_grads=2, fp32_master=4, adam_m=4, adam_v=4)
    return dict(fp32_params=4, fp32_grads=4, adam_m=4, adam_v=4)

mp = bytes_per_param(True)
print("Mixed-precision Adam:", mp, "total", sum(mp.values()), "B/param →", n_params * sum(mp.values()) / GB, "GB")
assert sum(mp.values()) == 16
for n_dev in [8, 64, 512]:
    print(f"  FSDP/ZeRO-3 across {n_dev:3d} devices: {n_params * 16 / n_dev / GB:7.2f} GB per device")

# Training compute: C ≈ 6 · Φ · T
tokens = 1e12
C_total = 6 * n_params * tokens
chip_flops, mfu, n_chips = 300e12, 0.4, 256        # assumptions: 300 TFLOP/s per chip, 40% utilization, 256 chips
days = C_total / (chip_flops * mfu * n_chips) / 86400
print(f"\n10B params × 1T tokens ≈ {C_total:.1e} FLOPs, about {days:.1f} days on 256 chips")

# Per-device traffic of a ring all-reduce
S = n_params * 4
for N in [8, 64]:
    print(f"N={N:3d}: each device sends/receives ≈ {2 * (N - 1) / N * S / GB:.1f} GB of gradient data per step")
Mixed-precision Adam: {'bf16_params': 2, 'bf16_grads': 2, 'fp32_master': 4, 'adam_m': 4, 'adam_v': 4} total 16 B/param → 160.0 GB
  FSDP/ZeRO-3 across   8 devices:   20.00 GB per device
  FSDP/ZeRO-3 across  64 devices:    2.50 GB per device
  FSDP/ZeRO-3 across 512 devices:    0.31 GB per device

10B params × 1T tokens ≈ 6.0e+22 FLOPs, about 22.6 days on 256 chips
N=  8: each device sends/receives ≈ 70.0 GB of gradient data per step
N= 64: each device sends/receives ≈ 78.8 GB of gradient data per step

6.2 Introduction to JAX ecosystem libraries (Flax, Haiku)

jit, grad, vmap and pmap are low-level primitives. If you write a large model directly with them, managing state such as parameters, optimizer statistics and random keys gets very tedious, and even more so across multiple devices. Neural network libraries provide higher-level abstractions on top of these primitives. The course introduces the two main libraries in the JAX ecosystem: Flax (Google) and Haiku (DeepMind). They differ in style but share the same goals.

Flax (Linen) as the course presents it: structured modules + explicit state

  1. Module system: subclass nn.Module, and define sublayers in setup or inline in a __call__ decorated with @nn.compact.
  2. Explicit initialization: parameters are created by module.init(rng, example_input), which separates the structure definition from the parameter values.
  3. Functional apply: model.apply({'params': params}, x) runs the forward pass with the given parameters, which makes it easy to combine with JAX transformations.
  4. PyTree state: parameters are nested dicts and can be passed straight to jit / grad.
  5. TrainState: flax.training.train_state.TrainState bundles apply_fn, the parameters and the optimizer state, which simplifies the training loop.

Haiku as the course presents it: object-oriented style + implicit state

  1. Module system: subclass hk.Module, and usually instantiate sublayers in __init__ or __call__.
  2. Implicit parameter creation: the first time a layer is called, Haiku creates and manages its parameters internally.
  3. hk.transform: turns a function that contains modules into a pair of pure functions, init (initialize parameters) and apply (forward pass).
  4. Internal state management: inside an hk.transform / hk.transform_with_state context, it manages RNG sequences and BatchNorm statistics behind the scenes.

Why large models use these libraries

  1. Organization: a fixed structure keeps even complex architectures readable and maintainable;
  2. State management: parameters, optimizer state, batch statistics and RNGs are handled in one place, which matters especially across devices;
  3. Less boilerplate;
  4. Composability: modules can be nested and reused;
  5. Ecosystem integration: matching tools for checkpointing, metric logging and training loops.

The course's conclusion: pick whichever style you prefer. Flax is more explicit and more functional; Haiku is more object-oriented, with implicit state.

⚠️ Note: Many sources describe Flax as "centered on flax.linen" and treat Haiku as an equally current alternative. As of November 2026 that has changed:

  • Flax's recommended API is now NNX. The front page of the official docs says "This site covers the new Flax NNX API". The Linen docs moved to a separate site (flax-linen.readthedocs.io), which says "new users are encouraged to use Flax NNX". Linen is the old API, though the maintainers have also said it won't be deprecated any time soon. As you'll see below, Linen still runs fine in Flax 0.12.12.
  • Haiku is in maintenance mode. The dm-haiku README says: "As of July 2023 Google DeepMind recommends that new projects adopt Flax instead of Haiku". The project only gets bug fixes and compatibility updates for new JAX versions, no new features.

So the main thread of this tutorial is always written in NNX. The course's Linen code is kept for comparison, and the Haiku code is for reading only (Haiku isn't installed in this environment).

First we run the course's Linen example for comparison, then show the NNX version.

# ❌ Old style (outdated): flax.linen is the old API (NNX is recommended for new projects); PRNGKey and jax.tree_util.tree_map are old style too
try:
    import flax.linen as nn

    class SimpleMLP(nn.Module):
        features: list[int]          # list of layer sizes, e.g. [128, 64, 10]

        @nn.compact                  # lets you define submodules inline in __call__
        def __call__(self, x):
            for i, feat in enumerate(self.features):
                x = nn.Dense(features=feat, name=f'dense_{i}')(x)
                if i != len(self.features) - 1:   # ReLU on every layer except the last
                    x = nn.relu(x)
            return x

    key = jax.random.PRNGKey(0)
    dummy_input = jnp.ones((1, 28 * 28))        # one flattened MNIST image (batch=1)
    linen_model = SimpleMLP(features=[128, 10])
    linen_params = linen_model.init(key, dummy_input)['params']
    print("Parameter PyTree structure:", jax.tree_util.tree_map(lambda x: x.shape, linen_params))
    linen_out = linen_model.apply({'params': linen_params}, dummy_input)
    print("Output shape:", linen_out.shape)
    print("Measured: Linen still runs in Flax", flax.__version__, "(old API, not deprecated)")
except Exception as e:
    print("Measured error:", type(e).__name__, e)
Parameter PyTree structure: {'dense_0': {'bias': (128,), 'kernel': (784, 128)}, 'dense_1': {'bias': (10,), 'kernel': (128, 10)}}
Output shape: (1, 10)
Measured: Linen still runs in Flax 0.12.12 (old API, not deprecated)
# ✅ Recommended: Flax NNX
class SimpleMLP(nnx.Module):
    def __init__(self, din: int, features: list[int], *, rngs: nnx.Rngs):
        dims = [din] + list(features)
        # A list of submodules must be wrapped in nnx.List so NNX recognizes it as a graph node
        self.layers = nnx.List([nnx.Linear(a, b, rngs=rngs) for a, b in zip(dims[:-1], dims[1:])])

    def __call__(self, x):
        for i, layer in enumerate(self.layers):
            x = layer(x)
            if i != len(self.layers) - 1:
                x = nnx.relu(x)
        return x

model = SimpleMLP(28 * 28, [128, 10], rngs=nnx.Rngs(0))     # constructing it initializes it
x_demo = jnp.ones((1, 28 * 28))
out = model(x_demo)                                           # call it directly
print("Output shape:", out.shape)

params_state = nnx.state(model, nnx.Param)                    # pull out every nnx.Param
shapes = jax.tree.map(jnp.shape, nnx.to_pure_dict(params_state))
print("Parameter structure:", shapes)

# Parameter shapes match the Linen version above one to one (Linen's dense_i ↔ NNX's layers[i])
assert shapes["layers"][0]["kernel"] == linen_params["dense_0"]["kernel"].shape == (784, 128)
assert shapes["layers"][1]["kernel"] == linen_params["dense_1"]["kernel"].shape == (128, 10)
assert out.shape == linen_out.shape == (1, 10)
Output shape: (1, 10)
Parameter structure: {'layers': {0: {'bias': (128,), 'kernel': (784, 128)}, 1: {'bias': (10,), 'kernel': (128, 10)}}}

Key points of the NNX style: an NNX module is just an ordinary mutable Python object. You create sublayers and parameters directly in __init__ (passing nnx.Rngs explicitly when you need randomness). There's no two-step init / apply. Once the object is constructed, the parameters are already inside it.

NNX's "functional exit": nnx.split / nnx.merge

NNX objects are mutable, but at any time you can split one into a static structure GraphDef + pure data State (a PyTree) and merge them back. This step plays the role of Haiku's init/apply and Linen's apply(params, x): the State can go straight into tools that only understand PyTrees, such as jax.jit, jax.grad and Orbax.

"Lifted" transformations like nnx.jit / nnx.grad do split → call → merge for you internally, and write back any changes the function makes to the module (BatchNorm statistics, optimizer step count, and so on) to the original object.

# ❌ Old style (outdated): Haiku has been in maintenance mode since 2023-07, and DeepMind recommends Flax for new projects; haiku isn't installed here
try:
    import haiku as hk

    class HkMLP(hk.Module):
        def __init__(self, features, name=None):
            super().__init__(name=name)
            self.features = features
        def __call__(self, x):
            for i, feat in enumerate(self.features):
                x = hk.Linear(output_size=feat, name=f'linear_{i}')(x)
                if i != len(self.features) - 1:
                    x = jax.nn.relu(x)
            return x

    model = hk.transform(lambda x: HkMLP([128, 10])(x))
    params = model.init(jax.random.PRNGKey(0), jnp.ones((1, 784)))
    print(model.apply(params, jax.random.PRNGKey(0), jnp.ones((1, 784))).shape)
except ImportError as e:
    print("Measured:", type(e).__name__, "-", e)
Measured: ModuleNotFoundError - No module named 'haiku'
# ✅ Recommended: Haiku's init / apply correspond to nnx.split / nnx.merge in NNX
graphdef, state = nnx.split(model)          # structure + data
print("GraphDef type:", type(graphdef).__name__, "| State leaf count:", len(jax.tree.leaves(state)))

@jax.jit                                     # plain jax.jit: only accepts PyTrees
def pure_apply(state, x):                    # equivalent to Linen's model.apply(params, x)
    return nnx.merge(graphdef, state)(x)

@nnx.jit                                     # NNX's lifted jit: you can pass the module object directly
def nnx_apply(model, x):
    return model(x)

y1, y2, y3 = pure_apply(state, x_demo), nnx_apply(model, x_demo), model(x_demo)
assert max_diff(y1, y3) == 0.0 and max_diff(y2, y3) == 0.0
print("jax.jit + split/merge, nnx.jit and a direct call all give identical results ✓")

# Merging back gives a new object that is independent of the original
model_copy = nnx.merge(graphdef, state)
assert model_copy is not model and max_diff(nnx.state(model_copy), nnx.state(model)) == 0.0
GraphDef type: GraphDef | State leaf count: 4
jax.jit + split/merge, nnx.jit and a direct call all give identical results ✓

Linen / NNX / Haiku side by side

ConceptFlax Linen (old)Flax NNX (recommended)Haiku (maintenance mode)
Define a moduleclass M(nn.Module) + dataclass fields + setup / @nn.compactclass M(nnx.Module), create sublayers directly in __init__class M(hk.Module), create in __call__
Initializeparams = model.init(key, x)['params']model = M(..., rngs=nnx.Rngs(0)) (constructing it initializes it)params = hk.transform(f).init(key, x)
Forward passmodel.apply({'params': params}, x)model(x)f.apply(params, key, x)
Where parameters liveexternal dict variables['params']nnx.Param attributes on the objectexternal dict
Mutable state (BN statistics)variables['batch_stats'] + mutable=['batch_stats']nnx.BatchStat, updated in place during the forward passhk.transform_with_state returns state
Randomnessrngs={'dropout': key} passed to applynnx.Rngs(params=0, dropout=1) stored in the modulehk.next_rng_key(), pass a key to apply
Train / inference modetrain argument / deterministicmodel.train() / model.eval()is_training argument
OptimizerTrainState.create(apply_fn, params, tx)nnx.Optimizer(model, tx, wrt=nnx.Param)optax by hand
jitjax.jitnnx.jit (or jax.jit + nnx.split/merge)jax.jit
Convert to a pure PyTreealready is onennx.split(model) / nnx.state(model, filter)already is one
Checkpointing (remat)wrap with remat (nn.remat)nnx.remathk.remat
Parameter sharding annotationsnn.with_partitioningnnx.with_partitioningnothing built in

🔍 Extra: a few "new rules" for NNX in Flax 0.12

🟢 Must-read

  1. .value is deprecated: to read or write the value of an nnx.Variable, use var[...] (read) and var[...] = new (write) for arrays, and get_value() / set_value() for other types. Old code like self.count.value += 1 triggers a DeprecationWarning.
  2. nnx.Optimizer requires wrt: nnx.Optimizer(model, tx, wrt=nnx.Param). wrt is a keyword argument with no default, and it declares "which variables to optimize". To update, call optimizer.update(model, grads) (you pass the module in too).
  3. Use nnx.List for lists of modules: lists / dicts of modules stored as attributes should be wrapped in nnx.List / nnx.Dict.
  4. Two modes: the default graph mode supports shared references (the same submodule attached in two places); tree mode with graph=False treats modules as plain PyTrees and has less overhead.
class Counter(nnx.Module):
    def __init__(self):
        self.count = nnx.Variable(jnp.array(0))

c = Counter()
with warnings.catch_warnings(record=True) as w:
    warnings.simplefilter("always")
    _ = c.count.value                              # old style
dep = [str(m.message) for m in w if issubclass(m.category, DeprecationWarning)]
print("Warning when reading .value:", dep[0].splitlines()[0] if dep else None)
assert dep and "deprecated" in dep[0]

c.count[...] = c.count[...] + 1                    # ✅ new style
assert int(c.count[...]) == 1

# wrt is a required keyword argument of nnx.Optimizer
import inspect
sig = inspect.signature(nnx.Optimizer.__init__)
print("nnx.Optimizer signature:", sig)
assert sig.parameters["wrt"].default is inspect.Parameter.empty
try:
    nnx.Optimizer(model, optax.sgd(0.1))
except TypeError as e:
    print("Without wrt →", type(e).__name__, ":", e)
Warning when reading .value: '.value' access is now deprecated. For Variable[Array] instances use:
nnx.Optimizer signature: (self, model: 'M', tx: 'optax.GradientTransformation', *, wrt: 'filterlib.Filter', graph: 'bool | None' = None)
Without wrt → TypeError : Missing required argument `wrt`. As of Flax 0.11.0 the `wrt` argument is required, if you want to keep the previous use nnx.ModelAndOptimizer instead of nnx.Optimizer.

6.3 Managing model parameters and state

Training a large model means managing state such as parameters, optimizer statistics and PRNG keys. JAX's functional style requires state to be passed in explicitly as inputs and returned as outputs. That feels a bit verbose at first, but the data flow is clear and easy to debug, which is especially useful in a distributed setting.

State as PyTrees

Parameters, optimizer state (momentum buffers, the step count for learning-rate schedules), BatchNorm statistics and PRNG keys are usually organized as PyTrees. JAX transformations all work on PyTrees: taking jax.grad of a function that takes a PyTree differentiates with respect to all of its numeric leaves, and pmap / jit automatically replicate or distribute data according to the PyTree. (For PyTree basics see Basics 06 State management; we won't repeat them here.)

The course's example packs parameters, Adam state, a key and a step count into one dict and counts the leaves.

# Course example (written in the recommended style: jax.random.key, jax.tree.*)
params = {
    'encoder': {
        'layer_1': {'w': jnp.ones((128, 256)), 'b': jnp.zeros(256)},
        'layer_norm': {'scale': jnp.ones(256), 'bias': jnp.zeros(256)},
    },
    'decoder': {
        'output': {'w': jnp.ones((256, 10)), 'b': jnp.zeros(10)},
    },
}
optimizer = optax.adam(learning_rate=1e-3)
opt_state = optimizer.init(params)          # structure mirrors the parameters
key = jax.random.key(0)

training_state = {'params': params, 'opt_state': opt_state, 'rng_key': key, 'step': 0}
leaves, treedef = jax.tree.flatten(training_state)
print("Number of leaves:", len(leaves))

# Break it down item by item
n_p   = len(jax.tree.leaves(params))
n_opt = len(jax.tree.leaves(opt_state))
print(f"params {n_p} + Adam state {n_opt} (count 1 + mu {n_p} + nu {n_p}, EmptyState 0) + key 1 + step 1")
print("opt_state structure:", [type(s).__name__ for s in opt_state])
assert n_p == 6 and n_opt == 1 + 2 * n_p
assert len(leaves) == 21
Number of leaves:
 21
params 6 + Adam state 13 (count 1 + mu 6 + nu 6, EmptyState 0) + key 1 + step 1
opt_state structure: ['ScaleByAdamState', 'EmptyState']

⚠️ Note: This state dict has 21 leaves (JAX 0.12 + Optax 0.2.9), and it's easy to miscount it as 11. The parameters have 6 leaves. The state of optax.adam is (ScaleByAdamState(count, mu, nu), EmptyState()), which is 1 + 6 + 6 = 13 leaves (EmptyState has no leaves). Then add 1 each for the key and the step.

State management in the ecosystem libraries

  • Flax (Linen): packs the related state into a TrainState object (params, opt_state, step, apply_fn, tx). To update, you call state.apply_gradients(grads=...), which produces a new object instead of modifying it in place. If you need BatchNorm, you can subclass TrainState and add a batch_stats field.
  • Haiku: uses hk.transform_with_state to separate the pure function logic from the parameters params and the mutable state state (such as BN statistics). init returns (params, state) and apply returns (output, new_state), and the caller passes them around explicitly.

Below we run the course's Linen TrainState example for comparison. The Haiku transform_with_state example isn't run (Haiku isn't installed here); its NNX counterpart is nnx.BatchStat (see below).

# ❌ Old style (outdated): Linen + flax.training.train_state is the old API; the NNX counterpart is nnx.Optimizer
try:
    from typing import Any
    import flax.linen as nn
    from flax.training import train_state

    class SimpleModel(nn.Module):
        features: int
        @nn.compact
        def __call__(self, x):
            return nn.Dense(features=self.features)(x)

    class TrainState(train_state.TrainState):
        batch_stats: Any = None              # optional: BatchNorm statistics

    key = jax.random.PRNGKey(0)
    lin_model = SimpleModel(features=10)
    lin_params = lin_model.init(key, jnp.ones([1, 128]))['params']
    ts = TrainState.create(apply_fn=lin_model.apply, params=lin_params, tx=optax.adam(1e-3))
    grads = jax.tree.map(jnp.ones_like, ts.params)
    ts2 = ts.apply_gradients(grads=grads)    # immutable update: returns a new object
    print("Measured: it runs; step", int(ts.step), "→", int(ts2.step), "; the old object is unchanged")
except Exception as e:
    print("Measured error:", type(e).__name__, e)
Measured: it runs; step 0 → 1 ; the old object is unchanged
# ✅ Recommended: in NNX, nnx.Optimizer takes over the role of TrainState (step count and optimizer state live inside it, updated in place)
class SimpleModelNNX(nnx.Module):
    def __init__(self, din, features, *, rngs):
        self.dense = nnx.Linear(din, features, rngs=rngs)
    def __call__(self, x):
        return self.dense(x)

nnx_model = SimpleModelNNX(128, 10, rngs=nnx.Rngs(0))
nnx_opt = nnx.Optimizer(nnx_model, optax.adam(1e-3), wrt=nnx.Param)
grads = jax.tree.map(jnp.ones_like, nnx.state(nnx_model, nnx.Param))
w_old = nnx_model.dense.kernel[...]
nnx_opt.update(nnx_model, grads)                 # equivalent to state.apply_gradients(grads=grads)
print("nnx.Optimizer.step:", int(nnx_opt.step[...]), "| parameters updated in place:", bool(jnp.any(nnx_model.dense.kernel[...] != w_old)))
assert int(nnx_opt.step[...]) == 1 and int(ts2.step) == 1
# Same effect as one update step of the Linen TrainState (same gradients, same Adam; the first update is -lr·sign(g) in both)
assert np.allclose(nnx_model.dense.kernel[...] - w_old, ts2.params["Dense_0"]["kernel"] - ts.params["Dense_0"]["kernel"], atol=1e-7)
nnx.Optimizer.step: 1 | parameters updated in place: True

✅ State in NNX: nnx.Variable and its subclasses

NNX uses variable types to tell different kinds of state apart, and filters to pull them out by type:

TypePurposeUpdated by the optimizer?
nnx.Paramtrainable parameters (kernel / bias of nnx.Linear, etc.)yes (wrt=nnx.Param)
nnx.BatchStatBatchNorm running mean / variance, updated in place during the forward pass in training modeno
nnx.RngState (RngKey / RngCount)the keys and counters in nnx.Rngsno
custom class Foo(nnx.Variable)any custom state (counters, EMA, etc.)no (unless you put it in wrt)

nnx.split(model, nnx.Param, nnx.BatchStat, ...) splits the state into several parts by type in one go (... means "everything else"), and nnx.state(model, filter) takes just one part.

class StepCounter(nnx.Variable):          # custom variable type
    pass

class BNNet(nnx.Module):
    def __init__(self, din, dh, dout, *, rngs: nnx.Rngs):
        self.l1 = nnx.Linear(din, dh, rngs=rngs)
        self.bn = nnx.BatchNorm(dh, momentum=0.9, rngs=rngs)
        self.drop = nnx.Dropout(0.1, rngs=rngs)
        self.l2 = nnx.Linear(dh, dout, rngs=rngs)
        self.calls = StepCounter(jnp.array(0, jnp.int32))

    def __call__(self, x):
        self.calls[...] += 1                           # modify custom state in place during the forward pass
        x = nnx.relu(self.bn(self.l1(x)))
        return self.l2(self.drop(x))

net = BNNet(16, 32, 4, rngs=nnx.Rngs(params=0, dropout=1))
graphdef, p_state, bn_state, rest = nnx.split(net, nnx.Param, nnx.BatchStat, ...)
print("Param     leaves:", jax.tree.map(jnp.shape, nnx.to_pure_dict(p_state)))
print("BatchStat leaves:", jax.tree.map(jnp.shape, nnx.to_pure_dict(bn_state)))
print("Remaining state:", list(nnx.to_pure_dict(rest).keys()))
assert set(nnx.to_pure_dict(bn_state)["bn"].keys()) == {"mean", "var"}
assert "calls" in nnx.to_pure_dict(rest) and "drop" in nnx.to_pure_dict(rest)
print("BatchStat is a subclass of Variable:", issubclass(nnx.BatchStat, nnx.Variable), "| so is Param:", issubclass(nnx.Param, nnx.Variable))
Param     leaves: {'bn': {'bias': (32,), 'scale': (32,)}, 'l1': {'bias': (32,), 'kernel': (16, 32)}, 'l2': {'bias': (4,), 'kernel': (32, 4)}}
BatchStat leaves: {'bn': {'mean': (32,), 'var': (32,)}}
Remaining state: ['calls', 'drop']
BatchStat is a subclass of Variable: True | so is Param: True

✅ Composing an Optax optimizer + an nnx.Optimizer training step

The most common optimizer setup in real large-model training is global-norm gradient clipping + AdamW + a warmup-cosine learning rate. Optax uses optax.chain to string several gradient transformations together, and the state is a tuple made of each part's state.

The formula for optax.warmup_cosine_decay_schedule(init_value, peak_value, warmup_steps, decay_steps, end_value) ($t$ is the step, $W$ the number of warmup steps, $D$ the total number of steps including warmup):

$$ \eta(t)=\begin{cases} \eta_0 + (\eta_{\max}-\eta_0)\,\dfrac{t}{W}, & t< W\\[2mm] \eta_{\text{end}} + \dfrac{\eta_{\max}-\eta_{\text{end}}}{2}\left(1+\cos\dfrac{\pi\,(t-W)}{D-W}\right), & W\le t\le D \end{cases} $$

schedule = optax.warmup_cosine_decay_schedule(init_value=0.0, peak_value=1e-3,
                                              warmup_steps=100, decay_steps=1000, end_value=1e-5)
for t in [0, 50, 100, 550, 1000, 2000]:
    print(f"step {t:5d}: lr = {float(schedule(t)):.3e}")

def cosine_ref(t, W=100, D=1000, peak=1e-3, end=1e-5):
    if t < W:
        return peak * t / W
    t = min(t, D)
    return end + (peak - end) / 2 * (1 + np.cos(np.pi * (t - W) / (D - W)))
for t in [0, 37, 100, 321, 550, 999, 1000, 2000]:
    assert np.isclose(float(schedule(t)), cosine_ref(t), rtol=1e-5, atol=1e-9), t

tx = optax.chain(
    optax.clip_by_global_norm(1.0),                       # clip first
    optax.adamw(learning_rate=schedule, weight_decay=1e-4),
)

# Check the clipping: scale the gradients by hand to global norm 1, feed them to adamw, and the result should match the chain exactly
p0 = {"w": jnp.arange(6.0).reshape(2, 3), "b": jnp.ones(3)}
g0 = {"w": jnp.full((2, 3), 10.0), "b": jnp.full(3, -5.0)}
u_chain, _ = tx.update(g0, tx.init(p0), p0)
gn = optax.global_norm(g0)
g_clipped = jax.tree.map(lambda g: g / gn, g0)
aw = optax.adamw(learning_rate=schedule, weight_decay=1e-4)
u_manual, _ = aw.update(g_clipped, aw.init(p0), p0)
print(f"\nGlobal norm before clipping {float(gn):.2f} → after clipping {float(optax.global_norm(g_clipped)):.2f}")
assert max_diff(u_chain, u_manual) < 1e-9
print("chain(clip, adamw) matches manual clipping + adamw ✓")
step     0: lr = 0.000e+00
step    50: lr = 5.000e-04
step   100: lr = 1.000e-03
step   550: lr = 5.050e-04
step  1000: lr = 1.000e-05
step  2000: lr = 1.000e-05
Global norm before clipping 25.98 → after clipping 1.00
chain(clip, adamw) matches manual clipping + adamw ✓
# Train BNNet with nnx.Optimizer: in one step the parameters, BatchStat, the custom counter and the optimizer step all get updated
rng = np.random.default_rng(0)
Xb = jnp.asarray(rng.normal(size=(64, 16)), jnp.float32)
Yb = jnp.asarray(rng.normal(size=(64, 4)), jnp.float32)

net = BNNet(16, 32, 4, rngs=nnx.Rngs(params=0, dropout=1))
opt = nnx.Optimizer(net, tx, wrt=nnx.Param)

@nnx.jit
def train_step(model, optimizer, x, y):
    def loss_fn(model):
        return jnp.mean((model(x) - y) ** 2)
    loss, grads = nnx.value_and_grad(loss_fn)(model)   # differentiate only w.r.t. nnx.Param (the default filter)
    optimizer.update(model, grads)                      # update parameters and optimizer state in place
    return loss

bn_before = nnx.to_pure_dict(nnx.state(net, nnx.BatchStat))
w_before = net.l1.kernel[...]
losses = [float(train_step(net, opt, Xb, Yb)) for _ in range(5)]
print("loss:", np.round(losses, 4))
print("optimizer.step =", int(opt.step[...]), "| custom counter calls =", int(net.calls[...]))
assert int(opt.step[...]) == 5 and int(net.calls[...]) == 5
assert max_diff(bn_before, nnx.to_pure_dict(nnx.state(net, nnx.BatchStat))) > 0     # BN statistics were updated
assert max_diff(w_before, net.l1.kernel[...]) > 0                                   # parameters were updated

net.eval()                        # inference mode: BN uses running statistics, Dropout is off
y_eval1, y_eval2 = net(Xb), net(Xb)
assert max_diff(y_eval1, y_eval2) == 0.0
net.train()
print("The forward pass is deterministic in eval mode ✓")
loss: [1.4642 1.4341 1.4713 1.4325 1.4904]
optimizer.step = 5 | custom counter calls = 5
The forward pass is deterministic in eval mode ✓

✅ Cross-checking a hand-written Adam against Optax (see Basics 07, topic 5)

A hand-written Adam (with an Optax-like interface: init / update returns (updates, new_state), $\Delta\theta=-\eta\,\hat m/(\sqrt{\hat v}+\epsilon)$) was already implemented in the Extra of section 6.6 in Basics 06 State management. The full version, with hand-written SGD / momentum / Adam / AdamW / chain cross-checked step by step against Optax, is in Basics 07 Practical topics, topic 5. We won't repeat it here and use Optax directly below.

State in a distributed setting

The course uses a diagram to show where each kind of state goes in pmap data parallelism:

  • Parameters and optimizer state: in standard data parallelism they are identical on every device. You initialize them once on the host and replicate them to each device. After each device computes its gradients, lax.pmean averages them across devices first, and the update then runs inside the pmap function, so all replicas stay in sync.
  • PRNG keys: a common and subtle source of bugs. If every device uses the same key for dropout or data augmentation, they generate the same random numbers. The course's approach: on the host, split the main key into one subkey per device and pass them to pmap as an array; inside the function, split again to get the key for this step, and return the new key for the next step.

Course code (splitting keys per device in pmap):

# ❌ Old style (outdated): pmap + one key per device; with jit + sharding's global view you only need one key (see the next cell)
try:
    num_devices = jax.local_device_count()
    key = jax.random.PRNGKey(42)
    device_keys = jax.random.split(key, num_devices)       # split once on the host

    @partial(jax.pmap, axis_name='devices')
    def dropout_pmap(local_key, x):
        dropout_key, new_local_key = jax.random.split(local_key)
        keep = jax.random.bernoulli(dropout_key, 0.5, x.shape)
        return jnp.where(keep, x / 0.5, 0.0), new_local_key

    xs = jnp.ones((num_devices, 4, 8))
    masked, device_keys = dropout_pmap(device_keys, xs)
    same = [bool(jnp.all(masked[0] == masked[i])) for i in range(1, num_devices)]
    print("Measured: pmap still runs; number of devices whose mask equals device 0's =", sum(same))
except Exception as e:
    print("Measured error:", type(e).__name__, e)
Measured: pmap still runs; number of devices whose mask equals device 0's = 0
# ✅ Recommended: jit + NamedSharding is a "global view": generate random numbers for the whole global array, and each device automatically gets a different slice
mesh_dp = jax.make_mesh((8,), ("data",))
x_global = jnp.ones((32, 8))
key = jax.random.key(42)

@jax.jit
def dropout_global(key, x):
    keep = jax.random.bernoulli(key, 0.5, x.shape)
    return jnp.where(keep, x / 0.5, 0.0)

with jax.set_mesh(mesh_dp):
    x_sharded = jax.device_put(x_global, P("data"))
    y_sharded = dropout_global(key, x_sharded)
    print("Input:", jax.typeof(x_sharded), "→ output:", jax.typeof(y_sharded))
y_single = dropout_global(key, x_global)

# jax_threefry_partitionable defaults to True: sharded and single-device computation produce exactly the same random mask
print("jax_threefry_partitionable =", jax.config.jax_threefry_partitionable)
assert jax.config.jax_threefry_partitionable
assert bool(jnp.all(y_sharded == y_single))
blocks = np.asarray(y_sharded).reshape(8, 4, 8)
assert not all(np.array_equal(blocks[0], blocks[i]) for i in range(1, 8))
print("The sharded mask matches the single-device mask element for element, and each device's slice is different ✓")
Input: float32[32@data,8] → output: float32[32@data,8]
jax_threefry_partitionable = True
The sharded mask matches the single-device mask element for element, and each device's slice is different ✓

Serialization and checkpoints

Training can run for hours or weeks, so you have to save the training state to disk regularly to be able to resume after an interruption. You need to save:

  1. Model parameters;
  2. Optimizer state (required for optimizers with momentum or adaptive learning rates, otherwise training behaves differently after resuming);
  3. Step counter (progress and learning-rate schedules depend on it);
  4. PRNG key (optional but recommended, for reproducibility);
  5. Other state, such as BatchNorm running statistics.

The course mentions that Orbax (orbax.checkpoint) is becoming the standard solution. It supports asynchronous checkpointing (writing to disk in the background without blocking training), and it's flexible about how things are saved. Flax also has an older flax.training.checkpoints. After restoring, the course code uses jax.device_put_replicated to copy the state back to each device:

# ❌ Wrong (outdated): jax.device_put_replicated was removed in JAX 0.12
try:
    restored_state = {"w": jnp.ones(3)}
    state = jax.device_put_replicated(restored_state, jax.local_devices())
    print("Measured: it runs")
except AttributeError as e:
    print("Measured error: AttributeError:", e)
Measured error: AttributeError: jax.device_put_replicated is deprecated; use jax.device_put instead. See https://docs.jax.dev/en/latest/migrate_pmap.html#drop-in-replacements for a drop-in replacement.
# ✅ Recommended: save / restore NNX state with Orbax; on restore, data goes straight to the target devices according to the "abstract state" (shape + dtype + sharding)
import orbax.checkpoint as ocp

ckpt_dir = tempfile.mkdtemp(prefix="adv06_ckpt_")
net_src = BNNet(16, 32, 4, rngs=nnx.Rngs(params=0, dropout=1))
opt_src = nnx.Optimizer(net_src, tx, wrt=nnx.Param)
train_step(net_src, opt_src, Xb, Yb)

def pack(model, optimizer):
    """Training state to save: parameters, BN statistics, optimizer (including the step). All pure PyTrees."""
    return {"params": nnx.state(model, nnx.Param),
            "batch_stats": nnx.state(model, nnx.BatchStat),
            "optimizer": nnx.state(optimizer)}

checkpointer = ocp.StandardCheckpointer()
checkpointer.save(os.path.join(ckpt_dir, "step_1"), pack(net_src, opt_src))
checkpointer.wait_until_finished()                 # saving is asynchronous; wait for it to finish before reading back

# Restore: first build an "abstract" model (shapes only, no memory allocated) and use its state as the restore template
abs_net = nnx.eval_shape(lambda: BNNet(16, 32, 4, rngs=nnx.Rngs(params=0, dropout=1)))
abs_opt = nnx.eval_shape(lambda: nnx.Optimizer(abs_net, tx, wrt=nnx.Param))
restored = checkpointer.restore(os.path.join(ckpt_dir, "step_1"), pack(abs_net, abs_opt))

net_new = BNNet(16, 32, 4, rngs=nnx.Rngs(params=123, dropout=1))   # arbitrary initialization
opt_new = nnx.Optimizer(net_new, tx, wrt=nnx.Param)
nnx.update(net_new, restored["params"], restored["batch_stats"])   # write the restored state back
nnx.update(opt_new, restored["optimizer"])
print("optimizer.step after restore =", int(opt_new.step[...]))
assert int(opt_new.step[...]) == 1
assert max_diff(nnx.state(net_new, nnx.Param), nnx.state(net_src, nnx.Param)) == 0.0
assert max_diff(nnx.state(net_new, nnx.BatchStat), nnx.state(net_src, nnx.BatchStat)) == 0.0
assert max_diff(nnx.state(opt_new), nnx.state(opt_src)) == 0.0
print("Orbax save → restore: parameters, BN statistics and optimizer state are bit-identical ✓")
optimizer.step after restore = 1
Orbax save → restore: parameters, BN statistics and optimizer state are bit-identical ✓

⚠️ Note: jax.device_put_replicated(restored_state, jax.local_devices()) no longer works in JAX 0.12. This function (along with jax.device_put_sharded) has been removed. Calling it raises AttributeError, and the error message tells you to use jax.device_put instead. With the jit + sharding style you don't even need a "replicate" step: put NamedSharding in the restore template and Orbax places the data directly on the right devices.

🔍 Extra: a few practical points about checkpoints

⚪ Optional: when you run long training jobs and need to resume from checkpoints

  • Always save the step. Learning-rate schedules and the accumulation counter of optax.MultiSteps depend on it. If you only save the parameters and not the optimizer state and step, the learning rate starts warming up from scratch after resuming.
  • Use nnx.eval_shape for the restore template: it only traces shapes and doesn't actually initialize anything, so even a model with tens of billions of parameters takes no memory. Set the shardings in the template and you can restore directly into shards (each process reads only its own part).
  • The data loader's position is part of the training state too (Grain iterators can be checkpointed). Otherwise you'll repeat or skip data after resuming.

6.4 Combining pmap with training frameworks

The framework (Flax / Haiku) handles model definition and state management, and pmap distributes the computation across devices. Put together, that's the data-parallel training the course describes. The core idea is still SPMD (single program, multiple data): write the function as if for a single device, and pmap runs it on every device at once, with each device processing a different slice of the input. (For the basics of pmap, axis_name and pmean, see Basics 05 Multi-device parallelism with pmap.)

The course's workflow has four steps:

  1. Replicate the model state: initialize parameters and optimizer state on the host, pack them into a TrainState, and replicate it to all devices. After replication, every leaf gets an extra leading dimension whose size equals the number of devices.
  2. Shard the data: split the global batch $B$ evenly over $N$ devices, $B/N$ examples each. Use tree_map to reshape every array in the batch to (N, B/N, ...). This has to be done before calling the pmapped function.
  3. Define and map the training step: each device computes loss and gradients on its local data. Before updating, you must average the gradients across devices with lax.pmean(grads, axis_name), then update with state.apply_gradients. jax.pmap(train_step, axis_name='batch_axis') ties the axis_name to the pmean inside the function. The returned state is identical on every device, while the loss is each device's own local value (shape (N,)). Metrics such as accuracy can also be pmeaned.
  4. Handle RNG: random operations like dropout must not use the same key on every device. split one subkey per device on the host and pass them in.

First we run the course's code for step 1.

import flax.linen as nn
from flax.training import train_state

class LinenCNN(nn.Module):               # a minimal model added here to serve as `model`
    @nn.compact
    def __call__(self, x):
        x = x.reshape((x.shape[0], -1))
        x = nn.relu(nn.Dense(32)(x))
        return nn.Dense(10)(x)

model = LinenCNN()
key = jax.random.key(0)
lparams = model.init(key, jnp.ones([1, 28, 28, 1]))['params']
tx_adam = optax.adam(learning_rate=1e-3)
state = train_state.TrainState.create(apply_fn=model.apply, params=lparams, tx=tx_adam)
num_devices = jax.local_device_count()
# ❌ Wrong (outdated): jax.device_put_replicate is misspelled, and the correctly spelled jax.device_put_replicated was also removed in 0.12
for fn_name in ["device_put_replicate", "device_put_replicated"]:
    try:
        replicated_state = getattr(jax, fn_name)(state, jax.local_devices())
        print(fn_name, "→ it runs")
    except AttributeError as e:
        print(f"jax.{fn_name} → AttributeError: {str(e)[:110]}")
jax.device_put_replicate → AttributeError: module 'jax' has no attribute 'device_put_replicate'
jax.device_put_replicated → AttributeError: jax.device_put_replicated is deprecated; use jax.device_put instead. See https://docs.jax.dev/en/latest/migrat
# ✅ Recommended (if you still want pmap): following the official migration guide, use jax.device_put to place a "stacked, split by device" array
def replicate_for_pmap(tree, devices):
    mesh = jax.sharding.Mesh(np.array(devices), ("d",))
    n = len(devices)
    return jax.tree.map(
        lambda x: jax.device_put(jnp.broadcast_to(x, (n,) + jnp.shape(x)), NamedSharding(mesh, P("d"))),
        tree)

replicated_state = replicate_for_pmap(state, jax.local_devices())
print(f"State replicated across {num_devices} devices.")
print("kernel shape after replication:", replicated_state.params['Dense_0']['kernel'].shape)
print("original kernel shape:         ", state.params['Dense_0']['kernel'].shape)
assert replicated_state.params['Dense_0']['kernel'].shape == (num_devices,) + state.params['Dense_0']['kernel'].shape
State replicated across 8 devices.
kernel shape after replication: (8, 784, 32)
original kernel shape:          (784, 32)
# Course steps 2 and 3: shard the data + pmapped training step
global_batch_size = 64
local_batch_size = global_batch_size // num_devices

def shard_batch(batch):
    """Reshape (B, ...) into (num_devices, B/num_devices, ...)."""
    return jax.tree.map(lambda x: x.reshape((num_devices, local_batch_size) + x.shape[1:]), batch)

rng = np.random.default_rng(0)
global_images = jnp.asarray(rng.normal(size=(global_batch_size, 28, 28, 1)), jnp.float32)
global_labels = jnp.asarray(rng.integers(0, 10, global_batch_size), jnp.int32)
batch = {'image': global_images, 'label': global_labels}
sharded_batch = shard_batch(batch)
print("Global image shape:", global_images.shape, "→ sharded image shape:", sharded_batch['image'].shape)

def compute_loss(params, batch, apply_fn):
    logits = apply_fn({'params': params}, batch['image'])
    one_hot = jax.nn.one_hot(batch['label'], num_classes=10)
    return jnp.mean(-jnp.sum(one_hot * jax.nn.log_softmax(logits), axis=-1))

def train_step(state, batch):
    loss, grads = jax.value_and_grad(compute_loss)(state.params, batch, state.apply_fn)
    grads = jax.lax.pmean(grads, axis_name='batch_axis')      # key step: average gradients across devices
    return state.apply_gradients(grads=grads), loss
Global image shape: (64, 28, 28, 1) → sharded image shape: (8, 8, 28, 28, 1)

Modern style: NNX + Optax + jit + NamedSharding

Below we first run the course's pmap training step for comparison, and right after it show the modern version. JAX 0.12 recommends writing data parallelism from a global view:

  • Create a Mesh and split the batch dimension across devices with P('data'). Parameters without sharding annotations are replicated.
  • The training step is ordinary single-device code compiled with nnx.jit (or jax.jit). When the compiler sees "data sharded, parameters replicated", it inserts the gradient all-reduce automatically. You don't need pmean, and you don't need to reshape out a device dimension or replicate state by hand.
  • The returned loss is the average over the global batch, not one value per device.

In 0.12, jax.make_mesh creates Explicit axes by default: sharding is part of the array's type, and jax.typeof(x) shows something like float32[64@data,28,28,1] (@data means this dimension is split along the data axis).

# ❌ Old style (outdated): pmap still works in 0.12, but the official advice is to migrate to jit + sharding (or shard_map)
try:
    p_train_step = jax.pmap(train_step, axis_name='batch_axis')
    pmap_state, local_losses = p_train_step(replicated_state, sharded_batch)
    print("Measured: pmap runs; local_losses shape", local_losses.shape, "| mean loss", f"{float(local_losses.mean()):.4f}")
    print("Sharding type of the pmap output:", type(local_losses.sharding).__name__, local_losses.sharding.spec)
except Exception as e:
    print("Measured error:", type(e).__name__, e)
Measured: pmap runs; local_losses shape (8,) | mean loss 2.4516
Sharding type of the pmap output: NamedSharding P('batch_axis',)
# ✅ Recommended: NNX + Optax + jit + NamedSharding data parallelism (no pmean / manual reshape / state replication needed)
class NNXMLP(nnx.Module):
    def __init__(self, *, rngs):
        self.l1 = nnx.Linear(784, 32, rngs=rngs)
        self.l2 = nnx.Linear(32, 10, rngs=rngs)
    def __call__(self, x):
        return self.l2(nnx.relu(self.l1(x.reshape(x.shape[0], -1))))

@nnx.jit
def nnx_train_step(model, optimizer, x, y):
    def loss_fn(model):
        return optax.softmax_cross_entropy_with_integer_labels(model(x), y).mean()
    loss, grads = nnx.value_and_grad(loss_fn)(model)
    optimizer.update(model, grads)
    return loss

mesh_dp = jax.make_mesh((8,), ("data",))
print("mesh:", mesh_dp)
with jax.set_mesh(mesh_dp):
    dp_model = NNXMLP(rngs=nnx.Rngs(0))
    dp_opt = nnx.Optimizer(dp_model, optax.adam(1e-3), wrt=nnx.Param)
    x_dp = jax.device_put(global_images, P("data"))          # shard the batch dimension
    y_dp = jax.device_put(global_labels, P("data"))
    print("x:", jax.typeof(x_dp), "| kernel:", jax.typeof(dp_model.l1.kernel[...]))
    dp_losses = [float(nnx_train_step(dp_model, dp_opt, x_dp, y_dp)) for _ in range(3)]
    jax.debug.visualize_array_sharding(x_dp.reshape(64, -1)[:, :8])

# Single-device reference: same initialization, same data
ref_model = NNXMLP(rngs=nnx.Rngs(0))
ref_opt = nnx.Optimizer(ref_model, optax.adam(1e-3), wrt=nnx.Param)
ref_losses = [float(nnx_train_step(ref_model, ref_opt, global_images, global_labels)) for _ in range(3)]
print("Data-parallel loss:", np.round(dp_losses, 6))
print("Single-device loss:", np.round(ref_losses, 6))
d = max_diff(nnx.state(dp_model, nnx.Param), nnx.state(ref_model, nnx.Param))
print(f"Max parameter difference after 3 steps: {d:.2e}")
assert np.allclose(dp_losses, ref_losses, rtol=1e-5) and d < 1e-5

# Take a look at the communication the compiler inserted automatically
with jax.set_mesh(mesh_dp):
    hlo = nnx_train_step.lower(dp_model, dp_opt, x_dp, y_dp).compile().as_text()
print("Automatically inserted collectives:", collectives(hlo))
assert collectives(hlo)["all-reduce"] >= 1
mesh: Mesh('data': 8, axis_types=(Explicit,))
x: float32[64@data,28,28,1] | kernel: float32[784,32]
  CPU 0  
         
  CPU 1  
         
  CPU 2  
         
  CPU 3  
         
  CPU 4  
         
  CPU 5  
         
  CPU 6  
         
  CPU 7  
         
Data-parallel loss: [2.569236 2.263043 1.98849 ]
Single-device loss: [2.569236 2.263043 1.98849 ]
Max parameter difference after 3 steps: 2.98e-08
Automatically inserted collectives: {'all-reduce': 1, 'all-gather': 0, 'reduce-scatter': 0, 'all-to-all': 0, 'collective-permute': 0}
# The pmap version also matches the modern style (the course's Linen model vs single-device jit)
@jax.jit
def single_step(state, batch):
    loss, g = jax.value_and_grad(compute_loss)(state.params, batch, state.apply_fn)
    return state.apply_gradients(grads=g), loss

s1, l1 = single_step(state, batch)
d = max_diff(jax.tree.map(lambda x: x[0], pmap_state.params), s1.params)
print(f"pmap replica 0 vs single device: parameter difference {d:.2e}; mean loss {float(local_losses.mean()):.6f} vs {float(l1):.6f}")
assert d < 1e-6 and np.isclose(float(local_losses.mean()), float(l1), rtol=1e-6)
# All replicas are exactly identical
assert all(max_diff(jax.tree.map(lambda x: x[i], pmap_state.params),
                    jax.tree.map(lambda x: x[0], pmap_state.params)) == 0 for i in range(num_devices))
pmap replica 0 vs single device: parameter difference 7.08e-08; mean loss 2.451626 vs 2.451626

✅ Parameter sharding annotations: nnx.with_partitioning

In data parallelism the parameters are replicated. When parameters get too big, you have to split the parameters themselves too (covered in detail in section 6.8). NNX does this by adding a sharding annotation to the initializer: nnx.with_partitioning(initializer, (axis_name, ...)). Each entry of the tuple corresponds to one dimension of the parameter, and None means that dimension isn't split.

  • When you construct a module inside a jax.set_mesh(mesh) context, NNX places the parameters on the corresponding devices right away according to the annotations (eager sharding). Without a mesh context, constructing annotated parameters raises an error.
  • nnx.get_partition_spec(state) returns the PartitionSpec of every variable.
  • Optimizer state automatically inherits the parameters' sharding: nnx.Optimizer creates $m$ and $v$ with zeros_like, so their shape and sharding match the parameters.
mesh2d = jax.make_mesh((2, 4), ("data", "model"))

class ShardedMLP(nnx.Module):
    def __init__(self, din, dh, dout, *, rngs):
        init = nnx.initializers.lecun_normal()
        # First layer is split by columns (output dim → model axis), second layer by rows (input dim → model axis)
        self.up = nnx.Linear(din, dh, use_bias=False,
                             kernel_init=nnx.with_partitioning(init, (None, "model")), rngs=rngs)
        self.down = nnx.Linear(dh, dout, use_bias=False,
                               kernel_init=nnx.with_partitioning(init, ("model", None)), rngs=rngs)
    def __call__(self, x):
        h = nnx.relu(self.up(x))
        # In Explicit mode, a matmul whose contracting dim is split needs an explicit output sharding (here: summed, then sharded by batch)
        return self.down(h, out_sharding=P("data", None))

try:
    ShardedMLP(64, 256, 64, rngs=nnx.Rngs(0))
except ValueError as e:
    print("No mesh context → ValueError:", str(e)[:100], "...")

with jax.set_mesh(mesh2d):
    smodel = ShardedMLP(64, 256, 64, rngs=nnx.Rngs(0))
    sopt = nnx.Optimizer(smodel, optax.adam(1e-3), wrt=nnx.Param)
    print("Sharding annotations:", jax.tree.map(str, nnx.to_pure_dict(nnx.get_partition_spec(nnx.state(smodel)))))
    print("up.kernel   :", jax.typeof(smodel.up.kernel[...]))
    print("down.kernel :", jax.typeof(smodel.down.kernel[...]))
    mu_up = sopt.opt_state[0].mu["up"]["kernel"][...]
    print("Adam mu(up) :", jax.typeof(mu_up))
    assert mu_up.sharding.spec == smodel.up.kernel[...].sharding.spec == P(None, "model")
    # Each device stores only 1/4 of up.kernel
    k = smodel.up.kernel[...]
    print(f"up.kernel: {k.nbytes} B global, {k.addressable_shards[0].data.nbytes} B per device")
    assert k.addressable_shards[0].data.nbytes * 4 == k.nbytes
    jax.debug.visualize_array_sharding(k)
No mesh context → ValueError: An auto mesh context or metadata is required if creating a variable with annotation out_sharding=P(N ...
Sharding annotations: {'down': {'kernel': "P('model', None)"}, 'up': {'kernel': "P(None, 'model')"}}
up.kernel   : float32[64,256@model]
down.kernel : float32[256@model,64]
Adam mu(up) : float32[64,256@model]
up.kernel: 65536 B global, 16384 B per device
                                                                                
                                                                                
                                                                                
                                                                                
                                                                                
         C…                  C…                  C…                  C…         
         0…                  1…                  2…                  3…         
                                                                                
                                                                                
                                                                                
                                                                                
                                                                                

🔍 Extra: key points for migrating from pmap to jit + sharding

⚪ Optional: when you maintain or migrate old pmap code

pmap stylejit + NamedSharding style
Manually reshape out a device dimension (N, B/N, ...)Global array (B, ...), device_put(x, P('data'))
Replicate state with device_put_replicated (removed in 0.12)Parameters without annotations are replicated by default
lax.pmean(grads, 'batch') inside the functionThe compiler inserts the all-reduce automatically based on the shardings
One PRNG key per deviceOne key, sample over the global shape (jax_threefry_partitionable=True)
Return values carry a device dimension; you need x[0] to get them backWhat comes back is the global value
Only one-dimensional device mappingA mesh with any number of dimensions (data × model × ...), and parameters can be sharded too
When you need hand-written per-device codeUse jax.shard_map (write collectives like psum explicitly)

One more thing: in 0.12 the output of pmap is already a NamedSharding (the P('batch_axis',) printed above), which shows that pmap is built on the same sharding machinery internally. New code can just use jit. If you translate a pmap training step one to one with shard_map, the easiest trap to fall into is "the gradients silently become $N$ times larger (N = number of devices)". The full pmap → shard_map migration checklist is in the "gradients become N times larger in shard_map" topic of Basics 05 Multi-device parallelism with pmap.


6.5 Gradient accumulation

When training a large model, the batch you want may not fit in a single accelerator's memory. Even if it fits, you may want a larger batch for training stability. The idea of gradient accumulation is: process several small micro-batches one after another, add up their gradients, and do only one optimizer update at the end. That separates "how many examples go into one update" from "how many examples sit in memory at the same time".

How it works

Let the effective batch be $B$ and the micro-batch size $b$. The number of accumulation steps $N$ satisfies $B=N\times b$.

  1. Initialize an accumulated gradient with the same structure as the parameters, all zeros;
  2. Loop $N$ times: take a micro-batch → forward pass → compute gradients → add them to the accumulated gradient;
  3. After the loop, divide the accumulated gradient by $N$, use it to update the parameters, and reset the accumulated gradient to zero.

$$ \nabla_\theta \mathcal L_{\text{eff}}=\frac1N\sum_{i=1}^{N}\nabla_\theta \mathcal L(\text{micro-batch}_i;\theta) $$

The main benefit: while computing gradients, memory only needs to hold the activations of one micro-batch at a time.

⚠️ Note: The accumulated average gradient doesn't just "approximate" the large-batch gradient. As long as the loss is a per-example mean, all micro-batches have the same size, and the model has no cross-example operations like BatchNorm, the two are exactly equal mathematically and differ numerically only by float32 rounding (about $10^{-7}$). Below we verify this with assert, then show the three cases that really do cause differences: a summed loss, unequal micro-batch sizes, and BatchNorm.

The course's implementation

The course first shows a plain train_step, then splits the accumulating version into two parts. One is a jit-compiled micro_batch_step that computes the gradient of one micro-batch and returns the logits via has_aux=True. The other is a Python loop, accumulated_train_step, that sums and averages with tree_map and defines a @jax.jit apply_update inside the function to apply the optimizer. The course's view is that data loading usually isn't put inside jit, so the loop isn't compiled as a whole. It also mentions that you can use lax.scan to put the accumulation loop inside one jit function, but that this "requires careful management of state and data loading".

First, set up a small model and data (both parameters and data are passed into the jit functions as arguments).

rng = np.random.default_rng(0)
X_acc = jnp.asarray(rng.normal(size=(64, 16)), jnp.float32)
Y_acc = jnp.asarray(rng.normal(size=(64, 1)), jnp.float32)
acc_params = {"w1": jnp.asarray(rng.normal(size=(16, 32)) * 0.3, jnp.float32), "b1": jnp.zeros(32),
              "w2": jnp.asarray(rng.normal(size=(32, 1)) * 0.3, jnp.float32), "b2": jnp.zeros(1)}

def mlp_apply(p, x):
    return jnp.tanh(x @ p["w1"] + p["b1"]) @ p["w2"] + p["b2"]

def loss_fn(p, x, y):
    return jnp.mean((mlp_apply(p, x) - y) ** 2)     # per-example mean

acc_tx = optax.adam(1e-2)
N_ACC = 4

def micro_batches(x, y, n):
    return [(x[i * len(x) // n:(i + 1) * len(x) // n], y[i * len(y) // n:(i + 1) * len(y) // n]) for i in range(n)]
# ❌ Wrong: apply_update is defined inside the function, so every call creates a new function object → it's retraced/recompiled every step
trace_count = {"n": 0}

@jax.jit
def micro_batch_step(params, x, y):
    return jax.value_and_grad(loss_fn)(params, x, y)

def accumulated_train_step(params, opt_state, data_iterator, accumulation_steps):
    accumulated_grads = jax.tree.map(jnp.zeros_like, params)
    total_loss = 0.0
    for _ in range(accumulation_steps):
        x, y = next(data_iterator)
        loss, grads = micro_batch_step(params, x, y)
        accumulated_grads = jax.tree.map(lambda a, g: a + g, accumulated_grads, grads)
        total_loss += loss
    averaged_grads = jax.tree.map(lambda g: g / accumulation_steps, accumulated_grads)

    @jax.jit
    def apply_update(p, o_state, avg_grads):
        trace_count["n"] += 1                       # only runs during tracing
        updates, new_o_state = acc_tx.update(avg_grads, o_state, p)
        return optax.apply_updates(p, updates), new_o_state

    params, opt_state = apply_update(params, opt_state, averaged_grads)
    return params, opt_state, total_loss / accumulation_steps

try:
    p_course, s_course = acc_params, acc_tx.init(acc_params)
    for step in range(5):
        p_course, s_course, _ = accumulated_train_step(p_course, s_course,
                                                       iter(micro_batches(X_acc, Y_acc, N_ACC)), N_ACC)
    print(f"Measured: in 5 training steps, apply_update was traced {trace_count['n']} times")
except Exception as e:
    print("Measured error:", type(e).__name__, e)
assert trace_count["n"] == 5
Measured: in 5 training steps, apply_update was traced 5 times
# ✅ Recommended: put the whole accumulation loop in one jit and iterate over micro-batches with lax.scan (reshape the data to (N, b, ...) first)
@partial(jax.jit, static_argnames="n_micro")
def accum_grads_scan(params, x, y, n_micro):
    xs = x.reshape(n_micro, -1, *x.shape[1:])
    ys = y.reshape(n_micro, -1, *y.shape[1:])
    def body(carry, mb):
        acc, loss_sum = carry
        loss, g = jax.value_and_grad(loss_fn)(params, *mb)
        return (jax.tree.map(jnp.add, acc, g), loss_sum + loss), None
    init = (jax.tree.map(jnp.zeros_like, params), 0.0)
    (acc, loss_sum), _ = lax.scan(body, init, (xs, ys))
    return jax.tree.map(lambda a: a / n_micro, acc), loss_sum / n_micro

@partial(jax.jit, static_argnames="n_micro", donate_argnums=(0, 1))
def accum_train_step(params, opt_state, x, y, n_micro):
    grads, loss = accum_grads_scan(params, x, y, n_micro)
    updates, opt_state = acc_tx.update(grads, opt_state, params)
    return optax.apply_updates(params, updates), opt_state, loss

g_big = jax.jit(jax.grad(loss_fn))(acc_params, X_acc, Y_acc)      # the whole batch at once
g_acc, _ = accum_grads_scan(acc_params, X_acc, Y_acc, N_ACC)
d = max_diff(g_big, g_acc)
print(f"lax.scan accumulating 4 micro-batches vs one large batch: max gradient difference {d:.2e}")
assert d < 1e-6

# Also matches the step-by-step accumulation above (5 training steps)
p_new, s_new = jax.tree.map(jnp.copy, acc_params), acc_tx.init(acc_params)
for step in range(5):
    p_new, s_new, _ = accum_train_step(p_new, s_new, X_acc, Y_acc, n_micro=N_ACC)
print(f"After 5 steps, scan version vs step-by-step version parameter difference: {max_diff(p_new, p_course):.2e}")
assert max_diff(p_new, p_course) < 1e-5
lax.scan accumulating 4 micro-batches vs one large batch: max gradient difference 3.58e-07
After 5 steps, scan version vs step-by-step version parameter difference: 2.98e-08

✅ optax.MultiSteps: make accumulation part of the optimizer

optax.MultiSteps(opt, every_k_schedule=k) can wrap any optimizer. For the first $k-1$ calls to update it only accumulates the gradient into its own state and returns an all-zero update. Only on the $k$-th call does it call the inner optimizer with the averaged gradient (use_grad_mean=True, the default) and reset the accumulator. The training loop doesn't change at all: you just call the training step once per micro-batch as usual. The cost is one extra accumulation buffer in the state, the same size as the parameters (you can set accumulator_dtype to bf16 to save memory).

ms = optax.MultiSteps(optax.adam(1e-2), every_k_schedule=N_ACC)
ms_state = ms.init(acc_params)
p_ms = acc_params

@jax.jit
def ms_step(p, s, x, y):
    g = jax.grad(loss_fn)(p, x, y)
    u, s = ms.update(g, s, p)
    return optax.apply_updates(p, u), s, u

for i, (xb, yb) in enumerate(micro_batches(X_acc, Y_acc, N_ACC)):
    p_ms, ms_state, u = ms_step(p_ms, ms_state, xb, yb)
    print(f"Micro-batch {i}: |update|max = {float(max(jnp.abs(l).max() for l in jax.tree.leaves(u))):.4f}, "
          f"mini_step = {int(ms_state.mini_step)}, gradient_step = {int(ms_state.gradient_step)}")
    if i < N_ACC - 1:
        assert all(float(jnp.abs(l).max()) == 0.0 for l in jax.tree.leaves(u))   # the first k-1 updates are zero

# Matches "one Adam update on the large batch"
adam = optax.adam(1e-2)
u_big, _ = adam.update(g_big, adam.init(acc_params), acc_params)
p_big = optax.apply_updates(acc_params, u_big)
print(f"MultiSteps(4) vs large-batch Adam: parameter difference {max_diff(p_ms, p_big):.2e}")
assert max_diff(p_ms, p_big) < 1e-6
print("MultiSteps extra state (accumulation buffer) bytes / parameter bytes =",
      tree_bytes(ms_state.acc_grads) / tree_bytes(acc_params))
Micro-batch 0: |update|max = 0.0000, mini_step = 1, gradient_step = 0
Micro-batch 1: |update|max = 0.0000, mini_step = 2, gradient_step = 0
Micro-batch 2: |update|max = 0.0000, mini_step = 3, gradient_step = 0
Micro-batch 3: |update|max = 0.0100, mini_step = 0, gradient_step = 1
MultiSteps(4) vs large-batch Adam: parameter difference 7.45e-09
MultiSteps extra state (accumulation buffer) bytes / parameter bytes = 1.0

Working with pmap, and the trade-offs

  • Combining with data parallelism: each device accumulates $N$ steps independently on its own data shard, then lax.pmean averages the accumulated gradients across devices (only one round of communication), and every device applies the same update with the globally averaged gradient. Each device only has to hold activations for $b/\text{num devices}$ examples, and the effective batch is still $B$.
  • Memory vs. compute: peak memory drops noticeably and the effective batch can be larger. The cost is $N$ forward and backward passes per update. If the larger batch means convergence needs fewer steps, total compute doesn't necessarily go up.
  • It doesn't save parameter or optimizer-state memory: accumulation only saves memory related to activations. For parameters and Adam moments you still need mixed precision or optimizer-state sharding (6.7, 6.8).
  • BatchNorm: statistics are computed per micro-batch, so they're noisy when micro-batches are small. You can switch to LayerNorm, or synchronize BN statistics across devices (more complicated).
  • Learning rate: the gradient is already averaged over $N$ steps, so in theory the learning rate usually doesn't need big changes, though you may still need to tune it empirically.

Below is "data parallelism + gradient accumulation" in the modern style: the data is sharded along the data axis, the micro-batch dimension comes first and is iterated by scan, and each micro-batch is itself sharded along its batch dimension. The whole function is still written as single-device code.

with jax.set_mesh(mesh_dp):
    # Shape (N, b, ...): the micro-batch dim isn't split, the batch dim is split along the data axis
    xs_dp = jax.device_put(X_acc.reshape(N_ACC, -1, 16), P(None, "data"))
    ys_dp = jax.device_put(Y_acc.reshape(N_ACC, -1, 1), P(None, "data"))
    print("Micro-batch data:", jax.typeof(xs_dp))

    @jax.jit
    def dp_accum(params, xs, ys):
        def body(acc, mb):
            return jax.tree.map(jnp.add, acc, jax.grad(loss_fn)(params, *mb)), None
        acc, _ = lax.scan(body, jax.tree.map(jnp.zeros_like, params), (xs, ys))
        return jax.tree.map(lambda a: a / xs.shape[0], acc)

    g_dp = dp_accum(acc_params, xs_dp, ys_dp)
    hlo = dp_accum.lower(acc_params, xs_dp, ys_dp).compile().as_text()
print(f"Data parallel + accumulation vs single-device large batch: gradient difference {max_diff(g_dp, g_big):.2e}; collectives {collectives(hlo)}")
assert max_diff(g_dp, g_big) < 1e-6
Micro-batch data: float32[4,16@data,16]
Data parallel + accumulation vs single-device large batch: gradient difference 2.38e-07; collectives {'all-reduce': 1, 'all-gather': 0, 'reduce-scatter': 0, 'all-to-all': 0, 'collective-permute': 0}

🔍 Extra: when is accumulation not equal to a large batch?

🟢 Must-read

  1. Summed loss: if the loss is a $\sum$ instead of a mean, the large-batch gradient is the sum of the micro-batch gradients. Dividing by $N$ is then wrong: the gradient becomes $N$ times smaller, which silently divides the learning rate by $N$.
  2. Unequal micro-batch sizes (for example a partial last batch, or a language model whose loss is computed per token, so each micro-batch has a different number of valid tokens): directly averaging "the mean of each micro-batch" gives more weight to micro-batches with fewer examples. The right way is to weight by the number of examples (or tokens): accumulate the $\sum$ (the summed loss), then divide by the total number of examples at the end.
  3. BatchNorm: statistics are computed per micro-batch, so the forward pass itself is different and the gradients can't match. BN's running statistics also get updated $N$ times instead of once.
mbs = micro_batches(X_acc, Y_acc, N_ACC)

# 1) Summed loss: after accumulating you should "add", not "average"
def loss_sum(p, x, y):
    return jnp.sum((mlp_apply(p, x) - y) ** 2)
gs_big = jax.grad(loss_sum)(acc_params, X_acc, Y_acc)
gs_list = [jax.grad(loss_sum)(acc_params, xb, yb) for xb, yb in mbs]
gs_sum = jax.tree.map(lambda *g: sum(g), *gs_list)
gs_mean = jax.tree.map(lambda *g: sum(g) / N_ACC, *gs_list)
print(f"Summed loss: adding {max_diff(gs_sum, gs_big):.2e}  |  wrongly averaging {max_diff(gs_mean, gs_big):.2e}")
assert max_diff(gs_sum, gs_big) < 1e-4 and max_diff(gs_mean, gs_big) > 1.0

# 2) Unequal micro-batches: 48 + 16 examples
uneq = [(X_acc[:48], Y_acc[:48]), (X_acc[48:], Y_acc[48:])]
g_list = [jax.grad(loss_fn)(acc_params, xb, yb) for xb, yb in uneq]
g_naive = jax.tree.map(lambda a, b: (a + b) / 2, *g_list)
g_weighted = jax.tree.map(lambda a, b: (48 * a + 16 * b) / 64, *g_list)
print(f"Unequal micro-batches: plain average {max_diff(g_naive, g_big):.2e}  |  weighted by example count {max_diff(g_weighted, g_big):.2e}")
assert max_diff(g_weighted, g_big) < 1e-6 and max_diff(g_naive, g_big) > 1e-3

# 3) BatchNorm: the gradients and running statistics of one large batch and 4 micro-batches both differ
class BNReg(nnx.Module):
    def __init__(self, *, rngs):
        self.l = nnx.Linear(16, 8, rngs=rngs); self.bn = nnx.BatchNorm(8, rngs=rngs); self.o = nnx.Linear(8, 1, rngs=rngs)
    def __call__(self, x):
        return self.o(nnx.relu(self.bn(self.l(x))))

@nnx.jit
def bn_grad(model, x, y):
    return nnx.grad(lambda m: jnp.mean((m(x) - y) ** 2))(model)

m_big, m_micro = BNReg(rngs=nnx.Rngs(0)), BNReg(rngs=nnx.Rngs(0))
gb = bn_grad(m_big, X_acc, Y_acc)
gm = jax.tree.map(lambda *g: sum(g) / N_ACC, *[bn_grad(m_micro, xb, yb) for xb, yb in mbs])
d_grad = max_diff(gb, gm)
d_stat = max_diff(m_big.bn.mean[...], m_micro.bn.mean[...])
print(f"BatchNorm: gradient difference {d_grad:.3f} (gradient magnitude {float(jnp.abs(gb['l']['kernel'][...]).max()):.3f}), running-mean difference {d_stat:.4f}")
assert d_grad > 1e-2 and d_stat > 1e-3
Summed loss: adding 2.29e-05  |  wrongly averaging 6.66e+01
Unequal micro-batches: plain average 2.89e-01  |  weighted by example count 3.58e-07
BatchNorm: gradient difference 0.161 (gradient magnitude 0.253), running-mean difference 0.0086

6.6 Gradient checkpointing (rematerialization)

Standard backprop stores all the intermediate activations from the forward pass for use in the backward pass. When a model is both deep and wide, these activations can exceed device memory. Gradient checkpointing (also called activation checkpointing, or rematerialization) works like this: during the forward pass, keep only some of the activations; when the backward pass needs an activation that wasn't kept, start from the nearest saved activation and recompute it.

  • Benefit: much lower memory;
  • Cost: more compute per step, because the backward pass reruns part of the forward pass.

How it works

  • Standard forward pass: input → block 1 → activation 1 → block 2 → activation 2 → … → block N → output, with all activations stored.
  • Checkpointed forward pass: store only selected activations (say activation 2, activation K and the output) and drop the rest.
  • Backward pass: to compute the gradient of block 3, use the stored activation 2 directly. Computing the gradients of blocks 1 and 2 needs activation 1, which wasn't stored, so block 1 is recomputed starting from an earlier checkpoint (the input). Once you specify the checkpoint boundaries, JAX handles the recomputation details automatically.

jax.checkpoint (jax.remat)

jax.checkpoint has the alias jax.remat (re-materialization). You can use it as a function wrapper or as a decorator. Intermediate results inside the wrapped function aren't saved for the backward pass. Autodiff knows these values weren't saved and recomputes them automatically during the backward pass. The course's example:

def compute_intensive_block(x, params):
    """A compute block standing in for several layers or a complex operation."""
    x = jnp.dot(x, params['w1']) + params['b1']
    x = jax.nn.relu(x)
    x = jnp.dot(x, params['w2']) + params['b2']
    return x

checkpointed_block = jax.checkpoint(compute_intensive_block)    # activations inside the block aren't stored

def model_fn(x, all_params):
    intermediate_output = x                                       # stands in for the output of earlier layers
    x = checkpointed_block(intermediate_output, all_params['block_params'])
    return x                                                      # stands in for the final output

grad_fn = jax.grad(lambda p, data: jnp.sum(model_fn(data, p)))
# ❌ Wrong: w1 and w2 are generated from the same key; both have the same number of elements (65536), so flattened they are exactly the same random sequence
key = jax.random.PRNGKey(0)
try:
    w1 = jax.random.normal(key, (128, 512))
    w2 = jax.random.normal(key, (512, 128))
    same = bool(jnp.all(w1.ravel() == w2.ravel()))
    print("Measured: are w1.ravel() and w2.ravel() identical?", same)
except Exception as e:
    print("Measured error:", type(e).__name__, e)
assert same
Measured: are w1.ravel() and w2.ravel() identical? True
# ✅ Recommended: give each parameter its own key from split
k1, k2 = jax.random.split(jax.random.key(0))
dummy_x = jnp.ones((1, 128))
dummy_params = {'block_params': {
    'w1': jax.random.normal(k1, (128, 512)), 'b1': jnp.zeros(512),
    'w2': jax.random.normal(k2, (512, 128)), 'b2': jnp.zeros(128)}}
assert not bool(jnp.all(dummy_params['block_params']['w1'].ravel() == dummy_params['block_params']['w2'].ravel()))

gradients = jax.jit(grad_fn)(dummy_params, dummy_x)
# Same result as without checkpointing
plain_grad = jax.jit(jax.grad(lambda p, d: jnp.sum(compute_intensive_block(d, p['block_params']))))(dummy_params, dummy_x)
assert max_diff(gradients, plain_grad) < 1e-4
print("Gradient computed successfully and matches the gradient without checkpointing ✓")

# jax.remat is an alias of jax.checkpoint (a thin wrapper, not the same object)
print("jax.remat is jax.checkpoint:", jax.remat is jax.checkpoint, "| docstring:", jax.remat.__doc__.strip())
assert max_diff(jax.grad(lambda p, d: jnp.sum(jax.remat(compute_intensive_block)(d, p['block_params'])))(dummy_params, dummy_x),
                plain_grad) < 1e-4
Gradient computed successfully and matches the gradient without checkpointing ✓
jax.remat is jax.checkpoint: False | docstring: Alias of :func:`jax.checkpoint`.

Using checkpoints strategically

  • Granularity: works best on large, compute-heavy parts. Checkpointing a tiny operation like an addition costs relatively too much recomputation for what it saves.
  • Logical blocks: a common choice is to checkpoint each Transformer layer or each large convolution block.
  • Find the bottleneck first: use profiling tools to find the parts that use the most activation memory.
  • Works with jit: jax.checkpoint and jax.jit compose smoothly.
  • When to consider it: activations cause OOM; you want to train a deeper or wider model; you can accept a longer step time (the course's example is "20–40% longer, depending on the model").

⚠️ Note: It's easy to assume that "checkpointing the whole model saves the most memory, but nearly doubles compute time". In JAX 0.12 neither half is true (see the two cells below):

  1. Memory: if you wrap the whole model in one jax.checkpoint, the backward pass starts by recomputing the entire forward pass, and that recomputation still saves every activation. So peak temporary memory is exactly the same as without checkpointing (83.25 MiB vs 83.25 MiB). What really saves memory is per-layer checkpointing (down to 37 MiB).
  2. Compute: the backward pass itself is about 2× the forward pass, so the total is about $3F$. Recomputation adds at most one more forward pass $F$, i.e. $3F\to 4F$, an increase of about 33%, not "double". The cell below counts FLOPs with cost_analysis: per-layer checkpointing adds only about 17%, because JAX recomputes only what the backward pass actually needs (the output of a block's last matmul doesn't need to be recomputed). With one checkpoint around the whole model and only jax.grad, FLOPs don't increase at all: grad doesn't use the output of the original forward pass, so that forward pass is removed as dead code and only the recomputation in the backward pass remains. Switch to jax.value_and_grad (common in training, since it returns the loss) and the original forward pass has to stay, so FLOPs go up by about 33%. Either way it's nowhere near "double", and it's in the same range as the "20–40% longer" above. How much wall-clock time actually goes up depends on the hardware; this machine's CPU timings are in the output below, for reference only.
L_W, B_W, D_W, H_W = 16, 256, 256, 1024
def res_block(x, w1, w2):
    return x + jax.nn.gelu(x @ w1) @ w2

def unrolled_loss(p, x, block=res_block):
    for i in range(L_W):
        x = block(x, p["w1"][i], p["w2"][i])
    return jnp.mean(x ** 2)

r = np.random.default_rng(0)
wp = {"w1": jnp.asarray(r.normal(size=(L_W, D_W, H_W)) * 0.02, jnp.float32),
      "w2": jnp.asarray(r.normal(size=(L_W, H_W, D_W)) * 0.02, jnp.float32)}
wx = jnp.asarray(r.normal(size=(B_W, D_W)), jnp.float32)

variants = {
    "no checkpoint":           unrolled_loss,
    "whole-model checkpoint":  jax.checkpoint(unrolled_loss),
    "per-layer checkpoint":    partial(unrolled_loss, block=jax.checkpoint(res_block)),
}
temp = {}
for name, f in variants.items():
    temp[name] = mem_of(jax.grad(f), wp, wx).temp_size_in_bytes / MiB
    print(f"{name:14s} temp memory {temp[name]:7.2f} MiB")
assert temp["whole-model checkpoint"] >= 0.95 * temp["no checkpoint"]       # almost no savings
assert temp["per-layer checkpoint"] < 0.6 * temp["no checkpoint"]           # only per-layer really saves memory
no checkpoint  temp memory   83.25 MiB
whole-model checkpoint temp memory   83.25 MiB
per-layer checkpoint temp memory   37.00 MiB
# Compute: count FLOPs with XLA's cost_analysis
def flops(f, *args):
    return jax.jit(f).lower(*args).compile().cost_analysis()["flops"]

F_fwd = flops(unrolled_loss, wp, wx)
F_bwd_plain = flops(jax.grad(unrolled_loss), wp, wx)
F_bwd_layer = flops(jax.grad(variants["per-layer checkpoint"]), wp, wx)
F_bwd_whole = flops(jax.grad(variants["whole-model checkpoint"]), wp, wx)
print(f"forward                    {F_fwd / 1e9:6.2f} GFLOP")
print(f"forward+backward           {F_bwd_plain / 1e9:6.2f} GFLOP  (= {F_bwd_plain / F_fwd:.2f} × forward)")
print(f"per-layer checkpoint       {F_bwd_layer / 1e9:6.2f} GFLOP  ({F_bwd_layer / F_bwd_plain - 1:.1%} more than no checkpoint)")
print(f"whole-model checkpoint     {F_bwd_whole / 1e9:6.2f} GFLOP  ({F_bwd_whole / F_bwd_plain - 1:.1%} more than no checkpoint)")
assert 2.8 < F_bwd_plain / F_fwd < 3.2                 # backward ≈ 2 × forward
assert F_bwd_layer / F_bwd_plain < 4 / 3 + 1e-3        # at most +33%, nowhere near "double"
assert F_bwd_whole / F_bwd_plain < 4 / 3 + 1e-3
# One checkpoint around the whole model: with grad alone, the original forward pass is removed as dead code; value_and_grad returns the loss, so the original forward pass stays
F_vg_plain = flops(jax.value_and_grad(unrolled_loss), wp, wx)
F_vg_whole = flops(jax.value_and_grad(variants["whole-model checkpoint"]), wp, wx)
print(f"whole-model checkpoint + value_and_grad: {F_vg_whole / F_vg_plain - 1:.1%} more than no checkpoint (the original forward pass can't be removed)")
assert F_bwd_whole <= 1.01 * F_bwd_plain
assert 1.2 < F_vg_whole / F_vg_plain < 1.4

for name, f in variants.items():
    print(f"{name:14s} measured time {bench(jax.jit(jax.grad(f)), wp, wx):7.2f} ms")
forward                      4.33 GFLOP
forward+backward            12.98 GFLOP  (= 3.00 × forward)
per-layer checkpoint        15.16 GFLOP  (16.8% more than no checkpoint)
whole-model checkpoint      12.98 GFLOP  (0.0% more than no checkpoint)
whole-model checkpoint + value_and_grad: 33.4% more than no checkpoint (the original forward pass can't be removed)
no checkpoint  measured time   23.90 ms
whole-model checkpoint measured time   31.04 ms
per-layer checkpoint measured time   28.46 ms

🔍 Extra: checkpoint policies

⚪ Optional: when you need a fine-grained trade-off between memory and recomputation

Without a policy, jax.checkpoint saves only the inputs of the wrapped function and recomputes everything inside. The policy argument gives you fine control over "which intermediate results are saved and which are recomputed". The policies available in jax.checkpoint_policies in JAX 0.12 are:

PolicyMeaning
everything_saveablesave everything (same as no checkpointing)
nothing_saveablerecompute everything (same as the default behavior)
dots_saveable (alias checkpoint_dots)save only matmul (dot_general) results and recompute elementwise ops (matmuls are expensive, elementwise ops are cheap)
dots_with_no_batch_dims_saveable (alias checkpoint_dots_with_no_batch_dims)save only matmul results without batch dimensions (batched matmuls in attention aren't saved)
save_only_these_names(*names)save only intermediate values tagged with checkpoint_name(x, "name")
save_anything_except_these_names(*names)save everything except these names
save_any_names_but_these(*names)save every named value except these
save_from_both_policies(p1, p2)the union of two policies
offload_dot_with_no_batch_dims(src, dst), save_and_offload_only_these_names(...)offload residuals to host memory (meaningful on GPU/TPU)

Two other arguments:

  • prevent_cse=True (the default): stops XLA's common subexpression elimination from optimizing the "recompute" back into a "reuse", which would make the checkpoint pointless. Inside lax.scan you can set it to False (scan already blocks CSE), which may be slightly faster.
  • static_argnums: treat some arguments as static values (for example a Python boolean that controls a branch).

Tools: from jax.ad_checkpoint import checkpoint_name, print_saved_residuals. print_saved_residuals(f, *args) prints which residuals the backward pass saves (shapes and where they come from), and it's available in 0.12. Note that the jax.ad_checkpoint submodule has to be imported explicitly: if you only import jax and write jax.ad_checkpoint.xxx, you get an AttributeError.

Below we compare the policies on a 16-layer residual MLP stacked with lax.scan.

from jax.ad_checkpoint import checkpoint_name, print_saved_residuals
cp = jax.checkpoint_policies
print("Policies available in 0.12:", [n for n in dir(cp) if not n.startswith("_")])

def named_block(x, w1, w2):
    h = checkpoint_name(x @ w1, "pre_act")       # give the output of the first matmul a name
    return x + jax.nn.gelu(h) @ w2

def make_scan_loss(policy=None, remat=True):
    blk = jax.checkpoint(named_block, policy=policy) if remat else named_block
    def loss(params, x):
        def body(x, p):
            return blk(x, p["w1"], p["w2"]), None
        x, _ = lax.scan(body, x, params)          # parameters are stacked along dim 0; scan walks through the layers
        return jnp.mean(x ** 2)
    return loss

policies = {
    "no checkpoint":                   make_scan_loss(remat=False),
    "everything_saveable":             make_scan_loss(cp.everything_saveable),
    "dots_saveable":                   make_scan_loss(cp.dots_saveable),
    "save_only_these_names('pre_act')": make_scan_loss(cp.save_only_these_names("pre_act")),
    "nothing_saveable":                make_scan_loss(cp.nothing_saveable),
    "default (policy=None)":           make_scan_loss(None),
}

g_ref = jax.jit(jax.grad(policies["no checkpoint"]))(wp, wx)
rows = []
for name, f in policies.items():
    gf = jax.grad(f)
    ma = mem_of(gf, wp, wx)
    g = jax.jit(gf)(wp, wx)
    ms_ = bench(jax.jit(gf), wp, wx)
    rows.append((name, ma.temp_size_in_bytes / MiB, ms_, max_diff(g, g_ref)))
base_temp = rows[0][1]
print(f"\n{'policy':34s} {'temp MiB':>10s} {'relative':>7s} {'time ms':>8s} {'grad err':>9s}")
for name, t, ms_, err in rows:
    print(f"{name:34s} {t:10.2f} {t / base_temp:7.1%} {ms_:8.2f} {err:9.1e}")
    assert err < 1e-5                                       # every policy gives the same gradient as no checkpointing
mem = {n: t for n, t, _, _ in rows}
assert mem["nothing_saveable"] == mem["default (policy=None)"]
assert mem["nothing_saveable"] < mem["dots_saveable"] < mem["no checkpoint"]
assert mem["save_only_these_names('pre_act')"] == mem["dots_saveable"]   # in this block both save the same things
assert cp.checkpoint_dots is cp.dots_saveable and cp.checkpoint_dots_with_no_batch_dims is cp.dots_with_no_batch_dims_saveable   # aliases
Policies available in 0.12: ['checkpoint_dots', 'checkpoint_dots_with_no_batch_dims', 'dots_saveable', 'dots_with_no_batch_dims_saveable', 'everything_saveable', 'nothing_saveable', 'offload_dot_with_no_batch_dims', 'save_and_offload_only_these_names', 'save_any_names_but_these', 'save_anything_except_these_names', 'save_from_both_policies', 'save_only_these_names']
policy                               temp MiB relative  time ms  grad err
no checkpoint                           88.50  100.0%    32.60   0.0e+00
everything_saveable                     92.50  104.5%    31.14   0.0e+00
dots_saveable                           25.50   28.8%    26.49   0.0e+00
save_only_these_names('pre_act')        25.50   28.8%    26.35   0.0e+00
nothing_saveable                         9.50   10.7%    30.00   0.0e+00
default (policy=None)                    9.50   10.7%    29.75   0.0e+00
print("—— no checkpoint: residuals each layer saves for the backward pass ——")
print_saved_residuals(policies["no checkpoint"], wp, wx)
print("\n—— dots_saveable: only matmul results are saved ——")
print_saved_residuals(policies["dots_saveable"], wp, wx)
print("\n—— nothing_saveable: only each layer's input is saved ——")
print_saved_residuals(policies["nothing_saveable"], wp, wx)
—— no checkpoint: residuals each layer saves for the backward pass ——
f32[16,256,1024] from the argument params['w1']
f32[16,1024,256] from the argument params['w2']
f32[16,256,256] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[256,256] output of mul from /tmp/ipykernel_66671/1797682936.py:15:24 (make_scan_loss.<locals>.loss)

—— dots_saveable: only matmul results are saved ——
f32[16,256,1024] from the argument params['w1']
f32[16,1024,256] from the argument params['w2']
f32[16,256,1024] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[16,256,256] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[256,256] output of mul from /tmp/ipykernel_66671/1797682936.py:15:24 (make_scan_loss.<locals>.loss)

—— nothing_saveable: only each layer's input is saved ——
f32[16,256,1024] from the argument params['w1']
f32[16,1024,256] from the argument params['w2']
f32[16,256,256] output of scan from /tmp/ipykernel_66671/1797682936.py:14:15 (make_scan_loss.<locals>.loss)
f32[256,256] output of mul from /tmp/ipykernel_66671/1797682936.py:15:24 (make_scan_loss.<locals>.loss)

Reading the output:

  • With no checkpoint, each layer stores 5 arrays of shape [16,256,1024]: the result of x @ w1, plus several intermediates of the tanh-approximated GELU ($x^3$, the output of tanh, $1-\tanh^2$, and so on). The width-$H$ intermediates are the bulk of it.
  • dots_saveable stores only x @ w1 (one [16,256,1024]) and each layer's input [16,256,256]; GELU is recomputed from x @ w1 in the backward pass. That adds almost no compute but drops 4/5 of the large arrays.
  • nothing_saveable stores only each layer's input (the scan carry) and recomputes all the width-$H$ intermediates.

Note that print_saved_residuals reflects residuals at the jaxpr level. XLA also does fusion when compiling, so the temporary memory reported by memory_analysis() is the final number. The two follow the same trend but the values differ. Also, on CPU, peak_memory_in_bytes from memory_analysis() is roughly "parameters + outputs" (off by only a few dozen bytes) and doesn't change with the policy. That's why this chapter always uses temp_size_in_bytes (temporary buffers) to measure activation memory.

ma0 = mem_of(jax.grad(policies["no checkpoint"]), wp, wx)
ma1 = mem_of(jax.grad(policies["nothing_saveable"]), wp, wx)
print("peak_memory_in_bytes:", ma0.peak_memory_in_bytes, ma1.peak_memory_in_bytes,
      "| argument+output:", ma0.argument_size_in_bytes + ma0.output_size_in_bytes)
print("temp_size_in_bytes  :", ma0.temp_size_in_bytes, ma1.temp_size_in_bytes)
assert ma0.peak_memory_in_bytes == ma1.peak_memory_in_bytes          # on CPU the peak field doesn't reflect temporary memory
assert ma1.temp_size_in_bytes < ma0.temp_size_in_bytes
peak_memory_in_bytes: 67371048 67371048 | argument+output: 67371024
temp_size_in_bytes  : 92799300 9961668

✅ Checkpointing in NNX: nnx.remat

nnx.remat is NNX's lifted version of jax.checkpoint. It can wrap "a function that takes a module" directly, and its policy, prevent_cse and static_argnums arguments work the same as in jax.checkpoint. A common pattern is to wrap the function that "calls one layer" and use it in the model's loop.

class MLPBlock(nnx.Module):
    def __init__(self, d, h, *, rngs):
        self.up = nnx.Linear(d, h, rngs=rngs)
        self.down = nnx.Linear(h, d, rngs=rngs)
    def __call__(self, x):
        return x + self.down(jax.nn.gelu(self.up(x)))

def call_block(block, x):
    return block(x)
call_block_remat = nnx.remat(call_block)                         # default: save only the inputs
call_block_dots = nnx.remat(call_block, policy=cp.dots_saveable)  # you can pass a policy too

class DeepMLP(nnx.Module):
    def __init__(self, n_layers, d, h, *, remat=None, rngs):
        self.blocks = nnx.List([MLPBlock(d, h, rngs=rngs) for _ in range(n_layers)])
        self.remat = remat
    def __call__(self, x):
        call = {None: call_block, "full": call_block_remat, "dots": call_block_dots}[self.remat]
        for blk in self.blocks:
            x = call(blk, x)
        return x

def nnx_loss(model, x):
    return jnp.mean(model(x) ** 2)

x_nnx = jnp.asarray(np.random.default_rng(1).normal(size=(256, 256)), jnp.float32)
res = {}
for mode in [None, "dots", "full"]:
    m = DeepMLP(8, 256, 1024, remat=mode, rngs=nnx.Rngs(0))
    gdef, st = nnx.split(m)
    # Use split/merge to get a pure function, then hand it to jax.jit(...).lower().compile().memory_analysis()
    pure_grad = lambda st, x: jax.grad(lambda s: nnx_loss(nnx.merge(gdef, s), x))(st)
    temp_mb = mem_of(pure_grad, st, x_nnx).temp_size_in_bytes / MiB
    g = nnx.jit(nnx.grad(nnx_loss))(m, x_nnx)
    res[mode] = (temp_mb, g)
    print(f"remat={str(mode):5s} temp memory {temp_mb:6.2f} MiB")
assert max_diff(res[None][1], res["full"][1]) < 1e-6 and max_diff(res[None][1], res["dots"][1]) < 1e-6
assert res["full"][0] <= res["dots"][0] < res[None][0]
print(f"nnx.remat gradients match no checkpointing ✓; full recomputation saves {1 - res['full'][0] / res[None][0]:.1%} of temp memory")
remat=None  temp memory  12.25 MiB
remat=dots  temp memory   2.25 MiB
remat=full  temp memory   2.25 MiB
nnx.remat gradients match no checkpointing ✓; full recomputation saves 81.6% of temp memory

🔍 Extra: stacking layers with lax.scan + checkpointing is the standard structure for large models

🟢 Must-read

  • If you unroll $L$ layers with a plain Python for loop, compile time grows linearly with $L$. Stack the parameters along dim 0 and use lax.scan (in NNX, create stacked layers with nnx.scan or nnx.vmap) and only one layer gets compiled, so compile time doesn't depend on the number of layers. Reference implementations such as MaxText all do it this way.
  • With scan + jax.checkpoint (the make_scan_loss pattern above), scan saves only the carry (the layer input) for each layer, which is exactly "one checkpoint per layer". Activation memory drops from $O(L\cdot\text{per-layer activations})$ to $O(L\cdot\text{layer input}+\text{one layer's activations})$.
  • Going further, the $\sqrt{L}$ strategy: checkpoint every $k$ layers, and you get $O(\sqrt L)$ memory for the price of just one extra forward pass (Chen et al. 2016, Training Deep Nets with Sublinear Memory Cost). See Exercise 2 at the end.

6.7 Mixed precision training

Mixed precision training uses 16-bit floats for part of the computation: activations take roughly half the memory, it's noticeably faster on supported hardware, and the accuracy loss is usually small. The approach: the expensive matmuls and convolutions in the forward and backward passes run in 16 bits, while the master copy of the model weights (and usually the optimizer state) stays in float32 for numerical stability.

The two 16-bit formats

  1. jnp.float16 (IEEE 754 half precision): 1 sign bit, 5 exponent bits, 10 mantissa bits.
    • Pros: clearly faster on hardware with fp16 support (such as NVIDIA Tensor Cores), half the memory of fp32.
    • Cons: very narrow dynamic range. Small gradients easily underflow to 0 and large values easily overflow to inf / NaN, so you usually need loss scaling.
  2. jnp.bfloat16 (Brain Float, from Google Brain): 1 sign bit, 8 exponent bits (same as fp32), 7 mantissa bits.
    • Pros: almost the same dynamic range as fp32, so it rarely overflows or underflows and usually doesn't need loss scaling; saves as much memory as fp16; faster on TPUs and newer GPUs.
    • Cons: even lower precision than fp16 (fewer mantissa bits), which can affect convergence for especially sensitive models.
Propertyfloat32float16bfloat16
Total bits321616
Exponent bits858
Mantissa bits23107
Dynamic rangewidenarrowwide (close to fp32)
Precisionhighmediumlow
Needs loss scaling?nousuallyusually not

The course's conclusion: when the hardware supports it (TPUs, NVIDIA Ampere and later GPUs), prefer bf16. If only fp16 is supported efficiently, you have to pair it with loss scaling and implement it carefully.

Below we verify this table with jnp.finfo.

print(f"{'dtype':9s} {'bits':>4s} {'exp':>4s} {'mant':>4s} {'eps':>10s} {'max':>11s} {'min normal':>11s} {'min subnormal':>11s}")
info = {}
for dt in [jnp.float32, jnp.float16, jnp.bfloat16]:
    fi = jnp.finfo(dt)
    info[fi.dtype.name] = fi
    print(f"{fi.dtype.name:9s} {fi.bits:4d} {fi.nexp:4d} {fi.nmant:4d} {float(fi.eps):10.3g} {float(fi.max):11.4g} "
          f"{float(fi.tiny):11.4g} {float(fi.smallest_subnormal):11.4g}")
assert (info["float16"].nexp, info["float16"].nmant) == (5, 10)
assert (info["bfloat16"].nexp, info["bfloat16"].nmant) == (8, 7)
assert info["bfloat16"].nexp == info["float32"].nexp                # same exponent bits as fp32 → same dynamic range
assert float(info["float16"].max) == 65504.0
assert float(info["bfloat16"].eps) > float(info["float16"].eps)     # bf16 has lower precision

# Underflow / overflow
print("\n1e-8  →  fp16:", jnp.array(1e-8, jnp.float16), "| bf16:", jnp.array(1e-8, jnp.bfloat16))
print("7e4   →  fp16:", jnp.array(7e4, jnp.float32).astype(jnp.float16), "| bf16:", jnp.array(7e4, jnp.bfloat16))
assert float(jnp.array(1e-8, jnp.float16)) == 0.0 and float(jnp.array(1e-8, jnp.bfloat16)) > 0
assert jnp.isinf(jnp.array(7e4, jnp.float32).astype(jnp.float16))
# Precision: bf16 has only 8 significant binary digits, so 256 + 1 == 256
print("256 + 1 →  bf16:", jnp.bfloat16(256) + jnp.bfloat16(1), "| fp16:", jnp.float16(256) + jnp.float16(1))
assert float(jnp.bfloat16(256) + jnp.bfloat16(1)) == 256.0 and float(jnp.float16(256) + jnp.float16(1)) == 257.0
dtype     bits  exp mant        eps         max  min normal min subnormal
float32     32    8   23   1.19e-07   3.403e+38   1.175e-38   1.401e-45
float16     16    5   10   0.000977    6.55e+04   6.104e-05    5.96e-08
bfloat16    16    8    7    0.00781    3.39e+38   1.175e-38   9.184e-41

1e-8  →  fp16: 0.0 | bf16: 1.00117e-08
7e4   →  fp16: inf | bf16: 70144
256 + 1 →  bf16: 256 | fp16: 257.0

Implementing it in JAX: fp32 parameters, bf16 compute

The course first shows the library approach: Flax lets you set the parameter dtype (param_dtype) and the compute dtype (dtype) separately. With param_dtype=jnp.float32 and dtype=jnp.bfloat16, parameters are initialized and stored in fp32, and during the forward pass the layer first casts parameters and inputs to bf16 before computing. The expected output: bf16 input, fp32 parameters, bf16 output.

# ❌ Old style (outdated): flax.linen is the old API (the NNX version is in the next cell)
try:
    import flax.linen as nn

    class SimpleDense(nn.Module):
        features: int
        param_dtype: jnp.dtype = jnp.float32     # master weights in fp32
        dtype: jnp.dtype = jnp.bfloat16          # compute in bf16

        @nn.compact
        def __call__(self, x):
            x = x.astype(self.dtype)
            y = nn.Dense(features=self.features, param_dtype=self.param_dtype,
                         dtype=self.dtype, name='dense_layer')(x)
            return nn.relu(y)

    key = jax.random.PRNGKey(0)
    dummy_input = jnp.zeros((1, 10), dtype=jnp.bfloat16)
    lin_mp = SimpleDense(features=5)
    lin_mp_params = lin_mp.init(key, dummy_input)['params']
    output = lin_mp.apply({'params': lin_mp_params}, dummy_input)
    print(f"Input dtype: {dummy_input.dtype}")
    print(f"Parameter dtype (example): {jax.tree.leaves(lin_mp_params)[0].dtype}")
    print(f"Output dtype: {output.dtype}")
    lin_grads = jax.grad(lambda p: lin_mp.apply({'params': p}, jnp.ones((1, 10))).astype(jnp.float32).sum())(lin_mp_params)
    print("Gradient dtype:", jax.tree.map(lambda g: g.dtype.name, lin_grads))
except Exception as e:
    print("Measured error:", type(e).__name__, e)
assert output.dtype == jnp.bfloat16 and jax.tree.leaves(lin_mp_params)[0].dtype == jnp.float32
Input dtype: bfloat16
Parameter dtype (example): float32
Output dtype: bfloat16
Gradient dtype: {'dense_layer': {'bias': 'float32', 'kernel': 'float32'}}
# ✅ Recommended: NNX's nnx.Linear also has dtype / param_dtype
lin_bf16 = nnx.Linear(10, 5, dtype=jnp.bfloat16, param_dtype=jnp.float32, rngs=nnx.Rngs(0))
x_mp = jnp.ones((4, 10), jnp.float32)
y_mp = lin_bf16(x_mp)
print("Parameter dtype:", lin_bf16.kernel[...].dtype, "| output dtype:", y_mp.dtype)
assert lin_bf16.kernel[...].dtype == jnp.float32 and y_mp.dtype == jnp.bfloat16

# Key point: differentiating w.r.t. fp32 parameters gives fp32 gradients (the transpose of the cast turns the bf16 cotangent back into fp32)
grads_mp = nnx.grad(lambda m: m(x_mp).astype(jnp.float32).sum())(lin_bf16)
print("Gradient dtype:", {k: v.dtype.name for k, v in nnx.to_pure_dict(grads_mp).items()})
assert all(g.dtype == jnp.float32 for g in jax.tree.leaves(grads_mp))
assert all(g.dtype == jnp.float32 for g in jax.tree.leaves(lin_grads))      # same for Linen
Parameter dtype: float32 | output dtype: bfloat16
Gradient dtype: {'bias': 'float32', 'kernel': 'float32'}

⚠️ Note: It's easy to assume the parameter gradients are also in the compute dtype (bfloat16 here), and easy to assume mixed precision cuts "the memory for activations and gradients roughly in half". In reality (JAX 0.12 + Flax 0.12.12, same for Linen and NNX), the gradient with respect to fp32 parameters is float32, because the transpose of astype(bf16) turns the cotangent back into the parameter's dtype. The cotangents flowing through backprop (the gradients of the activations) are bf16, so what you save is the memory for activations and activation gradients. Parameter gradients and optimizer state are still fp32.

Manual casting

Without a library you have to manage it yourself: cast the fp32 weights to bf16, do the matmuls and activations in bf16, get a bf16 output, and use the resulting gradients to update the fp32 master parameters. The course warns that this takes more care to make sure the types are right throughout the model and the training loop.

def predict_bf16(params_f32, inputs_bf16):
    """Manual casting example: params_f32 is a list [(W, b), ...]."""
    activations = inputs_bf16
    for W_f32, b_f32 in params_f32:
        W_bf16, b_bf16 = W_f32.astype(jnp.bfloat16), b_f32.astype(jnp.bfloat16)
        outputs = jnp.dot(activations, W_bf16) + b_bf16
        activations = jax.nn.relu(outputs)
    return activations                                          # the output is bf16

r = np.random.default_rng(0)
mp_params = [(jnp.asarray(r.normal(size=(64, 128)) / 8, jnp.float32), jnp.zeros(128)),
             (jnp.asarray(r.normal(size=(128, 32)) / 11, jnp.float32), jnp.zeros(32))]
x_in = jnp.asarray(r.normal(size=(16, 64)), jnp.float32)
out_bf16 = jax.jit(predict_bf16)(mp_params, x_in.astype(jnp.bfloat16))
print("Output dtype:", out_bf16.dtype)
assert out_bf16.dtype == jnp.bfloat16
Output dtype: bfloat16

🔍 Extra: preferred_element_type, or accumulating matmuls in fp32

⚪ Optional: when low-precision matmuls aren't accurate enough, or you want to understand what the accumulator dtype does

Each output element of a matmul is a sum of $K$ products. If the accumulator is also bf16, every addition is rounded to 8 significant binary digits (about 2–3 significant decimal digits), and the error grows with $K$. jnp.dot(a, b, preferred_element_type=jnp.float32) tells XLA: the inputs are bf16, but use fp32 for the output (the accumulator). On GPU / TPU that's how Tensor Cores / MXUs work natively anyway (multiply in bf16, add in fp32), so it adds almost no overhead. Flax's nnx.Linear also has a preferred_element_type argument.

Below we compare, on CPU, the relative error of several precision combinations against a float64 reference. Note: on this machine, jnp.dot with bf16 output is element-for-element identical to "accumulate in fp32, then round to bf16", which means XLA:CPU already accumulates in fp32 internally (there's an assert for this in the cell below). So on CPU the difference in error between those two rows comes only from the final rounding of the output, and most of the error comes from rounding the inputs to bf16. You only see the effect of accumulator precision when you really accumulate one value at a time in low precision; see the sequential-sum example at the end of the cell.

a_np = np.random.default_rng(0).normal(size=(256, 1024))
b_np = np.random.default_rng(1).normal(size=(1024, 256))
ref = a_np @ b_np                                    # float64 reference
a32, b32 = jnp.asarray(a_np, jnp.float32), jnp.asarray(b_np, jnp.float32)

matmuls = {
    "fp32":                            lambda a, b: jnp.dot(a, b),
    "bf16 in, bf16 out":               lambda a, b: jnp.dot(a.astype(jnp.bfloat16), b.astype(jnp.bfloat16)),
    "bf16 in, fp32 accum (preferred)": lambda a, b: jnp.dot(a.astype(jnp.bfloat16), b.astype(jnp.bfloat16),
                                                         preferred_element_type=jnp.float32),
    "fp16 in, fp16 out":               lambda a, b: jnp.dot(a.astype(jnp.float16), b.astype(jnp.float16)),
}
err = {}
for name, f in matmuls.items():
    out = jax.jit(f)(a32, b32)                      # data is passed in as arguments
    err[name] = np.abs(np.asarray(out, np.float64) - ref).max() / np.abs(ref).max()
    print(f"{name:32s} output {str(out.dtype):9s} relative error {err[name]:.2e}")
assert err["fp32"] < 1e-5
assert err["bf16 in, fp32 accum (preferred)"] < err["bf16 in, bf16 out"]
assert err["fp16 in, fp16 out"] < err["bf16 in, bf16 out"]          # fp16 has more mantissa bits, so it's more precise than bf16

hlo = jax.jit(matmuls["bf16 in, fp32 accum (preferred)"]).lower(a32, b32).as_text()
dot_line = [l.strip() for l in hlo.splitlines() if "dot_general" in l][0]
print("\nStableHLO:", dot_line[:160])
assert "xbf16>) -> tensor<256x256xf32>" in dot_line

# Measured on CPU: a dot with bf16 output is element-for-element identical to "accumulate in fp32, then round to bf16" → XLA:CPU accumulates in fp32 internally
a16, b16 = a32.astype(jnp.bfloat16), b32.astype(jnp.bfloat16)
out_bf16_dot = jax.jit(lambda a, b: jnp.dot(a, b))(a16, b16)
out_acc32_dot = jax.jit(lambda a, b: jnp.dot(a, b, preferred_element_type=jnp.float32))(a16, b16)
assert bool(jnp.all(out_bf16_dot == out_acc32_dot.astype(jnp.bfloat16)))

# The danger of accumulating one at a time: sequentially adding 1000 ones in bf16 gets stuck at 256
seq_sum = jax.jit(lambda x: lax.fori_loop(0, x.shape[0], lambda i, s: s + x[i], jnp.zeros((), x.dtype)))
print("\nSequentially adding 1000 ones: bf16 =", seq_sum(jnp.ones(1000, jnp.bfloat16)), "| fp32 =", seq_sum(jnp.ones(1000, jnp.float32)))
print("jnp.sum(10000 bf16 ones) =", jnp.sum(jnp.ones(10000, jnp.bfloat16)), "(XLA reduces in fp32 internally, then rounds to bf16)")
assert float(seq_sum(jnp.ones(1000, jnp.bfloat16))) == 256.0
fp32                             output float32   relative error 1.26e-06
bf16 in, bf16 out                output bfloat16  relative error 3.42e-03
bf16 in, fp32 accum (preferred)  output float32   relative error 2.36e-03
fp16 in, fp16 out                output float16   relative error 4.34e-04

StableHLO: %2 = stablehlo.dot_general %0, %1, contracting_dims = [1] x [0], precision = [DEFAULT, DEFAULT] : (tensor<256x1024xbf16>, tensor<1024x256xbf16>) -> tensor<256x2
Sequentially adding 1000 ones: bf16 = 256 | fp32 = 1000.0
jnp.sum(10000 bf16 ones) = 9984 (XLA reduces in fp32 internally, then rounds to bf16)
# A real training comparison: fp32 vs "fp32 parameters + bf16 compute", same initialization, same data
def make_mlp_loss(compute_dtype):
    def loss(p, x, y):
        h = x.astype(compute_dtype)
        for i, (w, b) in enumerate(p):
            h = jnp.dot(h, w.astype(compute_dtype), preferred_element_type=jnp.float32) + b
            h = (jnp.tanh(h) if i < len(p) - 1 else h).astype(compute_dtype)
        return jnp.mean((h.astype(jnp.float32) - y) ** 2)       # compute the loss in fp32
    return loss

r = np.random.default_rng(3)
Xmp = jnp.asarray(r.uniform(-2, 2, size=(512, 8)), jnp.float32)
Ymp = jnp.sin(Xmp.sum(axis=1, keepdims=True))
p_init = [(jnp.asarray(r.normal(size=(8, 64)) / np.sqrt(8), jnp.float32), jnp.zeros(64)),
          (jnp.asarray(r.normal(size=(64, 64)) / 8, jnp.float32), jnp.zeros(64)),
          (jnp.asarray(r.normal(size=(64, 1)) / 8, jnp.float32), jnp.zeros(1))]
mp_tx = optax.adam(3e-3)

def train_curve(compute_dtype, steps=400):
    loss = make_mlp_loss(compute_dtype)
    @jax.jit
    def run(p, x, y):
        def step(carry, _):
            p, s = carry
            l, g = jax.value_and_grad(loss)(p, x, y)
            u, s = mp_tx.update(g, s, p)
            return (optax.apply_updates(p, u), s), l
        (p, _), ls = lax.scan(step, (p, mp_tx.init(p)), None, length=steps)
        return p, ls
    return run(p_init, Xmp, Ymp)

p32, l32 = train_curve(jnp.float32)
pbf, lbf = train_curve(jnp.bfloat16)
print(f"fp32         : loss {float(l32[0]):.4f} → {float(l32[-1]):.5f}")
print(f"bf16 compute : loss {float(lbf[0]):.4f} → {float(lbf[-1]):.5f}")
print("Parameter dtype after bf16 training:", p32[0][0].dtype, pbf[0][0].dtype)
assert pbf[0][0].dtype == jnp.float32                    # master weights are always fp32
assert float(l32[-1]) < 0.2 * float(l32[0]) and float(lbf[-1]) < 0.2 * float(lbf[0])
fp32         : loss 0.7060 → 0.00015
bf16 compute : loss 0.7059 → 0.00054
Parameter dtype after bf16 training: float32 float32

Loss scaling for fp16

fp16 has a limited dynamic range, so very small gradients underflow to 0. The course gives four steps:

  1. Scale the loss: before backprop, multiply the loss by a scale factor $S$ (say $S=2^{15}$): $\mathcal L_{\text{scaled}}=\mathcal L\times S$;
  2. Compute gradients: differentiation is linear, so the gradients are also scaled up by $S$: $\nabla_\theta\mathcal L_{\text{scaled}}=S\,\nabla_\theta\mathcal L$, and small gradients that would have underflowed now land in the range fp16 can represent;
  3. Unscale: divide by $S$ before updating the fp32 master weights: $g=\nabla_\theta\mathcal L_{\text{scaled}}/S$;
  4. Update the weights: update the fp32 master weights with the unscaled gradients.

$S$ can be fixed or adjusted dynamically: start with a fairly large $S$; as soon as an overflow shows up (inf / NaN in the gradients), reduce $S$ and skip this update; after several normal steps in a row, increase $S$.

First, the problem itself: when computing in fp16, small gradients underflow during backprop.

def make_fp16_loss(scale, weight=1.0):
    """fp16 compute, fp32 master weights; returns weight × loss × scale."""
    def loss(p, x, y):
        h = x.astype(jnp.float16)
        for i, (w, b) in enumerate(p):
            h = h @ w.astype(jnp.float16) + b.astype(jnp.float16)
            if i < len(p) - 1:
                h = jnp.tanh(h)
        err = h.astype(jnp.float32) - y
        return weight * jnp.mean(err ** 2) * scale
    return loss

def fp32_loss(p, x, y, weight=1.0):
    h = x
    for i, (w, b) in enumerate(p):
        h = h @ w + b
        if i < len(p) - 1:
            h = jnp.tanh(h)
    return weight * jnp.mean((h - y) ** 2)

# Loss weight 1e-4: simulates very small gradients, as when "averaging over a huge batch / long sequence" or "an auxiliary loss with a tiny weight"
W_LOSS = 1e-4
g_true = jax.jit(jax.grad(partial(fp32_loss, weight=W_LOSS)))(p_init, Xmp, Ymp)
flat_true = np.concatenate([np.ravel(t) for t in jax.tree.leaves(g_true)])

def grad_quality(scale):
    g = jax.jit(jax.grad(make_fp16_loss(scale, W_LOSS)))(p_init, Xmp, Ymp)
    flat = np.concatenate([np.ravel(t) for t in jax.tree.leaves(g)]) / scale      # unscale
    zero_frac = np.mean((flat == 0) & (flat_true != 0))
    finite = bool(np.all(np.isfinite(flat)))
    rel = np.linalg.norm(flat - flat_true) / np.linalg.norm(flat_true) if finite else np.inf
    return zero_frac, rel, finite

print(f"Magnitude of the true gradient: max |g| = {np.abs(flat_true).max():.2e} (smallest fp16 normal number {float(jnp.finfo(jnp.float16).tiny):.1e})")
quality = {}
for S in [1.0, 2.0 ** 8, 2.0 ** 15, 2.0 ** 32]:
    zf, rel, fin = grad_quality(S)
    quality[S] = (zf, rel, fin)
    print(f"S = 2^{int(np.log2(S)):2d}: fraction underflowed to 0 {zf:6.1%}, relative error {rel:9.3e}, all finite: {fin}")
assert quality[1.0][0] > 0.05 and quality[1.0][1] > 0.05         # no scaling: lots of underflow, distorted gradient
assert quality[2.0 ** 15][0] == 0 and quality[2.0 ** 15][1] < 0.01 # 2^15: no underflow, the only error left is fp16 forward rounding
assert not quality[2.0 ** 32][2]                                  # too large: overflows to inf
Magnitude of the true gradient: max |g| = 3.79e-05 (smallest fp16 normal number 6.1e-05)
S = 2^ 0: fraction underflowed to 0  21.3%, relative error 1.557e-01, all finite: True
S = 2^ 8: fraction underflowed to 0   1.4%, relative error 8.996e-04, all finite: True
S = 2^15: fraction underflowed to 0   0.0%, relative error 6.778e-04, all finite: True
S = 2^32: fraction underflowed to 0   0.0%, relative error       inf, all finite: False

✅ Hand-written dynamic loss scaling

The state is just two numbers: the current scale factor scale and the number of consecutive good steps good_steps. Each step:

  • All gradients finite → normal update; good_steps += 1, and when it reaches period, scale *= 2 and the counter resets;
  • inf / NaN shows up → skip the update (parameters and optimizer state stay unchanged), scale /= 2, good_steps = 0.

Every branch is written as a data-dependent selection with jnp.where, so the whole training loop can go inside lax.scan.

from typing import NamedTuple

class DynScale(NamedTuple):
    scale: jax.Array
    good_steps: jax.Array

def all_finite(tree):
    return jnp.all(jnp.stack([jnp.all(jnp.isfinite(t)) for t in jax.tree.leaves(tree)]))

def dyn_scale_step(ls: DynScale, finite, period=50, factor=2.0, min_scale=1.0):
    grow = finite & (ls.good_steps + 1 >= period)
    new_scale = jnp.where(finite, jnp.where(grow, ls.scale * factor, ls.scale),
                          jnp.maximum(ls.scale / factor, min_scale))
    new_good = jnp.where(finite & ~grow, ls.good_steps + 1, 0)
    return DynScale(new_scale, new_good)

fp16_tx = mp_tx                    # same optimizer as the fp32 / bf16 training above (Adam 3e-3)

@partial(jax.jit, static_argnames="steps")
def train_fp16_dynamic(p, x, y, ls0, steps):
    def step(carry, _):
        p, s, ls = carry
        scaled_loss = lambda p_: make_fp16_loss(1.0)(p_, x, y) * ls.scale   # 1) scale the loss
        loss_scaled, g = jax.value_and_grad(scaled_loss)(p)                 # 2) the gradients are scaled up too
        g = jax.tree.map(lambda t: t / ls.scale, g)                         # 3) unscale
        finite = all_finite(g)
        u, s_new = fp16_tx.update(jax.tree.map(lambda t: jnp.where(finite, t, 0.0), g), s, p)
        p_new = optax.apply_updates(p, u)
        # Skip on overflow: parameters and optimizer state keep their old values
        p = jax.tree.map(lambda a, b: jnp.where(finite, a, b), p_new, p)
        s = jax.tree.map(lambda a, b: jnp.where(finite, a, b), s_new, s)
        return (p, s, dyn_scale_step(ls, finite)), (ls.scale, finite, loss_scaled / ls.scale)
    (p, s, ls), (scales, finites, losses) = lax.scan(step, (p, fp16_tx.init(p), ls0), None, length=steps)
    return p, ls, scales, finites, losses

ls0 = DynScale(jnp.float32(2.0 ** 24), jnp.int32(0))       # deliberately start with a scale that's too large
p_fp16, ls_end, scales, finites, losses_fp16 = train_fp16_dynamic(p_init, Xmp, Ymp, ls0, steps=400)
scales, finites = np.asarray(scales), np.asarray(finites)
first_ok = int(np.argmax(finites))
print("Scale for the first 12 steps:", [f"2^{int(np.log2(s))}" for s in scales[:12]])
print("Finite in the first 12 steps:", finites[:12].astype(int).tolist())
print(f"First finite gradient at step {first_ok}; total skipped steps {int((~finites).sum())}; final scale = 2^{int(np.log2(float(ls_end.scale)))}")
print(f"fp16 + dynamic scaling: loss {float(losses_fp16[first_ok]):.4f} → {float(losses_fp16[-1]):.5f} (fp32 reference {float(l32[-1]):.5f})")
assert not finites[0] and finites[first_ok] and scales[first_ok] < scales[0]
assert float(losses_fp16[-1]) < 0.2 * float(losses_fp16[first_ok])
assert p_fp16[0][0].dtype == jnp.float32
Scale for the first 12 steps: ['2^24', '2^23', '2^22', '2^21', '2^20', '2^19', '2^18', '2^17', '2^17', '2^17', '2^17', '2^17']
Finite in the first 12 steps: [0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1]
First finite gradient at step 7; total skipped steps 9; final scale = 2^21
fp16 + dynamic scaling: loss 0.7060 → 0.00025 (fp32 reference 0.00015)

⚠️ Note: Don't expect to find ready-made loss scaling in Optax. Optax 0.2.9 has no loss-scaling tool, only the companion optax.apply_if_finite / optax.skip_not_finite (skip this update when the gradients contain inf/NaN). Flax has flax.training.dynamic_scale.DynamicScale (a tool from the Linen era that can still be imported). DeepMind also has a separate small library, jmp, dedicated to mixed-precision policies and loss scaling. The hand-written version above puts both "scaling + skipping" inside one jit.

optax_names = [n for n in dir(optax) if "scal" in n.lower() and "loss" in n.lower()]
print("optax APIs whose names contain loss+scale:", optax_names)
print("optax APIs related to 'skipping non-finite updates':", [n for n in dir(optax) if "finite" in n])
from flax.training import dynamic_scale
print("Flax's DynamicScale:", dynamic_scale.DynamicScale)
assert optax_names == [] and hasattr(optax, "apply_if_finite")

# optax.apply_if_finite: on a NaN gradient the update is 0, and it's counted
safe = optax.apply_if_finite(optax.sgd(0.1), max_consecutive_errors=5)
st = safe.init({"w": jnp.ones(2)})
u, st = safe.update({"w": jnp.array([jnp.nan, 1.0])}, st, {"w": jnp.ones(2)})
print("NaN gradient → update", u["w"], "| notfinite_count =", int(st.notfinite_count))
assert float(jnp.abs(u["w"]).sum()) == 0.0 and int(st.notfinite_count) == 1
optax APIs whose names contain loss+scale: []
optax APIs related to 'skipping non-finite updates': ['apply_if_finite', 'skip_not_finite']
Flax's DynamicScale: <class 'flax.training.dynamic_scale.DynamicScale'>
NaN gradient → update [0. 0.] | notfinite_count = 1

Things to keep in mind

  • Hardware: how much faster it actually gets depends on GPU / TPU support for 16-bit math (Tensor Cores, bf16 units), so confirm with profiling. (This machine is a CPU with no bf16 matrix units, so this section compares accuracy only, not speed.)
  • Numerical stability: bf16 is generally very reliable; fp16 needs a careful implementation plus loss scaling. Sensitive operations such as variance computation in normalization layers, softmax and the final loss should usually stay in fp32.
  • Accuracy: usually close to pure fp32 training, but verify the final model quality.
  • Debugging: the usual steps are checking for NaNs and confirming that gradient scaling under fp16 is correct.

🔍 Extra: a practical checklist for mixed precision

🟢 Must-read

  1. What stays in fp32: master weights, optimizer state, the loss, LayerNorm / BatchNorm statistics, softmax reductions (or at least accumulate them in fp32), and the global gradient norm.
  2. Adam state can be lower precision too: optax.adam(..., mu_dtype=jnp.bfloat16) stores the first moment in bf16, bringing optimizer state from 2× the parameters down to 1.5× (see 6.9). The second moment is more sensitive to precision and is usually kept in fp32.
  3. Communication can be lower precision too: all-reducing bf16 gradients in data parallelism halves the communication volume, but the accumulation error grows with the number of devices, so verify it.
  4. bf16 isn't a cure-all: bf16 has only 8 significant binary digits (about 2–3 significant decimal digits). When the update $\eta g$ is smaller than about $2^{-9}$ to $2^{-8}$ of the parameter itself (half the spacing between neighboring values), storing parameters in bf16 means the update just gets rounded away. That's why you have to keep fp32 master weights (there's a demo in the exercises).

6.8 Model parallelism strategies

Data parallelism (with gradient accumulation and checkpointing) lets you scale the amount of data and the effective batch, but it assumes every device can hold the full model. If the parameters and intermediate activations are too big for a single accelerator even with mixed precision, you have to split the model itself across devices. That's model parallelism. Data parallelism is "the same model, different data shards"; model parallelism is "layers or parts of layers go to different devices, and the data flows across devices through the split-up model".

Tensor parallelism (intra-layer parallelism)

Split the computation inside a single layer across devices. This is especially effective for the huge linear layers in Transformers. Take $Y=XW$: split $W$ by columns into $W=[W_1, W_2]$. Send the input $X$ to both devices; device 1 computes $Y_1=XW_1$, device 2 computes $Y_2=XW_2$, and concatenating them gives $Y=[Y_1, Y_2]$. You can also split by rows, in which case each device gets a partial sum and you need one all-reduce to add them up. More complex operations like attention can be parallelized the same way.

Difficulties:

  • Communication overhead: both the forward and backward pass of a single layer need frequent communication (all-gather, reduce-scatter, all-reduce), which becomes a bottleneck on a slow interconnect;
  • Implementation complexity: you have to handle tensor splitting, communication primitives and gradient propagation across shards carefully.

Pipeline parallelism (inter-layer parallelism)

Split by layers: assign consecutive layers (a stage) to different devices, so the output activations of one stage are the input of the next. For example, 4 layers L1–L4 on 4 devices D1–D4. A naive implementation processes one batch at a time, so at any moment only one device is working while the rest sit idle; those idle periods are called pipeline bubbles. The fix is micro-batching: split the batch into micro-batches MB1, MB2, …; once D1 finishes MB1 it hands it to D2 and moves on to MB2, so several devices can work at the same time. The bubbles at the start and end are still there, but they're amortized over many micro-batches.

Difficulties: bubbles can't be avoided completely; the stages need load balancing, or the slowest stage becomes the bottleneck; and scheduling micro-batches, activations and backward gradient flow is complex.

Hybrid approaches

  • Data parallelism + pipeline parallelism: replicate the whole pipeline several times, each copy handling a different data shard;
  • Data parallelism + tensor parallelism: use tensor parallelism within layers, then replicate that tensor-parallel model for data parallelism;
  • Fully sharded data parallelism (FSDP): shard parameters, gradients and optimizer state across the data-parallel devices, and all-gather parameters only when the computation needs them. You can think of it as "parameters are split as in model parallelism, but used in a data-parallel setting" (i.e., ZeRO-3).

The course's conclusion: you can implement all these strategies with just primitives like pmap, the lax.p* collectives and device placement, but it's very hard, so in practice people usually rely on higher-level libraries.

⚠️ Note: Older material often says that tensor parallelism in JAX "uses pmap over the subset of devices assigned to the parallel layer, together with jax.lax.psum / all_gather", that pipeline parallelism "places things manually with jax.device_put", and that for FSDP "similar features are appearing in the ecosystem". That's the pmap-era approach. In JAX 0.12, tensor parallelism, FSDP and any combination of them are just "annotating arrays with a PartitionSpec": use jit + NamedSharding on a multi-dimensional Mesh, and the compiler (GSPMD / Shardy) inserts the all-gather, reduce-scatter and all-reduce automatically. When you need hand-written communication, use jax.shard_map. Everything below is written this way, with an assert that it matches the single-device result.

First, create a 2 × 4 2D mesh, with 2 devices on the data axis and 4 on the model axis, and see how different PartitionSpecs distribute a matrix over the 8 devices.

mesh2d = jax.make_mesh((2, 4), ("data", "model"))
print(mesh2d, "| device layout:\n", np.vectorize(lambda d: d.id)(mesh2d.devices))
assert mesh2d.axis_types == (AxisType.Explicit, AxisType.Explicit)     # in 0.12, make_mesh defaults to Explicit

demo = jnp.arange(8 * 8, dtype=jnp.float32).reshape(8, 8)
with jax.set_mesh(mesh2d):
    for spec in [P(), P("data", None), P(None, "model"), P("data", "model"), P(("data", "model"), None)]:
        arr = jax.device_put(demo, spec)
        shard_shape = arr.addressable_shards[0].data.shape
        print(f"\nPartitionSpec{str(tuple(spec)):24s} → jax.typeof: {jax.typeof(arr)}, per-device shard {shard_shape}")
        jax.debug.visualize_array_sharding(arr, use_color=False)
Mesh('data': 2, 'model': 4, axis_types=(Explicit, Explicit)) | device layout:
 [[0 1 2 3]
 [4 5 6 7]]

PartitionSpec()                       → jax.typeof: float32[8,8], per-device shard (8, 8)
┌───────────────────────┐
│                       │
│                       │
│                       │
│                       │
│  CPU 0,1,2,3,4,5,6,7  │
│                       │
│                       │
│                       │
│                       │
└───────────────────────┘
PartitionSpec('data', None)           → jax.typeof: float32[8@data,8], per-device shard (4, 8)
┌───────────────────────┐
│                       │
│      CPU 0,1,2,3      │
│                       │
│                       │
├───────────────────────┤
│                       │
│      CPU 4,5,6,7      │
│                       │
│                       │
└───────────────────────┘
PartitionSpec(None, 'model')          → jax.typeof: float32[8,8@model], per-device shard (8, 2)
┌───────┬───────┬───────┬───────┐
│       │       │       │       │
│       │       │       │       │
│       │       │       │       │
│       │       │       │       │
│CPU 0,4│CPU 1,5│CPU 2,6│CPU 3,7│
│       │       │       │       │
│       │       │       │       │
│       │       │       │       │
│       │       │       │       │
└───────┴───────┴───────┴───────┘
PartitionSpec('data', 'model')        → jax.typeof: float32[8@data,8@model], per-device shard (4, 2)
┌───────┬───────┬───────┬───────┐
│       │       │       │       │
│ CPU 0 │ CPU 1 │ CPU 2 │ CPU 3 │
│       │       │       │       │
│       │       │       │       │
├───────┼───────┼───────┼───────┤
│       │       │       │       │
│ CPU 4 │ CPU 5 │ CPU 6 │ CPU 7 │
│       │       │       │       │
│       │       │       │       │
└───────┴───────┴───────┴───────┘
PartitionSpec(('data', 'model'), None) → jax.typeof: float32[8@(data,model),8], per-device shard (1, 8)
┌───────────────────────┐
│         CPU 0         │
├───────────────────────┤
│         CPU 1         │
├───────────────────────┤
│         CPU 2         │
├───────────────────────┤
│         CPU 3         │
├───────────────────────┤
│         CPU 4         │
├───────────────────────┤
│         CPU 5         │
├───────────────────────┤
│         CPU 6         │
├───────────────────────┤
│         CPU 7         │
└───────────────────────┘

✅ Megatron-style tensor parallelism: column split + row split

A Transformer MLP is $Y=\mathrm{act}(XW_1)W_2$. The classic Megatron-LM split:

  • Split $W_1\in\mathbb R^{D\times H}$ by columns (P(None, 'model')): each device gets $H/4$ columns and computes $\mathrm{act}(XW_1^{(i)})$. The activation function is elementwise, so no communication is needed;
  • Split $W_2\in\mathbb R^{H\times D}$ by rows (P('model', None)): each device uses its own $H/4$ rows to compute a partial sum $Z^{(i)}=h^{(i)}W_2^{(i)}$, and finally $Y=\sum_i Z^{(i)}$, which takes just one all-reduce.

$$ Y=\sum_{i=1}^{4}\mathrm{act}\!\left(XW_1^{(i)}\right)W_2^{(i)} $$

So the whole MLP forward pass has only one round of communication. Split the batch dimension along the data axis at the same time and you get 2D "data parallel × tensor parallel" parallelism.

In Explicit mode, the sharding of the result of a matmul whose contracting dimension is split is ambiguous (after summing the partial sums, how should the result be split?). JAX requires you to spell it out with out_sharding, otherwise it raises ShardingTypeError. That conveniently makes the communication "explicitly visible".

r = np.random.default_rng(0)
D_TP, H_TP, B_TP = 64, 256, 32
w1_np = r.normal(size=(D_TP, H_TP)) * 0.1
w2_np = r.normal(size=(H_TP, D_TP)) * 0.1
x_np = r.normal(size=(B_TP, D_TP))
ref64 = np.maximum(x_np @ w1_np, 0) @ w2_np                 # NumPy float64 reference

def tp_mlp(w1, w2, x):
    h = jax.nn.relu(x @ w1)                                  # column parallel: no communication
    print("  [trace] h  :", jax.typeof(h))
    y = jnp.matmul(h, w2, out_sharding=P("data", None))      # row parallel: partial sums → all-reduce
    print("  [trace] y  :", jax.typeof(y))
    return y

with jax.set_mesh(mesh2d):
    w1_tp = jax.device_put(jnp.asarray(w1_np, jnp.float32), P(None, "model"))
    w2_tp = jax.device_put(jnp.asarray(w2_np, jnp.float32), P("model", None))
    x_tp = jax.device_put(jnp.asarray(x_np, jnp.float32), P("data", None))
    print("w1:", jax.typeof(w1_tp), "| w2:", jax.typeof(w2_tp), "| x:", jax.typeof(x_tp))
    tp_fn = jax.jit(tp_mlp)
    y_tp = tp_fn(w1_tp, w2_tp, x_tp)
    print("Collectives:", collectives(tp_fn.lower(w1_tp, w2_tp, x_tp).compile().as_text()))

    # What happens without out_sharding?
    try:
        jax.jit(lambda w1, w2, x: jax.nn.relu(x @ w1) @ w2)(w1_tp, w2_tp, x_tp)
    except Exception as e:
        print("No out_sharding →", type(e).__name__, ":", str(e)[:120], "...")

err = np.abs(np.asarray(y_tp, np.float64) - ref64).max()
print(f"Tensor-parallel result vs NumPy float64: max error {err:.2e}")
assert err < 1e-5
w1: float32[64,256@model] | w2: float32[256@model,64] | x: float32[32@data,64]
  [trace] h  : float32[32@data,256@model]
  [trace] y  : float32[32@data,64]
Collectives: {'all-reduce': 1, 'all-gather': 0, 'reduce-scatter': 0, 'all-to-all': 0, 'collective-permute': 0}
No out_sharding → ShardingTypeError : Contracting dimensions are sharded and it is ambiguous how the output should be sharded. Please specify the output shard ...
Tensor-parallel result vs NumPy float64: max error 6.49e-07

✅ The same thing written by hand with jax.shard_map

jit + sharding is a "global view": you write the global computation and the compiler handles communication. jax.shard_map is a "per-device view": the function body sees this device's shard, and you write the communication yourself (lax.psum and friends). It suits cases where you need precise control over communication (for example, hand-written ops that overlap communication with compute), and it shows the Megatron math most clearly: each device computes a partial sum, and psum sums along the model axis.

def tp_mlp_local(w1_blk, w2_blk, x_blk):
    # All shapes here are "this device's shard"
    h = jax.nn.relu(x_blk @ w1_blk)              # (B/2, H/4)
    partial_sum = h @ w2_blk                     # (B/2, D): just a partial sum
    return lax.psum(partial_sum, "model")        # sum along the model axis (all-reduce)

with jax.set_mesh(mesh2d):
    sm = jax.jit(jax.shard_map(tp_mlp_local,
                               in_specs=(P(None, "model"), P("model", None), P("data", None)),
                               out_specs=P("data", None)))
    y_sm = sm(w1_tp, w2_tp, x_tp)
    shapes = jax.shard_map(lambda a, b, c: (jnp.zeros(a.shape), jnp.zeros(b.shape), jnp.zeros(c.shape)),
                           in_specs=(P(None, "model"), P("model", None), P("data", None)),
                           out_specs=(P(None, "model"), P("model", None), P("data", None)))(w1_tp, w2_tp, x_tp)
print("Shard shapes each device sees: w1", shapes[0].addressable_shards[0].data.shape,
      "w2", shapes[1].addressable_shards[0].data.shape, "x", shapes[2].addressable_shards[0].data.shape)
assert max_diff(y_sm, y_tp) < 1e-6
print(f"Hand-written shard_map version matches the automatic jit version ✓ (difference {max_diff(y_sm, y_tp):.1e})")
Shard shapes each device sees: w1 (64, 64) w2 (64, 64) x (16, 64)
Hand-written shard_map version matches the automatic jit version ✓ (difference 0.0e+00)

✅ 2D parallel training: sharded training vs single-device training

Hook the MLP above up to Adam and train for 5 steps. In the sharded version, parameters are split the Megatron way and data is split along the data axis. Optax state is created with zeros_like, so it automatically inherits the parameters' sharding; the gradients are sharded like the parameters too (the all-reduce in the backward pass is inserted by the compiler). At the end we assert that the parameters match single-device training.

Y_np = r.normal(size=(B_TP, D_TP))
tx_mp = optax.adam(1e-2)
tx_strict = optax.sgd(0.1, momentum=0.9)       # the update is linear in the gradient, so we use it for a strict cross-check

def make_train_step(fwd, tx=tx_mp):
    def loss(p, x, y):
        return jnp.mean((fwd(p, x) - y) ** 2)
    @jax.jit
    def step(p, s, x, y):
        l, g = jax.value_and_grad(loss)(p, x, y)
        u, s = tx.update(g, s, p)
        return optax.apply_updates(p, u), s, l
    return step

def train_n(step, p, tx, x, y, n=5):
    s = tx.init(p)
    for _ in range(n):
        p, s, l = step(p, s, x, y)
    return p, s, l

def fwd_single(p, x):
    return jax.nn.relu(x @ p["w1"]) @ p["w2"]

def fwd_tp(p, x):
    h = jax.nn.relu(x @ p["w1"])
    return jnp.matmul(h, p["w2"], out_sharding=P("data", None))

params_np = {"w1": w1_np.astype(np.float32), "w2": w2_np.astype(np.float32)}
X32, Y32 = x_np.astype(np.float32), Y_np.astype(np.float32)

# Single-device reference
p_ref = jax.tree.map(jnp.asarray, params_np); s_ref = tx_mp.init(p_ref)
step_single = make_train_step(fwd_single)
for _ in range(5):
    p_ref, s_ref, l_ref = step_single(p_ref, s_ref, jnp.asarray(X32), jnp.asarray(Y32))

with jax.set_mesh(mesh2d):
    p_tp = {"w1": jax.device_put(params_np["w1"], P(None, "model")),
            "w2": jax.device_put(params_np["w2"], P("model", None))}
    s_tp = tx_mp.init(p_tp)
    xs, ys = jax.device_put(X32, P("data", None)), jax.device_put(Y32, P("data", None))
    step_tp = make_train_step(fwd_tp)
    print("Communication:", collectives(step_tp.lower(p_tp, s_tp, xs, ys).compile().as_text()))
    for _ in range(5):
        p_tp, s_tp, l_tp = step_tp(p_tp, s_tp, xs, ys)
    print("w1 after training:", jax.typeof(p_tp["w1"]), "| Adam nu(w2):", jax.typeof(s_tp[0].nu["w2"]))
    assert p_tp["w1"].sharding.spec == P(None, "model") and s_tp[0].nu["w2"].sharding.spec == P("model", None)

d = max_diff(p_tp, p_ref)
print(f"After 5 Adam steps: loss {float(l_tp):.6f} vs {float(l_ref):.6f}; max parameter difference {d:.2e}")
assert np.isclose(float(l_tp), float(l_ref), rtol=1e-5) and d < 1e-3

# Strict cross-check: momentum SGD (the update is a linear function of the gradient, so it doesn't amplify rounding error)
p_ref_s, _, _ = train_n(make_train_step(fwd_single, tx_strict), jax.tree.map(jnp.asarray, params_np), tx_strict,
                        jnp.asarray(X32), jnp.asarray(Y32))
with jax.set_mesh(mesh2d):
    p_tp_s, _, _ = train_n(make_train_step(fwd_tp, tx_strict),
                           {"w1": jax.device_put(params_np["w1"], P(None, "model")),
                            "w2": jax.device_put(params_np["w2"], P("model", None))}, tx_strict, xs, ys)
d_s = max_diff(p_tp_s, p_ref_s)
print(f"After 5 momentum-SGD steps: max parameter difference {d_s:.2e}")
assert d_s < 1e-6
Communication: {'all-reduce': 2, 'all-gather': 0, 'reduce-scatter': 0, 'all-to-all': 0, 'collective-permute': 0}
w1 after training: float32[64,256@model] | Adam nu(w2): float32[256@model,64]
After 5 Adam steps: loss 0.327869 vs 0.327869; max parameter difference 1.59e-04
After 5 momentum-SGD steps: max parameter difference 2.98e-08

✅ FSDP / ZeRO-3: parameters, gradients and optimizer state are all sharded

FSDP doesn't change how things are computed (each device still computes full layers on its own data shard). It only changes how things are stored:

  • Parameters are split along one dimension across all devices (here an 8-device fsdp axis, P('fsdp', None)), so each device stores only $1/8$;
  • When the forward pass reaches a layer, it first all-gathers the full weights and throws them away after use;
  • In the backward pass, the weight gradients go through a reduce-scatter (each device gets only the sum for its own slice of the gradient);
  • Optimizer state is sharded just like the parameters, so the update also happens only on the local shard.

In Explicit mode, "all-gather before use" is just reshard(w, P()); its transpose (in the backward pass) is exactly a reduce-scatter.

mesh_fsdp = jax.make_mesh((8,), ("fsdp",))

def fwd_fsdp(p, x):
    w1 = reshard(p["w1"], P())            # all-gather: temporarily get the full weights
    w2 = reshard(p["w2"], P())
    return jax.nn.relu(x @ w1) @ w2

with jax.set_mesh(mesh_fsdp):
    p_f = {"w1": jax.device_put(params_np["w1"], P("fsdp", None)),
           "w2": jax.device_put(params_np["w2"], P("fsdp", None))}
    s_f = tx_mp.init(p_f)
    xs_f, ys_f = jax.device_put(X32, P("fsdp", None)), jax.device_put(Y32, P("fsdp", None))
    step_f = make_train_step(fwd_fsdp)
    hlo_f = step_f.lower(p_f, s_f, xs_f, ys_f).compile().as_text()
    for _ in range(5):
        p_f, s_f, l_f = step_f(p_f, s_f, xs_f, ys_f)
    g_f = jax.jit(jax.grad(lambda p, x, y: jnp.mean((fwd_fsdp(p, x) - y) ** 2)))(p_f, xs_f, ys_f)
    print("Params        :", jax.typeof(p_f["w1"]), jax.typeof(p_f["w2"]))
    print("Grads         :", jax.typeof(g_f["w1"]), jax.typeof(g_f["w2"]))
    print("Adam mu       :", jax.typeof(s_f[0].mu["w1"]), "| nu:", jax.typeof(s_f[0].nu["w2"]))
    print("Communication :", collectives(hlo_f))

def per_device_bytes(tree):
    return sum(x.addressable_shards[0].data.nbytes for x in jax.tree.leaves(tree) if hasattr(x, "addressable_shards"))

full = tree_bytes(p_f) + tree_bytes(s_f[0].mu) + tree_bytes(s_f[0].nu)
local = per_device_bytes(p_f) + per_device_bytes(s_f[0].mu) + per_device_bytes(s_f[0].nu)
print(f"Params + Adam state: {full / 1024:.0f} KiB global, {local / 1024:.0f} KiB per device (1/{full // local})")
assert full == 8 * local
assert g_f["w1"].sharding.spec == P("fsdp", None)                 # gradients are sharded too
assert collectives(hlo_f)["all-gather"] >= 2

d = max_diff(p_f, p_ref)
print(f"FSDP + Adam after 5 steps vs single device: loss {float(l_f):.6f} vs {float(l_ref):.6f}, max parameter difference {d:.2e}")
assert np.isclose(float(l_f), float(l_ref), rtol=1e-5) and d < 1e-3
with jax.set_mesh(mesh_fsdp):
    p_f_s, _, _ = train_n(make_train_step(fwd_fsdp, tx_strict),
                          {"w1": jax.device_put(params_np["w1"], P("fsdp", None)),
                           "w2": jax.device_put(params_np["w2"], P("fsdp", None))}, tx_strict, xs_f, ys_f)
print(f"FSDP + momentum SGD after 5 steps: max parameter difference {max_diff(p_f_s, p_ref_s):.2e}")
assert max_diff(p_f_s, p_ref_s) < 1e-6
Params        : float32[64@fsdp,256] float32[256@fsdp,64]
Grads         : float32[64@fsdp,256] float32[256@fsdp,64]
Adam mu       : float32[64@fsdp,256] | nu: float32[256@fsdp,64]
Communication : {'all-reduce': 1, 'all-gather': 2, 'reduce-scatter': 0, 'all-to-all': 0, 'collective-permute': 0}
Params + Adam state: 384 KiB global, 48 KiB per device (1/8)
FSDP + Adam after 5 steps vs single device: loss 0.327869 vs 0.327869, max parameter difference 2.63e-05
FSDP + momentum SGD after 5 steps: max parameter difference 2.98e-08

🔍 Extra: why isn't sharded training with Adam "bit-identical" to single-device training? ⚪ Optional: when you're verifying that sharded training is correct

After sharding, the summation order of matmuls changes (partial sums first, then all-reduce), so the gradients pick up rounding differences on the order of $10^{-7}$. The SGD update is a linear function of the gradient, so those differences carry through unchanged (measured above at $<10^{-6}$). The Adam update is $\hat m/(\sqrt{\hat v}+\epsilon)$, which normalizes gradient components close to 0 (gradients around $10^{-8}$ caused by ReLU or symmetry) up to the order of $\pm\eta$, so the rounding differences get amplified into parameter differences around $10^{-4}$. This isn't a bug. To verify that sharding is correct, compare the gradients, use a linear optimizer, or look at the relative error of the loss curve.

Note: in this jit auto-sharded version, the CPU backend implements the reduce-scatter in the backward pass as an all-reduce + slice (reduce-scatter is 0 in the counts above). Still, the final sharding of the gradient (64@fsdp) shows that semantically it's a reduce-scatter. Exercise 3 at the end writes the same FSDP by hand with shard_map + all_gather, and a real reduce-scatter shows up in the compiled output.

✅ 2D parallel training with NNX, cross-checked against single-device NNX

The ShardedMLP from section 6.4 (nnx.with_partitioning annotations + out_sharding) can be trained directly. The control group is a plain model with the same initializers and the same Rngs, but no sharding annotations. The two start from exactly the same values (jax_threefry_partitionable=True guarantees that sharded random generation produces the same numbers as unsharded).

class PlainMLP(nnx.Module):
    def __init__(self, din, dh, dout, *, rngs):
        init = nnx.initializers.lecun_normal()
        self.up = nnx.Linear(din, dh, use_bias=False, kernel_init=init, rngs=rngs)
        self.down = nnx.Linear(dh, dout, use_bias=False, kernel_init=init, rngs=rngs)
    def __call__(self, x):
        return self.down(nnx.relu(self.up(x)))

@nnx.jit
def nnx_mse_step(model, optimizer, x, y):
    loss, grads = nnx.value_and_grad(lambda m: jnp.mean((m(x) - y) ** 2))(model)
    optimizer.update(model, grads)
    return loss

plain = PlainMLP(D_TP, H_TP, D_TP, rngs=nnx.Rngs(0))
plain_opt = nnx.Optimizer(plain, optax.adam(1e-2), wrt=nnx.Param)
with jax.set_mesh(mesh2d):
    sharded = ShardedMLP(D_TP, H_TP, D_TP, rngs=nnx.Rngs(0))
    sharded_opt = nnx.Optimizer(sharded, optax.adam(1e-2), wrt=nnx.Param)
    assert max_diff(nnx.state(sharded, nnx.Param), nnx.state(plain, nnx.Param)) == 0.0   # identical initial values
    xs, ys = jax.device_put(X32, P("data", None)), jax.device_put(Y32, P("data", None))
    l_sh = [float(nnx_mse_step(sharded, sharded_opt, xs, ys)) for _ in range(5)]
    print("Sharded up.kernel:", jax.typeof(sharded.up.kernel[...]),
          "| Adam mu:", jax.typeof(sharded_opt.opt_state[0].mu["up"]["kernel"][...]))
l_pl = [float(nnx_mse_step(plain, plain_opt, jnp.asarray(X32), jnp.asarray(Y32))) for _ in range(5)]
d = max_diff(nnx.state(sharded, nnx.Param), nnx.state(plain, nnx.Param))
print("Sharded loss:", np.round(l_sh, 6), "\nSingle-device loss:", np.round(l_pl, 6), f"\nMax parameter difference {d:.2e}")
assert np.allclose(l_sh, l_pl, rtol=1e-5) and d < 1e-5
Sharded up.kernel: float32[64,256@model] | Adam mu: float32[64,256@model]
Sharded loss: [1.62052  1.047089 0.719013 0.460227 0.314597] 
Single-device loss: [1.62052  1.047089 0.719013 0.460227 0.314597] 
Max parameter difference 2.68e-07

Pipeline parallelism: concept and schedule diagram

Pipeline parallelism has no "one-line annotation" version in JAX (you use shard_map + lax.ppermute to pass activations between stages, and libraries like MaxText have ready-made implementations), so here we only cover scheduling and bubbles.

GPipe schedule: $p$ stages, $m$ micro-batches. In the forward pass, stage $s$ processes micro-batch $t-s$ at time step $t$; the backward pass starts only after all forward passes are done. The total number of time steps is $m+p-1$ (counting only the forward pass), and each stage is busy for only $m$ of them, so the bubble fraction is

$$ \text{bubble}=\frac{p-1}{m+p-1}. $$

When $m\gg p$ the bubble goes to 0, but the activations of every micro-batch have to be kept until the backward pass (memory grows with $m$). The 1F1B schedule has each stage run the matching backward pass as soon as possible after each forward pass, which caps the number of in-flight activations at $p$; the bubble fraction stays the same.

def gpipe_forward_schedule(p, m):
    """Return grid[stage][t] = micro-batch index (-1 means idle)."""
    T = m + p - 1
    grid = [[-1] * T for _ in range(p)]
    for s in range(p):
        for mb in range(m):
            grid[s][s + mb] = mb
    return grid

def show(grid):
    for s, row in enumerate(grid):
        print(f"stage {s} (device {s}) | " + " ".join(f"F{mb}" if mb >= 0 else " ·" for mb in row))

for p_, m_ in [(4, 1), (4, 4), (4, 16)]:
    g = gpipe_forward_schedule(p_, m_)
    busy = sum(mb >= 0 for row in g for mb in row)
    bubble = 1 - busy / (p_ * len(g[0]))
    print(f"\np = {p_} stages, m = {m_} micro-batches: bubble fraction {bubble:.1%} (formula (p-1)/(m+p-1) = {(p_ - 1) / (m_ + p_ - 1):.1%})")
    if m_ <= 4:
        show(g)
    assert np.isclose(bubble, (p_ - 1) / (m_ + p_ - 1))
p = 4 stages, m = 1 micro-batches: bubble fraction 75.0% (formula (p-1)/(m+p-1) = 75.0%)
stage 0 (device 0) | F0  ·  ·  ·
stage 1 (device 1) |  · F0  ·  ·
stage 2 (device 2) |  ·  · F0  ·
stage 3 (device 3) |  ·  ·  · F0

p = 4 stages, m = 4 micro-batches: bubble fraction 42.9% (formula (p-1)/(m+p-1) = 42.9%)
stage 0 (device 0) | F0 F1 F2 F3  ·  ·  ·
stage 1 (device 1) |  · F0 F1 F2 F3  ·  ·
stage 2 (device 2) |  ·  · F0 F1 F2 F3  ·
stage 3 (device 3) |  ·  ·  · F0 F1 F2 F3

p = 4 stages, m = 16 micro-batches: bubble fraction 15.8% (formula (p-1)/(m+p-1) = 15.8%)
# A minimal "pipeline": the weights of 4 stages live on 4 different devices, and micro-batches flow through in turn (run sequentially, only to illustrate the data flow)
devs = jax.devices()[:4]
r = np.random.default_rng(7)
stage_w = [jax.device_put(jnp.asarray(r.normal(size=(32, 32)) / np.sqrt(32), jnp.float32), devs[s]) for s in range(4)]
stage_fn = jax.jit(lambda w, h: jnp.tanh(h @ w))       # one layer per stage; weights are passed in as arguments

X_pipe = jnp.asarray(r.normal(size=(64, 32)), jnp.float32)
micro = jnp.split(X_pipe, 8)                            # 8 micro-batches
outs = []
for mb in micro:
    h = mb
    for s in range(4):
        h = jax.device_put(h, devs[s])                  # "send" the activation to the device of the next stage
        h = stage_fn(stage_w[s], h)
    outs.append(h)
y_pipe = jnp.concatenate([jax.device_put(o, devs[0]) for o in outs])
print("Device of each stage's weights:", [list(w.devices())[0].id for w in stage_w], "| device of the last micro-batch's output:", list(outs[-1].devices())[0].id)

# Cross-check against the unsplit whole-model forward pass (NumPy float64)
h64 = np.asarray(X_pipe, np.float64)
for w in stage_w:
    h64 = np.tanh(h64 @ np.asarray(w, np.float64))
assert np.abs(np.asarray(y_pipe, np.float64) - h64).max() < 1e-5
print("The pipeline output (placed by stage + micro-batches) matches the whole-model forward pass ✓")
Device of each stage's weights: [0, 1, 2, 3] | device of the last micro-batch's output: 3
The pipeline output (placed by stage + micro-batches) matches the whole-model forward pass ✓

🔍 Extra: how to choose a parallel strategy

⚪ Optional: when you're designing the parallel setup for multi-node, multi-GPU training

StrategyWhat gets splitCommunication per layerWhen to use it
Data parallelism (DP)dataone gradient all-reduce per stepthe model fits on one device
FSDP / ZeRO-3data + parameters/gradients/optimizer stateall-gather in the forward pass and reduce-scatter in the backward pass, per layerthe model doesn't fit but a single layer does; about 1.5× the communication of DP, can overlap with compute
Tensor parallelism (TP)weight matrices within a layer1–2 all-reduces per layer in each of the forward and backward passesa single layer is too big, or you want less activation memory per device; communication is frequent, so keep it within a group of devices on a fast interconnect (e.g., the 8 GPUs in one machine)
Pipeline parallelism (PP)between layerspoint-to-point activation transfer between stagesscaling across slow interconnects; has bubbles
Sequence / context parallelismthe sequence dimensionexchange K/V in attentionvery long sequences

A common combination in practice is TP within a machine × FSDP (or DP) across machines, adding PP when needed. In JAX this is just a multi-dimensional mesh (say ('data', 'fsdp', 'tensor')) plus a PartitionSpec for each parameter. The same model code switches parallel strategy just by switching a set of annotations. MaxText manages these annotations with a rule table mapping "logical axis names → mesh axis names" (in Flax: write logical axis names in nnx.with_partitioning, then map them to mesh axes with the nnx.logical_axis_rules context; the Linen equivalent is nn.logical_axis_rules).


6.9 Optimization algorithms at scale

Data parallelism and gradient accumulation lead to larger effective batches, which changes how optimizers behave. SGD with momentum and Adam are the foundation, but at scale you often need optimizers designed specifically for very large batches or particular regularization needs.

AdamW: decoupled weight decay

Using "standard" L2 regularization with Adam (adding $\lambda\theta$ to the gradient) usually doesn't work well: the decay term goes into the adaptive moments $m_t$ and $v_t$, so weights with large historical gradients end up decaying less. AdamW decouples weight decay from the gradient update and applies it directly to the weights after the Adam step, more like weight decay in SGD:

  1. Compute the gradient $g_t=\nabla\mathcal L(\theta_{t-1})$;
  2. Update the moment estimates $m_t=\beta_1 m_{t-1}+(1-\beta_1)g_t$, $v_t=\beta_2 v_{t-1}+(1-\beta_2)g_t^2$;
  3. Bias correction $\hat m_t=m_t/(1-\beta_1^t)$, $\hat v_t=v_t/(1-\beta_2^t)$;
  4. Adam step: $\theta'_t=\theta_{t-1}-\eta\,\hat m_t/(\sqrt{\hat v_t}+\epsilon)$;
  5. Weight decay: $\theta_t=\theta'_t-\eta\lambda\,\theta_{t-1}$.

The course notes that this decoupling usually gives large models like Transformers better generalization and more stable training, and it's now the standard choice for training large language models.

LAMB: layer-wise adaptation

With very large batches (tens of thousands or more, usually achieved with data parallelism over many accelerators), Adam can become unstable or need very careful learning-rate tuning (in particular enough warmup). LAMB (Layer-wise Adaptive Moments optimizer for Batch training) was designed for this case: first compute each layer's update direction $\Delta\theta_l$ as in AdamW, then scale it by the ratio of that layer's weight norm to its update norm (the trust ratio):

$$ r_l=\frac{\lVert\theta_l\rVert}{\lVert\Delta\theta_l\rVert},\qquad \theta_{l,t}=\theta_{l,t-1}-\eta\, r_l\,\Delta\theta_l . $$

With large batches, where gradient variance gets smaller, it keeps some layers from updating too much or too little. The course mentions that LAMB let BERT train with very large batches and noticeably faster.

Other considerations

  • Learning-rate schedules: large-model training depends heavily on the schedule. Warmup: increase the learning rate linearly from a small value over the first few thousand steps to help early stability; decay: afterwards, decrease it linearly, with a cosine, with inverse square root, and so on.
  • Optimizer-state memory: Adam(W) and LAMB keep two pieces of state per parameter, which in a large model can be about as big as the parameters themselves or even bigger. In distributed training, optimizer state has to be sharded along with the parameters and gradients.

In practice with Optax (course code)

Typical Optax usage has two steps: first initialize the optimizer state from the parameters; then, in the training step, compute the updates and the new state from the gradients and the current state. In the course code, the scheduler's arguments are written literally as ...:

# ❌ Wrong: the arguments of warmup_cosine_decay_schedule(...) are an ellipsis placeholder, so it can't run
try:
    learning_rate_schedule = optax.warmup_cosine_decay_schedule(...)
    print("Measured: it runs")
except TypeError as e:
    print("Measured error: TypeError:", e)
Measured error: TypeError: warmup_cosine_decay_schedule() missing 3 required positional arguments: 'peak_value', 'warmup_steps', and 'decay_steps'
# ✅ Recommended: give the full arguments; update_step receives parameters and gradients as jit arguments
learning_rate_schedule = optax.warmup_cosine_decay_schedule(
    init_value=0.0, peak_value=1e-3, warmup_steps=1_000, decay_steps=100_000, end_value=1e-5)
optimizer = optax.adamw(learning_rate=learning_rate_schedule, weight_decay=0.01)

@jax.jit
def update_step(params, grads, opt_state):
    updates, opt_state = optimizer.update(grads, opt_state, params)
    return optax.apply_updates(params, updates), opt_state

p_demo = {"w": jnp.ones((4, 3)), "b": jnp.zeros(3)}
g_demo = {"w": jnp.full((4, 3), 0.5), "b": jnp.full(3, -0.2)}
s_demo = optimizer.init(p_demo)
p1, s1 = update_step(p_demo, g_demo, s_demo)
print("w[0] after one step:", p1["w"][0], "| learning rate at step 0 =", float(learning_rate_schedule(0)),
      "| parameters completely unchanged?", bool(jnp.all(p1["w"] == p_demo["w"])))
assert bool(jnp.all(p1["w"] == p_demo["w"]))          # with η=0 not even weight decay happens
p2, s2 = update_step(p1, g_demo, s1)
assert not bool(jnp.all(p2["w"] == p1["w"]))          # updates start from step 1
w[0] after one step: [1. 1. 1.] | learning rate at step 0 = 0.0 | parameters completely unchanged? True

The output above shows an easy-to-miss detail: in Optax's adamw, weight decay is also multiplied by the learning rate ($\eta\lambda\theta$), so at warmup step 0, where $\eta=0$, there's no weight decay either. Below we compute one AdamW step by hand with the course's 5-step formula, then one LAMB step by hand with the trust-ratio formula, and cross-check both against Optax.

r = np.random.default_rng(0)
theta = {"w": jnp.asarray(r.normal(size=(8, 4)), jnp.float32), "b": jnp.asarray(r.normal(size=4), jnp.float32)}
grad = {"w": jnp.asarray(r.normal(size=(8, 4)), jnp.float32), "b": jnp.asarray(r.normal(size=4), jnp.float32)}
eta, lam, b1, b2, eps = 1e-2, 0.1, 0.9, 0.999, 1e-8

# —— AdamW: formulas 1–5 ——
def adamw_one_step(th, g):
    m = (1 - b1) * g; v = (1 - b2) * g ** 2                  # m_0 = v_0 = 0
    m_hat, v_hat = m / (1 - b1), v / (1 - b2)
    th_prime = th - eta * m_hat / (jnp.sqrt(v_hat) + eps)    # 4) Adam step
    return th_prime - eta * lam * th                         # 5) decoupled weight decay
manual = jax.tree.map(adamw_one_step, theta, grad)
ow = optax.adamw(eta, b1=b1, b2=b2, eps=eps, weight_decay=lam)
u, _ = ow.update(grad, ow.init(theta), theta)
print(f"AdamW: hand-written formula vs optax.adamw, max difference {max_diff(manual, optax.apply_updates(theta, u)):.1e}")
assert max_diff(manual, optax.apply_updates(theta, u)) < 1e-6

# Not the same as "Adam + L2 regularization": L2 adds λθ to the gradient, which then gets normalized by 1/sqrt(v)
l2 = optax.adam(eta, b1=b1, b2=b2, eps=eps)
u_l2, _ = l2.update(jax.tree.map(lambda g, t: g + lam * t, grad, theta), l2.init(theta), theta)
print(f"Difference between Adam+L2 and AdamW after one step: {max_diff(optax.apply_updates(theta, u_l2), manual):.3e} (they're not the same thing)")
assert max_diff(optax.apply_updates(theta, u_l2), manual) > 1e-4

# —— LAMB: one trust ratio per parameter tensor ("layer") ——
def lamb_one_step(th, g, wd=0.01):
    m = (1 - b1) * g; v = (1 - b2) * g ** 2
    adam_dir = (m / (1 - b1)) / (jnp.sqrt(v / (1 - b2)) + 1e-6)     # optax.lamb defaults to eps=1e-6
    delta = adam_dir + wd * th                                     # AdamW direction
    r_l = jnp.linalg.norm(th) / jnp.linalg.norm(delta)             # trust ratio
    return th - eta * r_l * delta, r_l
lamb_manual = jax.tree.map(lambda t, g: lamb_one_step(t, g)[0], theta, grad)
ratios = jax.tree.map(lambda t, g: float(lamb_one_step(t, g)[1]), theta, grad)
ol = optax.lamb(eta, weight_decay=0.01)
u_l, _ = ol.update(grad, ol.init(theta), theta)
print(f"LAMB trust ratios: {ratios}; by hand vs optax.lamb max difference {max_diff(lamb_manual, optax.apply_updates(theta, u_l)):.1e}")
assert max_diff(lamb_manual, optax.apply_updates(theta, u_l)) < 1e-6
AdamW: hand-written formula vs optax.adamw, max difference 1.2e-07
Difference between Adam+L2 and AdamW after one step: 2.325e-03 (they're not the same thing)
LAMB trust ratios: {'b': 0.35019758343696594, 'w': 0.8128143548965454}; by hand vs optax.lamb max difference 7.5e-09

⚠️ Note: It's easy to think that Optax "automatically handles sharding of the optimizer state" inside a pmap function. In fact Optax itself has no sharding logic at all. Optimizer state is just a PyTree derived from the parameters with zeros_like and the like, so where it lives simply follows the parameters. In pmap the parameters are replicated, so the optimizer state is replicated too (a full copy on every device, no memory saved). Only when the parameters themselves are sharded (FSDP / tensor parallelism from section 6.8) does the state get sharded along with them. If you want "sharded optimizer state" (ZeRO-1/2), you have to annotate the state with shardings yourself.

big = {"w": jnp.zeros((1024, 256))}
# pmap style: parameters replicated on 8 devices → the state is 8 full copies too
rep = jax.device_put(big, NamedSharding(jax.make_mesh((8,), ("d",), axis_types=(AxisType.Auto,)), P()))
st_rep = optax.adam(1e-3).init(rep)
print("Replicated parameters → mu sharding:", st_rep[0].mu["w"].sharding.spec,
      "| mu bytes per device:", st_rep[0].mu["w"].addressable_shards[0].data.nbytes)
assert st_rep[0].mu["w"].addressable_shards[0].data.nbytes == big["w"].nbytes           # a full copy on every device

# Sharded parameters → the state is sharded along with them
with jax.set_mesh(mesh_fsdp):
    shd = {"w": jax.device_put(big["w"], P("fsdp", None))}
    st_shd = optax.adam(1e-3).init(shd)
    print("Sharded parameters → mu:", jax.typeof(st_shd[0].mu["w"]),
          "| mu bytes per device:", st_shd[0].mu["w"].addressable_shards[0].data.nbytes)
    assert st_shd[0].mu["w"].addressable_shards[0].data.nbytes * 8 == big["w"].nbytes

    # ZeRO-1: parameters replicated, only optimizer state sharded; you have to write the annotations yourself
    st_zero1 = jax.tree.map(lambda a: jax.device_put(a, P("fsdp", None)) if a.ndim == 2 else a,
                            optax.adam(1e-3).init(jax.device_put(big, P())))
    print("ZeRO-1 style (state annotated by hand) → mu:", jax.typeof(st_zero1[0].mu["w"]))
Replicated parameters → mu sharding: P() | mu bytes per device: 1048576
Sharded parameters → mu: float32[1024@fsdp,256] | mu bytes per device: 131072
ZeRO-1 style (state annotated by hand) → mu: float32[1024@fsdp,256]

✅ Common optimizers at scale: optimizer-state memory compared (all verified to work in Optax 0.2.9)

OptimizerOptax APIState per parameterCharacteristics / use cases
SGD + momentumoptax.sgd(lr, momentum=0.9)1 copyclassic for CV; pair with LARS for large batches
Adam / AdamWoptax.adam / optax.adamw2 copies ($m$, $v$)the default choice; standard for LLMs
AdamW (bf16 first moment)optax.adamw(..., mu_dtype=jnp.bfloat16)1.5 copiessaves 25% of state memory
Adafactoroptax.adafactorfor a matrix: two vectors, one for rows and one for columns ($n+m$ instead of $nm$), no momentum by defaulttiny state memory; used by T5, PaLM and others
LAMBoptax.lamb2 copiesvery large batches (BERT pretraining)
LARSoptax.lars1 copy (momentum)SGD variant for very large batches (large-batch ResNet training)
Lionoptax.lion1 copy (momentum), the update takes the signhalf the state of Adam; the learning rate should be 3–10× smaller than for AdamW, with weight decay scaled up accordingly
SM3optax.sm3second-order statistics approximated by rows + columns, but by default it also keeps 1 copy of momentum (momentum=0.9), about 1 copy in totalanother memory-saving adaptive method

optax.contrib also has newer methods such as Muon, Schedule-Free AdamW, Prodigy and DoG (all present in 0.2.9). Below we compute how many times the parameter bytes each optimizer's state takes, then look separately at the state shapes Adafactor uses for one large matrix.

opt_params = {"emb": jnp.zeros((4096, 512)), "w": jnp.zeros((512, 2048)), "b": jnp.zeros(2048)}
P_bytes = tree_bytes(opt_params)
candidates = {
    "sgd":                 optax.sgd(1e-2),
    "sgd+momentum":        optax.sgd(1e-2, momentum=0.9),
    "adamw":               optax.adamw(1e-3),
    "adamw (mu bf16)":     optax.adamw(1e-3, mu_dtype=jnp.bfloat16),
    "adafactor":           optax.adafactor(1e-3),
    "lamb":                optax.lamb(1e-3),
    "lars":                optax.lars(1e-2),
    "lion":                optax.lion(1e-4),
    "sm3":                 optax.sm3(1e-2),
}
state_ratio = {}
for name, o in candidates.items():
    state_ratio[name] = tree_bytes(o.init(opt_params)) / P_bytes
    print(f"{name:18s} state / params = {state_ratio[name]:.4f}")
assert state_ratio["sgd"] < 1e-3 and abs(state_ratio["adamw"] - 2) < 1e-3 and abs(state_ratio["lamb"] - 2) < 1e-3
assert abs(state_ratio["lion"] - 1) < 1e-3 and abs(state_ratio["lars"] - 1) < 1e-3
assert abs(state_ratio["adamw (mu bf16)"] - 1.5) < 1e-3
assert state_ratio["adafactor"] < 0.01

# For a (4096, 512) matrix, Adafactor stores statistics for just one row and one column
af_state = optax.adafactor(1e-3).init({"emb": jnp.zeros((4096, 512))})
shapes = [x.shape for x in jax.tree.leaves(af_state) if hasattr(x, "shape") and x.size > 1]
print("\nShapes of the non-scalar state Adafactor keeps for a (4096, 512) matrix:", shapes)
assert (4096,) in shapes and (512,) in shapes
sgd                state / params = 0.0000
sgd+momentum       state / params = 1.0000
adamw              state / params = 2.0000
adamw (mu bf16)    state / params = 1.5000
adafactor          state / params = 0.0029
lamb               state / params = 2.0000
lars               state / params = 1.0000
lion               state / params = 1.0000
sm3                state / params = 1.0029

Shapes of the non-scalar state Adafactor keeps for a (4096, 512) matrix: [(512,), (4096,)]

✅ A head-to-head run on a small problem

Same 4-layer MLP regression task, same initialization, same batch. Each optimizer uses a typical learning rate (not carefully tuned) and trains for 500 steps; we compare the final loss and the time per step. Note that this only shows "does it run, and roughly how does it do": rankings on a small problem don't carry over to large models. The advantages of Lion, LAMB and Adafactor show up mainly with large batches, large models and tight memory.

r = np.random.default_rng(42)
X_opt = jnp.asarray(r.uniform(-2, 2, size=(1024, 16)), jnp.float32)
Y_opt = jnp.sin(X_opt[:, :1] * X_opt[:, 1:2]) + 0.5 * jnp.cos(X_opt[:, 2:3])
sizes = [16, 128, 128, 128, 1]
p_opt0 = [{"w": jnp.asarray(r.normal(size=(a, b)) * np.sqrt(2 / a), jnp.float32), "b": jnp.zeros(b)}
          for a, b in zip(sizes[:-1], sizes[1:])]

def opt_loss(p, x, y):
    h = x
    for layer in p[:-1]:
        h = jax.nn.gelu(h @ layer["w"] + layer["b"])
    return jnp.mean((h @ p[-1]["w"] + p[-1]["b"] - y) ** 2)

def make_runner(tx, steps=500):
    @jax.jit
    def run(p, x, y):
        def step(carry, _):
            p, s = carry
            l, g = jax.value_and_grad(opt_loss)(p, x, y)
            u, s = tx.update(g, s, p)
            return (optax.apply_updates(p, u), s), l
        (p, s), ls = lax.scan(step, (p, tx.init(p)), None, length=steps)
        return ls
    return run

sched = lambda peak: optax.warmup_cosine_decay_schedule(0.0, peak, 50, 500, peak * 0.01)
contenders = {
    "sgd+momentum": optax.sgd(sched(3e-2), momentum=0.9),
    "adamw":        optax.adamw(sched(3e-3), weight_decay=1e-4),
    "adafactor":    optax.adafactor(sched(3e-3), multiply_by_parameter_scale=False),
    "lamb":         optax.lamb(sched(1e-2), weight_decay=1e-4),
    "lars":         optax.lars(sched(1.0), weight_decay=1e-4),
    "lion":         optax.lion(sched(3e-4), weight_decay=1e-3),
}
init_loss = float(opt_loss(p_opt0, X_opt, Y_opt))
print(f"Initial loss {init_loss:.4f}\n{'optimizer':14s} {'final loss':>10s} {'state/params':>9s} {'500 steps ms':>12s}")
final = {}
for name, tx_ in contenders.items():
    runner = make_runner(tx_)
    ls = runner(p_opt0, X_opt, Y_opt)
    final[name] = float(ls[-1])
    t_ms = bench(runner, p_opt0, X_opt, Y_opt, n=3)
    ratio = tree_bytes(tx_.init(p_opt0)) / tree_bytes(p_opt0)
    print(f"{name:14s} {final[name]:10.2e} {ratio:9.3f} {t_ms:12.1f}")
    assert np.isfinite(final[name]) and final[name] < 0.5 * init_loss, name     # every one brings the loss down
Initial loss 1.9866
optimizer      final loss state/params 500 steps ms
sgd+momentum     5.07e-04     1.000        885.8
adamw            1.06e-06     2.000        884.2
adafactor        1.88e-06     0.087       1007.6
lamb             1.68e-08     2.000        563.6
lars             1.98e-05     1.000        438.0
lion             1.97e-02     1.000        434.2

🔍 Extra: some lessons for optimization at scale

⚪ Optional: when tuning optimizer hyperparameters for large models

  1. Global gradient clipping is practically standard (optax.clip_by_global_norm(1.0)) and helps you survive loss spikes. Under FSDP / tensor parallelism, global_norm is computed across shards automatically (it's just a sum over sharded arrays, and the compiler inserts the all-reduce).
  2. Which parameters skip weight decay: LayerNorm scales and biases usually aren't decayed. Optax controls this with the mask argument (optax.adamw(..., mask=fn), where fn returns a boolean tree with the same structure as the parameters).
  3. Learning rate for large batches: the linear scaling rule (batch × k → lr × k) only holds within a certain range. Beyond that you need the layer-wise normalization of LARS / LAMB, or a longer warmup.
  4. Lion hyperparameters: its update is $\mathrm{sign}(\cdot)$, so every component takes a step of size $\eta$. The learning rate therefore needs to be small (about 1/3–1/10 of AdamW's), with weight decay scaled up accordingly so that $\eta\lambda$ stays the same.
  5. Optimizer-state dtype: mu_dtype=jnp.bfloat16 is the most cost-effective way to save memory. The second moment $v$ spans many orders of magnitude and usually stays in fp32 (or use Adafactor's factorization).
# Example: a mask that skips weight decay for LayerNorm / bias
p_mask_demo = {"dense": {"kernel": jnp.ones((3, 3)), "bias": jnp.ones(3)}, "ln": {"scale": jnp.ones(3), "bias": jnp.ones(3)}}
decay_mask = jax.tree.map_with_path(lambda path, _: path[-1].key == "kernel", p_mask_demo)
tx_masked = optax.adamw(1e-1, weight_decay=0.5, mask=decay_mask)
zero_g = jax.tree.map(jnp.zeros_like, p_mask_demo)
u_m, _ = tx_masked.update(zero_g, tx_masked.init(p_mask_demo), p_mask_demo)   # zero gradient, so only weight decay is left
print("mask:", decay_mask)
print("updates:", jax.tree.map(lambda x: float(x.ravel()[0]), u_m))
assert np.isclose(float(u_m["dense"]["kernel"][0, 0]), -0.05) and float(u_m["ln"]["scale"][0]) == 0.0
mask: {'dense': {'bias': False, 'kernel': True}, 'ln': {'bias': False, 'scale': False}}
updates: {'dense': {'bias': -0.0, 'kernel': -0.05000000074505806}, 'ln': {'bias': -0.0, 'scale': -0.0}}

6.10 Practice: implementing gradient checkpointing

This section gets hands-on with the memory–compute trade-off of jax.checkpoint.

How it works, again

Let $z=f(x)=f_2(f_1(x))$ and write $y=f_1(x)$.

  • Standard autodiff: the forward pass computes both $y$ and $z$ and keeps $y$ in memory so the backward pass can use it for the gradient of $f_2$;
  • Checkpointing $f_1$: after the forward pass computes $y$ it can be thrown away; when the backward pass needs $y$, it reruns $f_1(x)$ from the saved input $x$. That way you don't have to store a potentially large intermediate tensor like $y$.

To use it, pass the function you want to checkpoint to jax.checkpoint: checkpointed_f1 = jax.checkpoint(f1), then as usual y = checkpointed_f1(x), z = f2(y). When you take the gradient of the whole computation, autodiff knows the backward pass has to recompute checkpointed_f1.

The course example

A block with a large intermediate activation: $y=\mathrm{GELU}(xW_1)$, $z=yW_2$, with a mean squared error loss. hidden_dim = 8192 makes the intermediate activation $y$ (batch × hidden) relatively large.

def compute_block(x, W1, W2):
    """A block that may contain large intermediate activations."""
    y = jnp.dot(x, W1)
    y = jax.nn.gelu(y)            # GELU activation; y is the intermediate activation we'd like to avoid storing
    z = jnp.dot(y, W2)
    return z

def loss_fn_610(x, W1, W2, targets):
    z = compute_block(x, W1, W2)
    return jnp.mean((z - targets) ** 2)

grad_fn_standard = jax.jit(jax.value_and_grad(loss_fn_610, argnums=(1, 2)))

compute_block_checkpointed = jax.checkpoint(compute_block)
def loss_fn_checkpointed(x, W1, W2, targets):
    z = compute_block_checkpointed(x, W1, W2)
    return jnp.mean((z - targets) ** 2)
grad_fn_checkpointed = jax.jit(jax.value_and_grad(loss_fn_checkpointed, argnums=(1, 2)))

key = jax.random.key(42)
batch_size, input_dim, hidden_dim, output_dim = 64, 512, 8192, 512     # large hidden dimension
key, x_key, w1_key, w2_key, t_key = jax.random.split(key, 5)
x610 = jax.random.normal(x_key, (batch_size, input_dim))
W1_610 = jax.random.normal(w1_key, (input_dim, hidden_dim)) * 0.02
W2_610 = jax.random.normal(w2_key, (hidden_dim, output_dim)) * 0.02
t610 = jax.random.normal(t_key, (batch_size, output_dim))
# ❌ Wrong: timing with time.time() and timing the very first call → what you measure is mostly compile time
try:
    start_time = time.time()
    loss_std, (dW1_std, dW2_std) = grad_fn_standard(x610, W1_610, W2_610, t610)
    loss_std.block_until_ready(); dW1_std.block_until_ready(); dW2_std.block_until_ready()
    time_std = time.time() - start_time
    print("Running the standard version (compile + execute)...")
    print(f"Standard loss: {loss_std:.4f}\nStandard time: {time_std:.4f} s")

    start_time = time.time()
    loss_ckpt, (dW1_ckpt, dW2_ckpt) = grad_fn_checkpointed(x610, W1_610, W2_610, t610)
    loss_ckpt.block_until_ready(); dW1_ckpt.block_until_ready(); dW2_ckpt.block_until_ready()
    time_ckpt = time.time() - start_time
    print("Running the checkpointed version (compile + execute)...")
    print(f"Checkpointed loss: {loss_ckpt:.4f}\nCheckpointed time: {time_ckpt:.4f} s")

    print("Comparing gradients...")
    print(f"W1 max absolute difference: {jnp.max(jnp.abs(dW1_std - dW1_ckpt)):.2e}")
    print(f"W2 max absolute difference: {jnp.max(jnp.abs(dW2_std - dW2_ckpt)):.2e}")
except Exception as e:
    print("Measured error:", type(e).__name__, e)
Running the standard version (compile + execute)...
Standard loss: 1.2202
Standard time: 0.0482 s
Running the checkpointed version (compile + execute)...
Checkpointed loss: 1.2202
Checkpointed time: 0.0447 s
Comparing gradients...
W1 max absolute difference: 0.00e+00
W2 max absolute difference: 0.00e+00
# ✅ Recommended: time compilation and execution separately; warm up and average several runs for execution; quantify memory with memory_analysis()
def compile_time(f, *args):
    t0 = time.perf_counter()
    compiled = jax.jit(f).lower(*args).compile()
    return (time.perf_counter() - t0) * 1e3, compiled

args610 = (x610, W1_610, W2_610, t610)
report = {}
for name, f in [("standard", jax.value_and_grad(loss_fn_610, argnums=(1, 2))),
                ("checkpointed", jax.value_and_grad(loss_fn_checkpointed, argnums=(1, 2)))]:
    ct, compiled = compile_time(f, *args610)
    rt = bench(compiled, *args610, n=20)
    temp = compiled.memory_analysis().temp_size_in_bytes / MiB
    report[name] = (ct, rt, temp)
    print(f"{name:4s}: compile {ct:7.1f} ms | run {rt:6.2f} ms/step | temp memory {temp:6.3f} MiB")
saving_610 = 1 - report["checkpointed"][2] / report["standard"][2]
print(f"Temp memory saved for a single block: {saving_610:.1%}; run-time ratio: {report['checkpointed'][1] / report['standard'][1]:.2f}×")

# Gradient consistency (the course expects them to be "numerically very close")
assert float(jnp.max(jnp.abs(dW1_std - dW1_ckpt))) < 1e-6 and float(jnp.max(jnp.abs(dW2_std - dW2_ckpt))) < 1e-6
assert report["checkpointed"][2] < report["standard"][2]

print("\n—— Residuals saved by the standard version ——");   print_saved_residuals(loss_fn_610, *args610)
print("\n—— Residuals saved by the checkpointed version ——"); print_saved_residuals(loss_fn_checkpointed, *args610)
standard: compile    34.9 ms | run   3.32 ms/step | temp memory  6.125 MiB
checkpointed: compile    34.1 ms | run   3.66 ms/step | temp memory  4.125 MiB
Temp memory saved for a single block: 32.7%; run-time ratio: 1.10×

—— Residuals saved by the standard version ——
f32[] from a literal
f32[] from a literal
f32[64,512] from the argument x
f32[512,8192] from the argument W1
f32[8192,512] from the argument W2
f32[64,8192] output of dot_general from /tmp/ipykernel_66671/2742922787.py:3:8 (compute_block)
f32[64,8192] output of mul from /tmp/ipykernel_66671/2742922787.py:4:8 (compute_block)
f32[64,8192] output of one_minus_square from /tmp/ipykernel_66671/2742922787.py:4:8 (compute_block)
f32[64,8192] output of mul from /tmp/ipykernel_66671/2742922787.py:4:8 (compute_block)
f32[64,8192] output of mul from /tmp/ipykernel_66671/2742922787.py:4:8 (compute_block)
f32[64,512] output of mul from /tmp/ipykernel_66671/2742922787.py:10:20 (loss_fn_610)

—— Residuals saved by the checkpointed version ——
f32[64,512] from the argument x
f32[512,8192] from the argument W1
f32[8192,512] from the argument W2
f32[64,512] output of mul from /tmp/ipykernel_66671/2742922787.py:17:20 (loss_fn_checkpointed)

⚠️ Note:

  1. You can measure peak memory directly. jax.jit(f).lower(*args).compile().memory_analysis() reports the size of the temporary buffers XLA allocates without running anything. Measured that way, the single block above saves about 33% of temporary memory. print_saved_residuals lists the 5 f32[64,8192] residuals that are no longer saved (x @ W1 and the GELU intermediates).
  2. The first call's time is mostly compile time and can't be used to compare run speed, so don't time a single time.time() call that includes compilation. Warm up first, then time several runs and average them (with block_until_ready).

Analyzing the results, and when to use it

  1. Memory: the checkpointed version doesn't need to store the large activation $y$ (batch_size × hidden_dim) for the backward pass; it recomputes it with compute_block during the backward pass. The bigger hidden_dim, the more you save.
  2. Compute time: the checkpointed version should be slower, because the backward pass reruns part of the forward pass. How much slower depends on the relative cost of forward and backward and on the hardware.
  3. Gradient accuracy: mathematically identical, with at most tiny differences from the order of floating-point operations (more noticeable under mixed precision).

When to use it: intermediate activations are large (wide linear layers, Transformer self-attention); recomputation costs less than the memory it saves is worth; memory is tight and you want to train a bigger model or use a bigger batch within a fixed budget. You can apply it to selected layers or blocks only, and finding the balance between memory and compute takes experimentation. The wrapper in Flax is flax.linen.remat (Linen), and in NNX it's nnx.remat.

🔍 Extra: why does a single block save only about 33%?

⚪ Optional: when you want to understand where the benefit of checkpointing actually comes from

In a single-block model, the backward pass starts as soon as the forward pass ends, and the first thing the backward pass does is recompute $y$. The recomputed $y$ and the originally stored $y$ stay alive for almost the same amount of time, so the peak barely drops. The benefit of checkpointing comes from multiple layers: without checkpointing, the activations of $L$ layers have to be kept from the forward pass all the way to the backward pass, so the peak is $\propto L$; with per-layer checkpointing, you only store $L$ layer inputs plus the activations of "the layer currently being recomputed". The full exercise below verifies this with a multi-layer Transformer.

✅ Full practice: a checkpointed deep Transformer

Write a Pre-LN Transformer block in pure JAX (multi-head causal self-attention + GELU MLP), stack the parameters of 8 layers along dim 0, and run the layers one by one with lax.scan. Compare five approaches: no checkpointing, full per-layer recomputation, dots_saveable, dots_with_no_batch_dims_saveable, and saving only the named MLP hidden layer. Each one asserts that its gradient matches no checkpointing, and reports how much temporary memory it saves and how long it takes.

TF = dict(L=8, B=8, T=128, D=256, H=4, F=1024)

def layer_norm(x, g, b, eps=1e-5):
    mu = x.mean(-1, keepdims=True)
    var = ((x - mu) ** 2).mean(-1, keepdims=True)
    return (x - mu) * lax.rsqrt(var + eps) * g + b

def causal_attention(x, p, n_heads):
    B, T, D = x.shape
    hd = D // n_heads
    q, k, v = jnp.split(x @ p["wqkv"], 3, axis=-1)
    q, k, v = (t.reshape(B, T, n_heads, hd) for t in (q, k, v))
    s = jnp.einsum("bthd,bshd->bhts", q, k) / np.sqrt(hd)
    mask = jnp.tril(jnp.ones((T, T), bool))
    a = jax.nn.softmax(jnp.where(mask, s, -1e30), axis=-1)
    o = jnp.einsum("bhts,bshd->bthd", a, v).reshape(B, T, D)
    return o @ p["wo"]

def transformer_block(x, p):
    h = x + causal_attention(layer_norm(x, p["ln1_g"], p["ln1_b"]), p, TF["H"])
    m = jax.nn.gelu(layer_norm(h, p["ln2_g"], p["ln2_b"]) @ p["w1"])
    m = checkpoint_name(m, "mlp_hidden")                  # give the MLP hidden layer a name
    return h + m @ p["w2"]

def init_transformer(key, L, D, F):
    ks = jax.random.split(key, 4)
    n = lambda k, shape, fan_in: jax.random.normal(k, (L,) + shape) / np.sqrt(fan_in)
    return {"wqkv": n(ks[0], (D, 3 * D), D), "wo": n(ks[1], (D, D), D),
            "w1": n(ks[2], (D, F), D), "w2": n(ks[3], (F, D), F),
            "ln1_g": jnp.ones((L, D)), "ln1_b": jnp.zeros((L, D)),
            "ln2_g": jnp.ones((L, D)), "ln2_b": jnp.zeros((L, D))}

def make_tf_loss(remat=False, policy=None):
    blk = jax.checkpoint(transformer_block, policy=policy) if remat else transformer_block
    def loss(params, x, y):
        out, _ = lax.scan(lambda h, p: (blk(h, p), None), x, params)
        return jnp.mean((out - y) ** 2)
    return loss

k_p, k_x, k_y = jax.random.split(jax.random.key(610), 3)
tf_params = init_transformer(k_p, TF["L"], TF["D"], TF["F"])
tf_x = jax.random.normal(k_x, (TF["B"], TF["T"], TF["D"]))
tf_y = jax.random.normal(k_y, (TF["B"], TF["T"], TF["D"]))
print("Parameter count:", sum(x.size for x in jax.tree.leaves(tf_params)), "| parameter bytes:", tree_bytes(tf_params) / MiB, "MiB")
Parameter count: 6299648 | parameter bytes: 24.03125 MiB
tf_variants = {
    "no checkpoint":                        make_tf_loss(False),
    "per-layer full remat (nothing_saveable)": make_tf_loss(True, cp.nothing_saveable),
    "dots_saveable":                      make_tf_loss(True, cp.dots_saveable),
    "dots_with_no_batch_dims_saveable":   make_tf_loss(True, cp.dots_with_no_batch_dims_saveable),
    "save_only_these_names('mlp_hidden')": make_tf_loss(True, cp.save_only_these_names("mlp_hidden")),
}
tf_args = (tf_params, tf_x, tf_y)
g_tf_ref = jax.jit(jax.grad(tf_variants["no checkpoint"]))(*tf_args)
tf_rows = {}
for name, f in tf_variants.items():
    gf = jax.grad(f)
    temp = mem_of(gf, *tf_args).temp_size_in_bytes / MiB
    g = jax.jit(gf)(*tf_args)
    err = max_diff(g, g_tf_ref) / max(float(jnp.abs(l).max()) for l in jax.tree.leaves(g_tf_ref))
    tf_rows[name] = (temp, bench(jax.jit(gf), *tf_args, n=5), err)

base_t, base_ms, _ = tf_rows["no checkpoint"]
print(f"{'policy':38s} {'temp MiB':>10s} {'saved':>7s} {'time ms':>8s} {'ratio':>6s} {'grad rel err':>11s}")
for name, (temp, ms_, err) in tf_rows.items():
    print(f"{name:38s} {temp:10.2f} {1 - temp / base_t:7.1%} {ms_:8.1f} {ms_ / base_ms:6.2f} {err:11.1e}")
    assert err < 1e-5, name                                   # gradient matches no checkpointing

tf_saving = 1 - tf_rows["per-layer full remat (nothing_saveable)"][0] / base_t
print(f"\n8-layer Transformer: per-layer checkpointing saves {tf_saving:.1%} of temp memory (the single-block example saved only {saving_610:.1%})")
assert tf_saving > 0.5 and tf_saving > saving_610
assert tf_rows["per-layer full remat (nothing_saveable)"][0] <= tf_rows["dots_saveable"][0] <= base_t
policy                                   temp MiB   saved  time ms  ratio grad rel err
no checkpoint                              318.92    0.0%    128.1   1.00     0.0e+00
per-layer full remat (nothing_saveable)      50.07   84.3%    110.2   0.86     0.0e+00
dots_saveable                              133.04   58.3%     94.4   0.74     0.0e+00
dots_with_no_batch_dims_saveable           108.04   66.1%    139.8   1.09     0.0e+00
save_only_these_names('mlp_hidden')         81.07   74.6%    275.5   2.15     0.0e+00

8-layer Transformer: per-layer checkpointing saves 84.3% of temp memory (the single-block example saved only 32.7%)
# What does dots_with_no_batch_dims_saveable actually save? Projection matmuls (no batch dim) are saved; the attention einsums (with a batch dim) are recomputed
print_saved_residuals(tf_variants["dots_with_no_batch_dims_saveable"], *tf_args)
f32[8,256] from the argument params['ln1_b']
f32[8,256] from the argument params['ln1_g']
f32[8,256] from the argument params['ln2_b']
f32[8,256] from the argument params['ln2_g']
f32[8,256,1024] from the argument params['w1']
f32[8,1024,256] from the argument params['w2']
f32[8,256,256] from the argument params['wo']
f32[8,256,768] from the argument params['wqkv']
f32[8,8,128,768] output of scan from /tmp/ipykernel_66671/2744042544.py:36:17 (make_tf_loss.<locals>.loss)
f32[8,8,128,256] output of scan from /tmp/ipykernel_66671/2744042544.py:36:17 (make_tf_loss.<locals>.loss)
f32[8,8,128,1024] output of scan from /tmp/ipykernel_66671/2744042544.py:36:17 (make_tf_loss.<locals>.loss)
f32[8,8,128,256] output of scan from /tmp/ipykernel_66671/2744042544.py:36:17 (make_tf_loss.<locals>.loss)
f32[8,128,256] output of mul from /tmp/ipykernel_66671/2744042544.py:37:24 (make_tf_loss.<locals>.loss)

Reading the results:

  • Full per-layer recomputation saves the most (each layer stores only its input [8,128,256]), at the cost of rerunning every layer's forward pass during the backward pass;
  • dots_saveable saves every matmul result (including attention's $QK^\top$ and $AV$) and recomputes the elementwise ops (LayerNorm, GELU, softmax), so its memory lands between the first two. In theory its recomputation also lands in between, but actual CPU timings don't necessarily follow that order (see the note below);
  • dots_with_no_batch_dims_saveable saves only the weight projections (dots like x @ W with no batch dimension) and recomputes the attention einsums that have a batch dimension. This is a common compromise in LLM training, because the $O(T^2)$ attention score matrix is the biggest memory hog;
  • Name-based policies let you save just one specific tensor (here the MLP hidden layer), which is handy for manual tuning based on profiling results.

The timing differences on CPU are for reference only (in the table above, no checkpointing may even be slower than full recomputation, which has to do with memory allocation and cache overhead on CPU). On GPU / TPU, matmuls are faster and memory bandwidth matters more, so the relative ranking of the policies may be different.

✅ The same Transformer in NNX, with per-layer checkpointing via nnx.remat

class NNXBlock(nnx.Module):
    def __init__(self, D, H, F, *, rngs):
        self.n_heads = H
        self.ln1, self.ln2 = nnx.LayerNorm(D, rngs=rngs), nnx.LayerNorm(D, rngs=rngs)
        self.qkv = nnx.Linear(D, 3 * D, use_bias=False, rngs=rngs)
        self.proj = nnx.Linear(D, D, use_bias=False, rngs=rngs)
        self.fc1, self.fc2 = nnx.Linear(D, F, rngs=rngs), nnx.Linear(F, D, rngs=rngs)

    def attn(self, x):
        B, T, D = x.shape
        hd = D // self.n_heads
        q, k, v = (t.reshape(B, T, self.n_heads, hd) for t in jnp.split(self.qkv(x), 3, axis=-1))
        s = jnp.einsum("bthd,bshd->bhts", q, k) / np.sqrt(hd)
        a = jax.nn.softmax(jnp.where(jnp.tril(jnp.ones((T, T), bool)), s, -1e30), axis=-1)
        return self.proj(jnp.einsum("bhts,bshd->bthd", a, v).reshape(B, T, D))

    def __call__(self, x):
        h = x + self.attn(self.ln1(x))
        return h + self.fc2(jax.nn.gelu(self.fc1(self.ln2(h))))

def run_block(block, x):
    return block(x)
run_block_remat = nnx.remat(run_block)              # per-layer checkpointing

class NNXTransformer(nnx.Module):
    def __init__(self, L, D, H, F, *, use_remat, rngs):
        self.blocks = nnx.List([NNXBlock(D, H, F, rngs=rngs) for _ in range(L)])
        self.use_remat = use_remat
    def __call__(self, x):
        f = run_block_remat if self.use_remat else run_block
        for blk in self.blocks:
            x = f(blk, x)
        return x

def nnx_tf_loss(model, x, y):
    return jnp.mean((model(x) - y) ** 2)

nnx_tf = {}
for use_remat in [False, True]:
    m = NNXTransformer(TF["L"], TF["D"], TF["H"], TF["F"], use_remat=use_remat, rngs=nnx.Rngs(0))
    gdef, st = nnx.split(m)
    pure = lambda st, x, y: jax.grad(lambda s: nnx_tf_loss(nnx.merge(gdef, s), x, y))(st)
    temp = mem_of(pure, st, tf_x, tf_y).temp_size_in_bytes / MiB
    grad_step = nnx.jit(nnx.grad(nnx_tf_loss))
    g = grad_step(m, tf_x, tf_y)
    nnx_tf[use_remat] = (temp, bench(grad_step, m, tf_x, tf_y, n=5), g)
    print(f"nnx.remat={use_remat!s:5s}: temp memory {temp:7.2f} MiB, time {nnx_tf[use_remat][1]:6.1f} ms")
nnx_saving = 1 - nnx_tf[True][0] / nnx_tf[False][0]
err = max_diff(nnx_tf[True][2], nnx_tf[False][2])
print(f"NNX per-layer checkpointing: memory saved {nnx_saving:.1%}, max gradient difference {err:.1e}")
assert err < 1e-5 and nnx_saving > 0.5
nnx.remat=False: temp memory  227.19 MiB, time  154.2 ms
nnx.remat=True : temp memory   34.03 MiB, time  187.1 ms
NNX per-layer checkpointing: memory saved 85.0%, max gradient difference 0.0e+00

✍️ Hands-on exercises

The course's exercise for this chapter is the 6.10 practice (implement gradient checkpointing and compare memory, time and gradients), which we've already done in full above and extended to an 8-layer Transformer. Here are 4 more substantial exercises. For each one, write it yourself first, then look at the reference solution.

Exercise 1: gradient accumulation weighted by token count. In a language model, each micro-batch has a different number of valid tokens (because of padding). The loss is defined as "the average loss over all valid tokens in the global batch". Write a lax.scan accumulation function whose result is exactly the same as the gradient computed on the whole batch at once; then verify that "directly averaging the per-micro-batch means" is wrong.

Self-check interface: tok_loss_sum_count(w, x, y, m) returns (the sum of the loss over valid tokens, the number of valid tokens); accum_token_weighted(w, x, y, m, n_micro) accumulates with lax.scan and returns the gradient of the global per-token mean loss. Shapes: x (B, T, D), y and m are (B, T), w (D,).

Exercise 2: $\sqrt L$ grouped checkpointing. A 64-layer residual network (each layer does $x\leftarrow x+\tanh(xW)$) where every $k$ layers form a group wrapped in one jax.checkpoint (activations are saved normally inside a group), i.e., two nested scans. Measure the temporary memory for $k=1,2,4,\dots,64$, find the $k$ that uses the least memory, and verify the gradients match. If each layer input takes $a_{\text{in}}$ and the activations saved inside one layer take $a_{\text{layer}}$, then memory $\approx \frac{L}{k}a_{\text{in}}+k\,a_{\text{layer}}$, and the optimum is $k^{*}=\sqrt{L\,a_{\text{in}}/a_{\text{layer}}}$. When one layer's activations are about the same size as its input, that's the classic $k\sim\sqrt L$.

Self-check interface: grouped_loss(params, x, k=None): params = {"w": (64, D, D)}; k=None means no checkpointing, otherwise every k layers form a group with one jax.checkpoint per group; returns jnp.mean(x_out ** 2).

Exercise 3: FSDP by hand with jax.shard_map. Split the parameters along dim 0 across 8 devices; in the function body, fetch the full weights with lax.all_gather(..., tiled=True) before computing. The batch is also split across devices, and the loss is averaged globally with lax.pmean. Verify: the forward loss matches a single device; the gradient from jax.grad of it also matches a single device, and the gradient is still sharded (the transpose of all_gather is a reduce-scatter).

Self-check interface: fsdp_loss(p, x, y): a function wrapped with jax.shard_map, p = {"w1": (D, H), "w2": (H, D)}, model $\mathrm{mean}\big((\mathrm{relu}(xW_1)W_2-y)^2\big)$; parameters and batch are both sharded with P("fsdp", None), it's called under jax.set_mesh(mesh_fsdp), and it returns a scalar loss.

Exercise 4: why do you have to keep fp32 master weights? Start a parameter at 1.0 and update it for 100 steps with SGD, learning rate $10^{-3}$, and a constant gradient of 1. Compare: (a) storing the parameter directly in bf16; (b) fp32 master weights, cast to bf16 for compute. Explain the result of (a).

Self-check interface: w_bf16: the parameter after 100 steps of scheme (a) (a bf16 scalar); w_master: the fp32 master weight after 100 steps of scheme (b).

# ✍️ Your answer: Exercise 1
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 1
try:
    tok_loss_sum_count, accum_token_weighted
except NameError as e:
    print(f"⏳ Exercise 1: not answered yet: {e}")
else:
    _r = np.random.default_rng(99)
    _B, _T, _D = 24, 10, 5
    _x = jnp.asarray(_r.normal(size=(_B, _T, _D)), jnp.float32)
    _y = jnp.asarray(_r.normal(size=(_B, _T)), jnp.float32)
    _len = np.where(np.arange(_B) % 3 == 0, 1, _T)                  # very different valid lengths: 1 or T
    _m = jnp.asarray(np.arange(_T)[None, :] < _len[:, None], jnp.float32)
    _w = jnp.asarray(_r.normal(size=(_D,)), jnp.float32)
    _s, _n = tok_loss_sum_count(_w, _x, _y, _m)
    assert float(_n) == float(_len.sum()), f"The number of valid tokens should be {_len.sum()}, got {float(_n)}"
    _s1, _ = tok_loss_sum_count(_w, _x[:10], _y[:10], _m[:10])
    _s2, _ = tok_loss_sum_count(_w, _x[10:], _y[10:], _m[10:])
    assert np.isclose(float(_s1 + _s2), float(_s), rtol=1e-5), "The first return value should be the *sum* of the loss (so it can be split by example and added up), not the mean"
    _g_full = jax.grad(lambda w: tok_loss_sum_count(w, _x, _y, _m)[0] / tok_loss_sum_count(w, _x, _y, _m)[1])(_w)
    for _nm in (2, 4, 8):
        _g = accum_token_weighted(_w, _x, _y, _m, n_micro=_nm)
        assert max_diff(_g, _g_full) < 1e-5, f"With n_micro={_nm} the result doesn't match the full-batch gradient (difference {max_diff(_g, _g_full):.2e})"
    assert "scan" in str(jax.make_jaxpr(partial(accum_token_weighted, n_micro=4))(_w, _x, _y, _m)), "The exercise asks you to accumulate with lax.scan"
    print("✅ Exercise 1: self-check passed")
⏳ Exercise 1: not answered yet: name 'tok_loss_sum_count' is not defined
# ✍️ Your answer: Exercise 2
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 2
try:
    grouped_loss
except NameError as e:
    print(f"⏳ Exercise 2: not answered yet: {e}")
else:
    _r = np.random.default_rng(21)
    _D = 128
    _p = {"w": jnp.asarray(_r.normal(size=(64, _D, _D)) / np.sqrt(_D), jnp.float32)}
    _x = jnp.asarray(_r.normal(size=(128, _D)), jnp.float32)
    _xs = np.asarray(_x, np.float64)                              # independent forward reference: 64 layers of x ← x + tanh(x W)
    for _l in range(64):
        _xs = _xs + np.tanh(_xs @ np.asarray(_p["w"][_l], np.float64))
    assert np.isclose(float(grouped_loss(_p, _x)), np.mean(_xs ** 2), rtol=1e-4), "With k=None the loss doesn't match the definition of the 64-layer residual network"
    _g0 = jax.jit(jax.grad(grouped_loss))(_p, _x)
    _mem = {}
    for _k in (1, 8, 64):
        _f = partial(grouped_loss, k=_k)
        assert "checkpoint" in str(jax.make_jaxpr(_f)(_p, _x)) or "remat" in str(jax.make_jaxpr(_f)(_p, _x)), f"With k={_k} you should use jax.checkpoint"
        assert max_diff(jax.jit(jax.grad(_f))(_p, _x), _g0) < 1e-4, f"With k={_k} the gradient doesn't match the one without checkpointing"
        _mem[_k] = mem_of(jax.grad(_f), _p, _x).temp_size_in_bytes
    print({k: f"{v / MiB:.2f} MiB" for k, v in _mem.items()})
    assert _mem[8] < _mem[1] and _mem[8] < _mem[64], "k=8 (≈√L) should use less temporary memory than k=1 and k=64"
    print("✅ Exercise 2: self-check passed")
⏳ Exercise 2: not answered yet: name 'grouped_loss' is not defined
# ✍️ Your answer: Exercise 3
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 3
try:
    fsdp_loss
except NameError as e:
    print(f"⏳ Exercise 3: not answered yet: {e}")
else:
    _r = np.random.default_rng(77)
    _D, _H, _B = 32, 64, 32
    _p = {"w1": jnp.asarray(_r.normal(size=(_D, _H)) * 0.2, jnp.float32), "w2": jnp.asarray(_r.normal(size=(_H, _D)) * 0.2, jnp.float32)}
    _x = jnp.asarray(_r.normal(size=(_B, _D)), jnp.float32)
    _y = jnp.asarray(_r.normal(size=(_B, _D)), jnp.float32)
    _ref = lambda p, x, y: jnp.mean((jax.nn.relu(x @ p["w1"]) @ p["w2"] - y) ** 2)
    _l_ref, _g_ref = jax.value_and_grad(_ref)(_p, _x, _y)
    with jax.set_mesh(mesh_fsdp):
        _ps = {k: jax.device_put(v, P("fsdp", None)) for k, v in _p.items()}
        _xs, _ys = jax.device_put(_x, P("fsdp", None)), jax.device_put(_y, P("fsdp", None))
        _l = jax.jit(fsdp_loss)(_ps, _xs, _ys)
        _g = jax.jit(jax.grad(fsdp_loss))(_ps, _xs, _ys)
        _hlo = jax.jit(jax.grad(fsdp_loss)).lower(_ps, _xs, _ys).compile().as_text()
    assert np.isclose(float(_l), float(_l_ref), rtol=1e-5), f"Loss {float(_l):.6f} doesn't match single-device {float(_l_ref):.6f} (the per-device local losses need a global average)"
    assert max_diff(_g, _g_ref) < 1e-5, f"Gradient doesn't match single-device (difference {max_diff(_g, _g_ref):.2e})"
    for _k in ("w1", "w2"):
        assert _g[_k].sharding.spec == P("fsdp", None), f"Gradient {_k} should stay sharded along dim 0, got {_g[_k].sharding.spec}"
    _cc = collectives(_hlo)
    assert _cc["all-gather"] >= 2, f"The forward pass should all_gather both weights: {_cc}"
    assert _cc["reduce-scatter"] >= 2, f"A reduce-scatter (the transpose of all_gather) should appear in the backward pass: {_cc}"
    print("✅ Exercise 3: self-check passed")
⏳ Exercise 3: not answered yet: name 'fsdp_loss' is not defined
# ✍️ Your answer: Exercise 4
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 4
try:
    w_bf16, w_master
except NameError as e:
    print(f"⏳ Exercise 4: not answered yet: {e}")
else:
    assert jnp.asarray(w_bf16).dtype == jnp.bfloat16, f"In scheme (a) the parameter should be stored in bf16 throughout; got dtype {jnp.asarray(w_bf16).dtype}"
    assert jnp.asarray(w_master).dtype == jnp.float32, f"In scheme (b) the master weight should be fp32; got dtype {jnp.asarray(w_master).dtype}"
    assert np.isclose(float(w_master), 1.0 - 100 * 1e-3, atol=1e-5), f"The fp32 master weight should be about 0.9 after 100 steps, got {float(w_master)}"
    assert float(w_bf16) == 1.0, f"The bf16 parameter should still be 1.0 after 100 steps (every update gets rounded away), got {float(w_bf16)}"
    print("✅ Exercise 4: self-check passed")
⏳ Exercise 4: not answered yet: name 'w_bf16' is not defined

Explanation for Exercise 4: bf16 has only 8 significant binary digits, so near 1.0 neighboring representable numbers are $2^{-8}\approx0.0039$ apart. Each update of $10^{-3}$ is less than half that spacing, so 1.0 - 0.001 rounds back to 1.0, and after 100 steps the parameter hasn't changed at all. fp32 master weights can accumulate these small updates ($1-100\times10^{-3}=0.9$), and you just cast to bf16 for compute. That's why mixed precision has to be "fp32 parameters, bf16 compute".


📝 Quiz

8 questions below; answers and explanations are in the collapsible section at the end.

1. A model with 1 billion parameters is trained with fp32 parameters + fp32 gradients + Adam (fp32 $m$ and $v$). Ignoring activations, what's the minimum memory needed?

  • A. 4 GB · B. 8 GB · C. 12 GB · D. 16 GB

2. Which of the following does not reduce peak activation memory?

  • A. Applying jax.checkpoint to every Transformer layer
  • B. Wrapping the whole model's loss function in one jax.checkpoint
  • C. Using gradient accumulation to split the batch into 4 micro-batches
  • D. Computing activations in bf16

3. Which statement about jax.checkpoint(f, policy=jax.checkpoint_policies.dots_saveable) is correct?

  • A. It saves all intermediate results, equivalent to no checkpointing
  • B. It saves only the outputs of matmuls (dot_general); elementwise ops are recomputed in the backward pass
  • C. It saves only values tagged with checkpoint_name
  • D. It offloads residuals to host memory

4. In NNX in Flax 0.12.12, which of these lines raises a TypeError?

  • A. nnx.Optimizer(model, optax.adam(1e-3), wrt=nnx.Param)
  • B. nnx.Optimizer(model, optax.adam(1e-3))
  • C. optimizer.update(model, grads)
  • D. nnx.state(model, nnx.BatchStat)

5. With fp32 parameters and bf16 compute (matmul after astype(bfloat16)), what is the dtype of jax.grad with respect to the parameters?

  • A. bfloat16 · B. float16 · C. float32 · D. depends on preferred_element_type

6. A Megatron-style MLP on a ('data', 'model') 2D mesh: $W_1$ is P(None, 'model'), $W_2$ is P('model', None), the input is P('data', None). How many rounds of cross-device communication does the forward pass need?

  • A. 0 · B. 1 all-reduce (after the second matmul) · C. 2 all-gathers · D. one all-reduce per matmul

7. Which statement about FSDP (ZeRO-3) is wrong?

  • A. Parameters, gradients and optimizer state are all sharded across devices
  • B. When the forward pass reaches a layer, it first all-gathers the full weights
  • C. Each device computes only some of the model's layers
  • D. In JAX it can be implemented by annotating parameters with P('fsdp', None) and calling reshard(w, P()) before use

8. You train with optax.MultiSteps(optax.adam(1e-3), every_k_schedule=4), calling update once per micro-batch. Which statement is correct?

  • A. Every call updates the parameters, just with the learning rate divided by 4
  • B. The first 3 calls return all-zero updates; the 4th does one Adam update with the average gradient of the 4 micro-batches
  • C. You need to divide the gradient by 4 by hand
  • D. The optimizer state doesn't change, since it's just a scheduler
Click to expand: answers and explanations
  1. D. Per parameter 4 (params) + 4 (gradients) + 8 ($m$, $v$) = 16 bytes, so $10^9\times16$ B = 16 GB (section 6.1).
  2. B. With only one checkpoint around the whole model, the backward pass starts by recomputing the entire forward pass and saves all activations, so the peak is the same as without checkpointing (the note in section 6.6 measured 83.25 MiB vs 83.25 MiB). A, per-layer checkpointing, works; C holds only one micro-batch's activations at a time; D halves the bytes per activation.
  3. B. A is everything_saveable, C is save_only_these_names, D is the offload_* family (the policy table in the 6.6 Extra).
  4. B. Since Flax 0.11, wrt is a required keyword argument; leaving it out raises TypeError: Missing required argument wrt (6.2 Extra).
  5. C. The transpose of astype turns the cotangent back into the parameter's dtype, so parameter gradients are fp32; only the activation gradients flowing through the backward pass are bf16 (the note in section 6.7).
  6. B. The column-split first layer needs no communication; the row-split second layer produces partial sums that one all-reduce adds up (section 6.8 counted all-reduce: 1).
  7. C. In FSDP every device still computes all layers (on its own data shard); the parameters are just stored sharded and all-gathered when needed. "Each device computes only some layers" is pipeline parallelism.
  8. B. MultiSteps only accumulates for the first $k-1$ calls and updates with the average gradient (use_grad_mean=True) on the $k$-th, matching one Adam update on the large batch (section 6.5 measured a parameter difference $<10^{-6}$); its state carries an extra accumulation buffer.
# Verify quiz questions 1, 4 and 5
print("Q1:", 1e9 * (4 + 4 + 8) / 1e9, "GB")
try:
    nnx.Optimizer(nnx.Linear(2, 2, rngs=nnx.Rngs(0)), optax.adam(1e-3))
except TypeError as e:
    print("Q4: B →", type(e).__name__)
w_q5 = jnp.ones((3, 2))
print("Q5:", jax.grad(lambda w: jnp.sum(jnp.ones((1, 3), jnp.bfloat16) @ w.astype(jnp.bfloat16)).astype(jnp.float32))(w_q5).dtype)
assert jax.grad(lambda w: jnp.sum(jnp.ones((1, 3), jnp.bfloat16) @ w.astype(jnp.bfloat16)).astype(jnp.float32))(w_q5).dtype == jnp.float32
Q1: 16.0 GB
Q4: B → TypeError
Q5: float32

Version notes

This table collects the patterns and conclusions in this chapter worth knowing under JAX 0.12.0.1 / Flax 0.12.12 / Optax 0.2.9 (CPU). Apart from hardware specs and timings, each item is verified with an assert at the corresponding place in the text.

No.SectionKey pointWhere
16.1Adam/AdamW optimizer state is exactly 2× the parameters ($m$, $v$); it only reaches 3× if you also count the fp32 master weights of mixed precisionafter the 6.1 memory-estimate cell
26.2Flax's recommended API is NNX (Linen is the old API, but not deprecated); Haiku has been in maintenance mode since 2023-07, and DeepMind recommends switching to Flax. The main thread is written in NNX, with a Linen / NNX / Haiku comparison tablenote at the start of 6.2 + comparison table
36.2/6.3Replace jax.random.PRNGKey and jax.tree_util.tree_map with jax.random.key and jax.tree.map; in NNX, a Variable's .value is deprecated (DeprecationWarning), use var[...] instead6.2 Extra
46.3The packed training state (params + Adam state + key + step) has 21 leaves: params 6 + Adam state 13 (count 1 + mu 6 + nu 6) + key 1 + step 16.3 PyTree example
56.3jax.device_put_replicated was removed in 0.12 (AttributeError); after restoring a checkpoint, use jax.device_put + NamedSharding, or let Orbax restore directly from a sharded abstract state6.3 checkpoints
66.4jax.device_put_replicate is a typo, and the correctly spelled device_put_replicated has been removed too; replicate model state with jax.device_put(stacked_array, NamedSharding(...)) instead6.4 replicating model state
76.4Data parallelism no longer needs pmap + manually reshaping out a device dim + lax.pmean: use jit + NamedSharding (batch dim P('data')), the compiler inserts the gradient all-reduce automatically, and the result matches a single device6.4 modern style
86.5With a mean loss, equal-size micro-batches and no BatchNorm, the averaged accumulated gradient is exactly equal to the full-batch gradient (difference $\sim10^{-7}$); only a summed loss, unequal micro-batches or BN make a differencenote at the start of 6.5 + Extra
96.5Defining @jax.jit def apply_update inside accumulated_train_step creates a new function object on every call, so 5 steps retrace 5 times; define the jit function at the outer level, or put the whole accumulation in one jit + lax.scan (for the full table of recompilation triggers see Basics 02 JIT compilation §2.3)6.5 course implementation
106.6w1 and w2 must not be generated from the same key: they have the same number of elements, so flattened they are exactly the same random sequence; split first6.6 course example
116.6Wrapping the whole model in one checkpoint doesn't lower peak memory (83.25 vs 83.25 MiB); only per-layer checkpointing helps (37 MiB). FLOPs go up by at most about +33% (one extra forward pass, $3F\to4F$), never double: per-layer +16.8%; one whole-model checkpoint with only grad +0% (the original forward pass is removed by DCE), about +33% with value_and_grad6.6 note
126.6Advanced checkpoint tools: all the jax.checkpoint_policies available in 0.12, checkpoint_name, print_saved_residuals (needs an explicit import jax.ad_checkpoint), memory_analysis(), nnx.remat6.6 Extra
136.7Under mixed precision, differentiating w.r.t. fp32 parameters gives fp32 gradients (same in Linen and NNX); what gets halved is the memory for activations and activation gradients6.7 note
146.7Optax 0.2.9 has no loss-scaling tool, only apply_if_finite / skip_not_finite; Flax has flax.training.dynamic_scale.DynamicScale; this chapter writes dynamic loss scaling by hand6.7 loss scaling
156.8Tensor parallelism, FSDP and their combinations are all a multi-dimensional Mesh + PartitionSpec annotations, with jit inserting communication automatically; no pmap over device subsets or manual device_put placement. Use jax.shard_map for hand-written communication. In 0.12, make_mesh defaults to Explicit axes, and a matmul whose contracting dim is split needs out_sharding6.8 note
166.9optax.warmup_cosine_decay_schedule(...) with an ellipsis placeholder fails immediately with TypeError; give init_value, peak_value, warmup_steps, decay_steps6.9 course code
176.9Optax has no sharding logic; where optimizer state lives simply follows the parameters: in pmap the parameters are replicated, so the state is a full copy per device too; the state is sharded only when the parameters are6.9 note
186.10jit(f).lower(...).compile().memory_analysis() gives temporary memory directly (32.7% saved for a single block, 84.3% for the 8-layer Transformer); to compare speed, don't use a single time.time() that includes compilation; warm up, block_until_ready, and average several runs6.10 note
196.1Single-accelerator HBM is no longer limited to 16–80 GB: newer GPUs reach 141–192 GB (H200 / B200 / MI300X, hardware specs, not measured); but the 120 GB of static memory for 10 billion parameters + Adam is still close to or beyond single-device capacity, so the conclusion stands6.1 memory limits

📌 Chapter summary

  • Three kinds of bottleneck: memory (params + gradients + optimizer state + activations + workspace), compute ($\approx6\Phi T$ FLOPs), communication (data parallelism all-reduces about 2× the model size every step). Mixed-precision Adam takes about 16 B/param, so mixed precision alone can't lower static memory.
  • Ecosystem libraries: Flax recommends NNX. Modules are mutable objects; nnx.Param / nnx.BatchStat / custom nnx.Variable distinguish kinds of state; nnx.split/merge/state turn them into pure PyTrees at any time; nnx.Optimizer(model, tx, wrt=nnx.Param) + optimizer.update(model, grads). Haiku is in maintenance mode.
  • Optax: chain(clip_by_global_norm, adamw(schedule)) is the standard setup for large models; MultiSteps builds gradient accumulation into the optimizer; state sharding follows the parameters.
  • Data parallelism: the global view of jit + NamedSharding replaces pmap. Data uses P('data'), parameters without annotations are replicated, and the compiler inserts the all-reduce automatically; you need just one PRNG key.
  • Gradient accumulation: iterate over micro-batches with lax.scan; with a mean loss and equal-size micro-batches it matches the large-batch gradient exactly; a summed loss or unequal micro-batches need weighting, and BatchNorm can't be made equivalent.
  • Gradient checkpointing: do it per layer (or every $k$ layers); one checkpoint around the whole model doesn't lower the peak. Use policies to trade memory against compute (dots_saveable, dots_with_no_batch_dims_saveable, save_only_these_names), and verify quantitatively with memory_analysis().temp_size_in_bytes and print_saved_residuals. The FLOP overhead is at most about +33% (one extra forward pass).
  • Mixed precision: bf16 has the same dynamic range as fp32 but low precision; fp16 has higher precision but a narrow range and needs loss scaling (dynamic scaling + skipping on overflow). fp32 parameters, bf16 compute, accumulate with preferred_element_type=float32; parameter gradients are fp32.
  • Model parallelism: Megatron MLP = column split (no communication) + row split (one all-reduce); FSDP = sharded params/gradients/optimizer state + all-gather on use; the pipeline-parallel bubble fraction is $(p-1)/(m+p-1)$. When verifying sharded correctness, remember that Adam amplifies rounding differences.
  • Optimizers: AdamW is the default; LAMB / LARS target very large batches; Adafactor uses factorization to cut a matrix's second moment from $nm$ to $n+m$; Lion has only one copy of momentum; mu_dtype=bf16 saves 25% of state.

🗂️ Cheat sheet

APIWhat it does
nnx.Module / nnx.Rngs(params=0, dropout=1)define an NNX module / named random streams
nnx.Param / nnx.BatchStat / class X(nnx.Variable)trainable parameters / BN statistics / custom state
var[...], var[...] = vread/write a Variable (replaces the deprecated .value)
nnx.split(m, nnx.Param, ...) / nnx.merge(gdef, state)split into GraphDef + State / merge back into a module
nnx.state(m, filter) / nnx.update(m, state) / nnx.to_pure_dictget state / write state back / convert to a plain dict
nnx.Optimizer(m, tx, wrt=nnx.Param), opt.update(m, grads)NNX optimizer (wrt is required)
nnx.jit / nnx.grad / nnx.value_and_grad / nnx.rematNNX lifted transformations (write module state back automatically)
nnx.with_partitioning(init, (None, 'model'))add a sharding annotation to a parameter (construct inside jax.set_mesh)
nnx.eval_shape(lambda: Model(...))build only an abstract model (checkpoint restore template)
optax.chain(optax.clip_by_global_norm(1.0), optax.adamw(sched, weight_decay=...))clipping + AdamW
optax.warmup_cosine_decay_schedule(init, peak, warmup, decay, end)warmup + cosine decay
optax.MultiSteps(opt, every_k_schedule=k)gradient accumulation
optax.adafactor / lamb / lars / lion / adamw(mu_dtype=bf16)optimizers for scale / saving memory
optax.apply_if_finite(opt, max_consecutive_errors)skip inf/NaN updates
jax.make_mesh((2, 4), ('data', 'model'))create a device mesh (Explicit axes by default in 0.12)
jax.set_mesh(mesh) / jax.device_put(x, P('data', None))set the current mesh / place an array by PartitionSpec
jax.typeof(x) / jax.debug.visualize_array_sharding(x)show the type including sharding (f32[8@data,4]) / visualize sharding
jnp.matmul(a, b, out_sharding=P(...))specify the output sharding in Explicit mode (partial sums in row parallelism)
jax.sharding.reshard(x, P())change the sharding (e.g., the all-gather in FSDP)
jax.shard_map(f, in_specs, out_specs) + lax.psum/all_gather/pmeanper-device view, hand-written communication
jax.checkpoint(f, policy=..., prevent_cse=..., static_argnums=...)gradient checkpointing (jax.remat is an alias)
jax.checkpoint_policies.dots_saveable, etc.choose which residuals to save
jax.ad_checkpoint.checkpoint_name(x, 'name')name an intermediate value, for use with save_only_these_names
jax.ad_checkpoint.print_saved_residuals(f, *args)print the residuals the backward pass will save
jax.jit(f).lower(*a).compile().memory_analysis()compile-time memory analysis (temp_size_in_bytes)
jax.jit(f).lower(*a).compile().cost_analysis()['flops']compile-time FLOP count
jnp.dot(a, b, preferred_element_type=jnp.float32)low-precision inputs, fp32 accumulation
jnp.finfo(dtype)inspect a float format's bits, eps and range
orbax.checkpoint.StandardCheckpointer().save/restoresave / restore a training-state PyTree

🧭 Where to go next

TopicResourceNotes
Scaling overviewHow to Scale Your Model (the official JAX scaling book)From rooflines and TPU/GPU architecture to the communication costs of data/tensor/pipeline/expert parallelism; the quantitative version of every parallel strategy in this chapter
Sharding and parallelismJAX docs Distributed arrays and automatic parallelization, Explicit sharding, Manual parallelism with shard_map, Migrating from pmapThe authoritative explanation of jit + NamedSharding / shard_map in 0.12
Checkpointing (rematerialization)JAX docs Gradient checkpointing with jax.checkpointPolicies, checkpoint_name, offloading to host memory
Training checkpointsOrbax (orbax.checkpoint)Asynchronous saving, multi-host sharded checkpoints, retention policies with CheckpointManager
Data loadingGrainDeterministic, checkpointable (resumable) data pipelines for the JAX ecosystem, with multi-host sharding
Neural network libraryFlax NNX docs (including Scale up on multiple devices and the Linen to NNX migration guide)NNX sharding annotations, stacking layers with nnx.scan, logical axis rules
OptimizersOptax docsOptimizers, schedules, gradient transformations, newer methods in optax.contrib (Muon, Schedule-Free, etc.)
Reference implementationsMaxText, EasyLM, the JAX implementations of Gemma / PaliGemmaProduction LLM training code: logical-axis sharding rules, scan + remat, mixed precision, Orbax and Grain all put together
Profilingjax.profiler, TensorBoard / Perfetto, XLA dumpsFind activation-memory and communication bottlenecks, then decide on the checkpointing policy and parallelism

Appendix: reference solutions

Write your own answer in the "✍️ Your answer" cells first and check it with the "🧪 Self-check"; come here to compare once the self-check passes or if you're really stuck. Running the solutions below in order overwrites any functions with the same names you defined in the answer cells.

Exercise 1

See "✍️ Hands-on exercises" in the main text for the problem.

# Reference solution: Exercise 1
r = np.random.default_rng(11)
NB, T_, Dm = 32, 12, 8
tok_x = jnp.asarray(r.normal(size=(NB, T_, Dm)), jnp.float32)
tok_y = jnp.asarray(r.normal(size=(NB, T_)), jnp.float32)
lengths = r.integers(1, T_ + 1, size=NB)
tok_mask = jnp.asarray(np.arange(T_)[None, :] < lengths[:, None], jnp.float32)   # 1 for valid tokens
tok_w = jnp.asarray(r.normal(size=(Dm,)), jnp.float32)

def tok_loss_sum_count(w, x, y, m):
    per_tok = (jnp.tanh(x @ w) - y) ** 2
    return jnp.sum(per_tok * m), jnp.sum(m)          # returns (loss sum, number of valid tokens)

def global_mean_loss(w, x, y, m):
    s, n = tok_loss_sum_count(w, x, y, m)
    return s / n

@partial(jax.jit, static_argnames="n_micro")
def accum_token_weighted(w, x, y, m, n_micro):
    split = lambda a: a.reshape(n_micro, -1, *a.shape[1:])
    def body(carry, mb):
        g_sum, n_sum = carry
        # differentiate the "loss sum" (has_aux carries out the token count); accumulate gradients and token counts
        g, n = jax.grad(tok_loss_sum_count, has_aux=True)(w, *mb)
        return (g_sum + g, n_sum + n), None
    (g_sum, n_sum), _ = lax.scan(body, (jnp.zeros_like(w), 0.0), (split(x), split(y), split(m)))
    return g_sum / n_sum                                                 # divide by the total token count at the end

g_full = jax.grad(global_mean_loss)(tok_w, tok_x, tok_y, tok_mask)
g_acc = accum_token_weighted(tok_w, tok_x, tok_y, tok_mask, n_micro=4)
# The wrong way: take each micro-batch's per-token mean, then a plain average
mb_slices = [slice(i * 8, (i + 1) * 8) for i in range(4)]
g_wrong = sum(jax.grad(global_mean_loss)(tok_w, tok_x[sl], tok_y[sl], tok_mask[sl]) for sl in mb_slices) / 4
print(f"Token-weighted accumulation vs full batch: {max_diff(g_acc, g_full):.2e}  |  plain average: {max_diff(g_wrong, g_full):.2e}")
assert max_diff(g_acc, g_full) < 1e-6 and max_diff(g_wrong, g_full) > 1e-4
Token-weighted accumulation vs full batch: 4.47e-08  |  plain average: 1.02e-02

Exercise 2

See "✍️ Hands-on exercises" in the main text for the problem.

# Reference solution: Exercise 2
L2, B2, D2 = 64, 256, 256
r = np.random.default_rng(12)
ex2_p = {"w": jnp.asarray(r.normal(size=(L2, D2, D2)) / np.sqrt(D2), jnp.float32)}
ex2_x = jnp.asarray(r.normal(size=(B2, D2)), jnp.float32)

def one_layer(x, p):
    return x + jnp.tanh(x @ p["w"]), None          # each layer stores x@w and the tanh output, about 2–3 × a_in

def grouped_loss(params, x, k=None):
    if k is None:                                   # no checkpointing
        x, _ = lax.scan(one_layer, x, params)
        return jnp.mean(x ** 2)
    grouped = jax.tree.map(lambda a: a.reshape(L2 // k, k, *a.shape[1:]), params)
    @jax.checkpoint                                 # one checkpoint per group: store only the group's input
    def group(x, gp):
        x, _ = lax.scan(one_layer, x, gp)           # activations are saved normally inside a group
        return x
    x, _ = lax.scan(lambda x, gp: (group(x, gp), None), x, grouped)
    return jnp.mean(x ** 2)

g_ref2 = jax.jit(jax.grad(grouped_loss))(ex2_p, ex2_x)
mem2 = {"none": mem_of(jax.grad(grouped_loss), ex2_p, ex2_x).temp_size_in_bytes / MiB}
for k in [1, 2, 4, 8, 16, 32, 64]:
    f = partial(grouped_loss, k=k)
    mem2[k] = mem_of(jax.grad(f), ex2_p, ex2_x).temp_size_in_bytes / MiB
    assert max_diff(jax.jit(jax.grad(f))(ex2_p, ex2_x), g_ref2) < 1e-4
for k, v in mem2.items():
    print(f"k = {str(k):>2s}: temp memory {v:7.2f} MiB")
best_k = min((k for k in mem2 if k != "none"), key=lambda k: mem2[k])
print(f"The k that uses the least memory = {best_k} (√L = {np.sqrt(L2):.0f}; k=L means one checkpoint around everything, about the same as no checkpointing)")
assert 1 < best_k < L2 and mem2[best_k] < mem2[1] and mem2[best_k] < mem2[L2]
assert mem2[L2] >= 0.95 * mem2["none"]
k = none: temp memory   33.25 MiB
k =  1: temp memory   17.25 MiB
k =  2: temp memory   11.25 MiB
k =  4: temp memory    9.25 MiB
k =  8: temp memory   11.25 MiB
k = 16: temp memory   18.25 MiB
k = 32: temp memory   33.75 MiB
k = 64: temp memory   33.25 MiB
The k that uses the least memory = 4 (√L = 8; k=L means one checkpoint around everything, about the same as no checkpointing)

Exercise 3

See "✍️ Hands-on exercises" in the main text for the problem.

# Reference solution: Exercise 3
D3, H3, B3 = 64, 256, 64
r = np.random.default_rng(13)
ex3_p = {"w1": jnp.asarray(r.normal(size=(D3, H3)) * 0.1, jnp.float32),
         "w2": jnp.asarray(r.normal(size=(H3, D3)) * 0.1, jnp.float32)}
ex3_x = jnp.asarray(r.normal(size=(B3, D3)), jnp.float32)
ex3_y = jnp.asarray(r.normal(size=(B3, D3)), jnp.float32)

def single_loss(p, x, y):
    return jnp.mean((jax.nn.relu(x @ p["w1"]) @ p["w2"] - y) ** 2)

def fsdp_local_loss(p_blk, x_blk, y_blk):
    w1 = lax.all_gather(p_blk["w1"], "fsdp", tiled=True)     # (D/8, H) → (D, H)
    w2 = lax.all_gather(p_blk["w2"], "fsdp", tiled=True)     # (H/8, D) → (H, D)
    local = jnp.mean((jax.nn.relu(x_blk @ w1) @ w2 - y_blk) ** 2)
    return lax.pmean(local, "fsdp")                          # every device has an equal-size batch → global mean

with jax.set_mesh(mesh_fsdp):
    spec_p = {"w1": P("fsdp", None), "w2": P("fsdp", None)}
    fsdp_loss = jax.shard_map(fsdp_local_loss, in_specs=(spec_p, P("fsdp", None), P("fsdp", None)), out_specs=P())
    p3 = {k: jax.device_put(v, spec_p[k]) for k, v in ex3_p.items()}
    x3, y3 = jax.device_put(ex3_x, P("fsdp", None)), jax.device_put(ex3_y, P("fsdp", None))
    l3 = jax.jit(fsdp_loss)(p3, x3, y3)
    g3 = jax.jit(jax.grad(fsdp_loss))(p3, x3, y3)
    print("shard_map FSDP gradients:", jax.typeof(g3["w1"]), jax.typeof(g3["w2"]))
    hlo3 = jax.jit(jax.grad(fsdp_loss)).lower(p3, x3, y3).compile().as_text()
l_ref3 = single_loss(ex3_p, ex3_x, ex3_y)
g_ref3 = jax.grad(single_loss)(ex3_p, ex3_x, ex3_y)
print(f"Loss {float(l3):.6f} vs {float(l_ref3):.6f}; max gradient difference {max_diff(g3, g_ref3):.2e}; communication {collectives(hlo3)}")
assert np.isclose(float(l3), float(l_ref3), rtol=1e-6) and max_diff(g3, g_ref3) < 1e-6
assert g3["w1"].sharding.spec == P("fsdp", None)          # gradients are still sharded
shard_map FSDP gradients: float32[64@fsdp,256] float32[256@fsdp,64]
Loss 1.798891 vs 1.798891; max gradient difference 7.45e-09; communication {'all-reduce': 0, 'all-gather': 2, 'reduce-scatter': 2, 'all-to-all': 0, 'collective-permute': 0}

Exercise 4

See "✍️ Hands-on exercises" in the main text for the problem.

# Reference solution: Exercise 4
lr4, steps4 = 1e-3, 100
w_bf16 = jnp.bfloat16(1.0)
w_master = jnp.float32(1.0)
for _ in range(steps4):
    w_bf16 = (w_bf16 - jnp.bfloat16(lr4) * jnp.bfloat16(1.0)).astype(jnp.bfloat16)   # (a) parameter stored in bf16
    w_master = w_master - lr4 * jnp.float32(1.0)                                        # (b) fp32 master weight
print(f"(a) bf16 parameter: {float(w_bf16):.4f}   (b) fp32 master weight: {float(w_master):.4f} (as bf16 for compute: {float(w_master.astype(jnp.bfloat16)):.4f})")
print(f"bf16 spacing at 1.0 = 2^-7 = {2**-7:.4f}, just below 1.0 = 2^-8 = {2**-8:.4f}; a single update of {lr4} < half the spacing → it gets rounded back to the old value every time")
assert float(w_bf16) == 1.0                         # all 100 updates are lost
assert np.isclose(float(w_master), 0.9, atol=1e-5)
(a) bf16 parameter: 1.0000   (b) fp32 master weight: 0.9000 (as bf16 for compute: 0.8984)
bf16 spacing at 1.0 = 2^-7 = 0.0078, just below 1.0 = 2^-8 = 0.0039; a single update of 0.001 < half the spacing → it gets rounded back to the old value every time