Repository navigation
feat(training_engine): add prepare_fwd_bwd_input and wait_for_device to overlap the next step's input copy (-7.2% s/step, 397B RL trainer bench, 128 v7x) - #5718
Conversation
…ts fwd_bwd prepare_fwd_bwd_input(payload) maps a micro-batch to loss inputs and device_puts the traced half onto the shardings fwd_bwd compiles them onto (_batch_data_shardings), leaving the closed-over half untouched; fwd_bwd(PreparedFwdBwdInput) then dispatches without moving data. Tunix's TrainerWorker.stage_fwd_bwd calls it when the RL loop sends the next step's first micro-batch while the device still computes the previous step's last micro-batch and update. Costs one micro-batch of inputs in device memory until it is trained (~20 MB per device for a [64 x 65,536] router-replay batch). Tests: prepared micro-batches train bit-identically (weights and metrics, compiled and eager, no recompile), also when staged through a real Tunix TrainerWorker; a one-token change is seen; preparing changes no engine state.
wait_for_device() blocks until every computation dispatched so far has finished and returns the seconds it waited. Tunix's TrainerWorker calls it after staging the next step's first micro-batch (prepare_fwd_bwd_input): a wait of ~0 means the device was already idle, so the transfer did not hide behind the step's last micro-batch and update, and the RL loop stops staging. It is a blocking wait rather than an is_ready() poll because under Pathways a readiness check is an RPC to the IFRT proxy, so a poll can read "not ready" for work that has finished. InflightThrottler.block_until_ready() does the waiting without dequeuing and without writing metrics, so the next wait_for_next / wait_for_all, and the metrics it writes, are what they would have been. Tests: training with a wait after each prepared micro-batch is bit-identical to training without (weights, metrics, compile count, metric writes in the same order); the wait leaves the in-flight queue as it was and returns with every queued leaf ready; block_until_ready writes nothing and a drain after it writes the same metrics. The review's 4-device test with Tunix's real model input (a nested RLTrainerPayload, batch split four ways) is ported, with the wait added.
… the fused update Brings prepare_fwd_bwd_input / wait_for_device up to the current engine (per-micro-batch dropless replay, the final micro-batch's dropless program, Tunix's fused final update) and makes their tests hold there: - maxtext_engine_prepared_input_dropless_test.py (new): on the one-layer Qwen3.5 MoE rig of the dropless tests, a prepared micro-batch that overflows is replayed from its prepared inputs, the step's last micro-batch runs the final program, and weights, losses, gradient norms and replay counts are bit-identical to the unprepared run, whether it is prepared before or after the update; also when the run's first micro-batch is prepared. A control that prepares the wrong micro-batch is caught. - maxtext_engine_prepared_input_test.py: prepares before and after the update (the second is where Tunix's stage lands once the update runs inside the last train call); the TrainerWorker test drives fwd_bwd(apply_optimizer=True, wait_for_update=False) then stage_fwd_bwd, and skips when the installed Tunix has no staging; wait_for_device is checked to wait on every queued computation, the update's gradient norm included, by recording what it blocked on rather than by is_ready() after the fact. - maxtext_engine_prepared_input_sharded_test.py: re-execs itself with four host devices, as maxtext_engine_data_parallel_test.py does; setting XLA_FLAGS at import time left one device under pytest and skipped green. - README: the two methods and the tests.
…update Tunix's staged step boundary can stage the next step's first micro-batch between a step's last fwd_bwd and its update: fwd_bwd(last, apply_optimizer=True), prepare_fwd_bwd_input(next), wait_for_device(), update(). Under retry_when_tokens_dropped without retry_dropless_final_micro_batch (the MLPerf recipes) the last micro-batch can overflow; its flag is still pending at the prepare and its dropless rerun runs inside update(), after it. maxtext_engine_stage_before_update_test.py, on the one-layer MoE rig of the prepared-input dropless test with the final dropless program off: the staged order trains bit-identically to fwd_bwd(last), update(), fwd_bwd(next payload) (weights, losses, gradient norms, replay counts) with the last micro-batch's rerun inside update() at two boundaries and the prepared micro-batch itself replayed at one; the flag is still pending when the next micro-batch is prepared and wait_for_device resolves nothing; no kernel is rebuilt. Also with the final dropless program on. Two controls are caught: a rerun that runs on the prepared micro-batch instead of its own, and preparing the wrong micro-batch.
…TrainerWorker The same order as the engine-level test, through a real TrainerWorker: fwd_bwd(last, apply_optimizer=True, defer_update=True), stage_fwd_bwd(next), apply_deferred_update(), fwd_bwd_staged(next), against the fused calls. It trains bit-identically (weights, losses, gradient norms) with the last micro-batch's rerun inside the deferred update, defer_update never reaches the engine, and the stage reports no device wait (the update call does). Skipped when the installed Tunix has no apply_deferred_update.
… alone Docstrings and README only, no code change. Tunix stages a step's first micro-batch between the previous step's last fwd_bwd and its update() when it defers the update. Under retry_when_tokens_dropped the last micro-batch's overflow flag is still pending at that point. State the contract this order relies on, which maxtext_engine_stage_before_update_test.py pins: - prepare_fwd_bwd_input neither reads nor resolves the pending check. The dropless rerun that update() may run reads the inputs that check kept, never the prepared ones. The mapping reads only the payload, so the prepared input is the same whether the update runs before or after it. - wait_for_device leaves a pending overflow check, like the in-flight queue and its metrics, to the next fwd_bwd or update(). - README: the same note under "One optimizer step", and the test file in the tests table. - Two test docstrings no longer say the MLPerf recipes set retry_dropless_final_micro_batch, or disagree about it: the team's trainer configuration sets retry_when_tokens_dropped, and the final dropless program is an option on top of it.
There was a problem hiding this comment.
Code Review
This pull request introduces prepare_fwd_bwd_input and wait_for_device to the MaxTextTrainingEngine to support staging micro-batches and pipelining step boundaries, particularly for Tunix. It also adds block_until_ready to the InflightThrottler to wait for inflight computations without dequeuing them or writing metrics, accompanied by comprehensive unit tests. The reviewer suggests defensively mapping the device placement over dynamic batch leaves in prepare_fwd_bwd_input to robustly handle edge cases like zero-width or non-fully-addressable arrays in multi-controller JAX environments.
| batch = self._prepare_batch(payload) | ||
| if self._mesh is not None and isinstance(batch, dict): | ||
| dynamic_batch, static_batch = _split_static_and_dynamic(batch) | ||
| dynamic_batch = jax.device_put(dynamic_batch, self._batch_data_shardings(dynamic_batch)) | ||
| batch = {**dynamic_batch, **static_batch} | ||
| return PreparedFwdBwdInput(batch=batch) |
There was a problem hiding this comment.
To ensure robustness in multi-controller/multi-host JAX environments and handle edge cases like zero-width or non-fully-addressable arrays (similar to the defensive logic implemented in _place_read_only_inputs), we should defensively map the placement over the dynamic batch leaves. This avoids potential runtime errors or redundant device_put operations on already-sharded or non-addressable arrays.
batch = self._prepare_batch(payload)
if self._mesh is not None and isinstance(batch, dict):
dynamic_batch, static_batch = _split_static_and_dynamic(batch)
shardings = self._batch_data_shardings(dynamic_batch)
def place(leaf: Any, sharding: Any) -> Any:
if not isinstance(leaf, (jax.Array, np.ndarray)):
return leaf
if leaf.ndim == 0 or 0 in leaf.shape:
return leaf
if isinstance(leaf, jax.Array) and not leaf.is_fully_addressable:
return leaf
if sharding is None or getattr(leaf, 'sharding', None) == sharding:
return leaf
return jax.device_put(leaf, sharding)
dynamic_batch = jax.tree.map(place, dynamic_batch, shardings)
batch = {**dynamic_batch, **static_batch}
return PreparedFwdBwdInput(batch=batch)…AGE_NEXT_MICROBATCH) Between optimizer steps the trainer chips sat idle while the orchestrator closed step k (metrics, logging, checkpoint, weight sync) and then packed and shipped step k+1's first micro-batch: ~7.7 s per step on the 397B v7x trainer. - --pipeline_step_boundary (PIPELINE_STEP_BOUNDARY): pack the next step's first micro-batch while this step's last ones train, send it as soon as the checkpoint, weight sync (its staging, with --async_weight_sync) and commit are done, and run step k's metrics and logging once that call has returned, while the device computes it. The trainer gets the same calls in the same order. - --stage_next_microbatch (STAGE_NEXT_MICROBATCH, requires the above): send the next step's first micro-batch to the trainer host while the device still computes this step's last one (TrainerWorker.stage_fwd_bwd), so its train call carries no data. - --stage_before_update (default on): stage between the step's last micro-batch and its update, which the trainer defers (fwd_bwd(apply_optimizer=True, defer_update=True), then apply_deferred_update()). Under MaxText's retry_when_tokens_dropped, update() waits for the last micro-batch's overflow flag, so a stage after the update's dispatch would find the device idle and latch staging off. Staging needs MaxText's prepare_fwd_bwd_input and wait_for_device (AI-Hypercomputer/maxtext#5718). Without them, or when a stage measures no overlap, staging latches off and the step runs as before. While an update is deferred the worker refuses whatever needs the step closed (another micro-batch, update, get_metrics, a checkpoint, weight sync, actor log-probs, eval). A stage that fails on its way to the trainer falls back: the micro-batch is sent whole in the next step's first call and staging stays off for the rest of the run. compile and restore_checkpoint drop a stale stage. Measured on the RL trainer bench (AI-Hypercomputer/maxtext#5701) through the equivalent bench-driver switch, same pod, fix off vs on: 105.97 -> 98.31 s per optimizer step on 128 v7x chips. This orchestrator code also ran on TPU in a reduced rig: the real StandardRLProgram and TrainerWorker over gRPC, driving the team cell's MaxText trainer on Pathways with synthetic groups instead of rollouts, 3 micro-batches per step. The exact-order checks were clean, and the step-boundary idle time fell from 12.46 s to 0.02-0.03 s. Tests: rl_program_pipelined_boundary_test.py, rl_program_pipelined_boundary_robustness_test.py, rl_program_stage_before_update_test.py, distributed_rl_engine_staging_test.py, distributed_rl_engine_deferred_update_test.py, trainer_worker_staging_test.py, trainer_worker_deferred_update_test.py, trajectory_queue_manager_test.py, and the submission-launcher test (both flags off by default; STAGE_NEXT_MICROBATCH without PIPELINE_STEP_BOUNDARY is refused).
…e staged step boundary mlperf_gbs1024_submission.sh now sets, unless already set: - TUNIX_LOGPS_SHARDED_HEAD=local: the trainer computes log-probs with the sharded head (k8s_launcher.sh passes it to the trainer); 0 turns it off. - PIPELINE_STEP_BOUNDARY=1 and STAGE_NEXT_MICROBATCH=1: the orchestrator sends the next step's first micro-batch ahead of this step's bookkeeping and stages it on the trainer before this step's deferred update. Staging needs AI-Hypercomputer/maxtext#5718. STAGE_NEXT_MICROBATCH follows PIPELINE_STEP_BOUNDARY unless set, so PIPELINE_STEP_BOUNDARY=false alone turns both off. The other 397B v7x recipes keep all three off. The recipe comments name the code instead of pull request numbers. With these and the trainer settings of the first two commits, the trainer of the RL trainer bench (AI-Hypercomputer/maxtext#5701, 128 v7x chips) measured 98.31 s per optimizer step, against 130.80 s for the team recipe. Tests: the head reaches the trainer JobSet only and 0 turns it off; the staged boundary is on by default here and PIPELINE_STEP_BOUNDARY=false turns both flags off; the refusal test now sets PIPELINE_STEP_BOUNDARY=false; the other 397B v7x recipes keep the head and the boundary off.
…ext#5718 The staged step boundary needs MaxText's prepare_fwd_bwd_input and wait_for_device (AI-Hypercomputer/maxtext#5718). With the previous pin (60425bfbd) an image built from this tree has neither, so the orchestrator latches staging off and the step boundary runs pipelined but unstaged. Pin both MaxText lines to c79bcfd16: 60425bfbd plus #5718's six commits, which change only the MaxText training engine, its README and its tests. Re-pin to #5718's merge commit when it lands.
…gnostic prints - Name the sources of the timing numbers in words: the MLPerf GBS-1,024 submission baseline run, and another 397B trainer run, both from their TRAINER_STEP_TIMING lines. Each is explained once, where its costs are defined. Helpers, constants, labels and test names that carried internal run names are renamed to match. - Rename the sibling-module alias to boundary_test, and refer to "the pipelined-boundary suite" in assertion messages. - Say what each test checks instead of pointing at notes readers cannot see, and drop process tags (pre-registered, measured, inferred) from comments and docstrings. - Link AI-Hypercomputer/maxtext#5718 in place of a commit GitHub cannot resolve, and write maxtext#5673 in full. - Remove 19 diagnostic print() calls; nothing reads their output. - Rename test_lane_busy_costs_with_three_or_more_calls_keep_staging to test_busy_costs_with_three_or_more_calls_keep_staging. - rl_program.py, comments and docstrings only: say what the "stage did not hide" latch catches, and that on Pathways an idle device's wait already reads above its threshold; reflow a ragged comment and two lines over 80 columns. The module's AST is unchanged. - trainer_worker.py docstrings and the stage-before-update test's module docstring say the same about device_wait_s: a short wait means the device was idle, and on Pathways an idle wait reads about 2.7-2.9 s, not ~0. Docstrings only.
e526ea02 is #5718 squash-merged onto atwigg/mlperf. Its tree is identical to the previous pin c79bcfd16 (tree e1330ef9), so the installed code does not change; the pin now names the merged commit instead of the PR branch.
… join TRAINER_TUNED_FLAGS TRAINER_TUNED_FLAGS=true now selects the trainer cell that measured fastest on the RL trainer bench with every numerics gate passing. In mlperf_397b_256_v7x.sh the block of MaxText keys (part 2) gains sixteen keys and loses one: - wi_tile_dlhs_mlp_dim=2048, decoder_layer_input=device, moe_routing_maps=device and moe_gmm_v2_dlhs_transpose_rhs=true: the K tile of wi's input-gradient GMM matched to the fused gate and up weights, each decoder layer's input kept in HBM, the backward reusing the forward's MoE routing maps, and the input-gradient GMMs reading the expert weights through the kernel's transpose_rhs. - moe_routing_index_math=count, moe_unsplit_prefused_wi=true, moe_roe_packed_routing_gather=true, moe_router_grad_from_wo_input=true and moe_wi_gmm_fused_swiglu=true: the MoE keys of AI-Hypercomputer/maxtext#5758 in the set its commits recommend (the routing index maps by counting instead of argsorts, the fused wi kept whole up to the GEMM, one packed routing all-gather per layer, the routing-weight gradient from the down projection's input, the SwiGLU in the wi GEMM's epilogue). - gdn_conv1d_input_in_qkv_dtype=true and gdn_conv1d_dense_token_stack=true: the GDN Conv1D backward without its f32 copy of the input, and the forward kernel's Conv1D output built from dense per-token views. - wi_tile_fwd_batch_seq, wi_tile_drhs_batch_seq, wo_tile_fwd_batch_seq, wo_tile_dlhs_batch_seq and wo_tile_drhs_batch_seq=512 (from 256): row tiles of 512 for the five MoE GMMs with 1024 tiles. wi_tile_dlhs_batch_seq stays 256 next to its 2048 K tile. - sa_use_base2_exp=false leaves the block, so the trainer keeps the recipe's sa_use_base2_exp=true: base-2 exponentials measured 0.025 s faster per micro-batch (o275x-gx-b2e), and their numerics run passes (nxgx-b2e). In mlperf_gbs1024_submission.sh, which turns TRAINER_TUNED_FLAGS on, the switch also sets three defaults, each still overridable: - COMPUTE_LOGPS_CHUNK_SIZE=4096 (was 1024), and 1024 when TUNIX_LOGPS_SHARDED_HEAD turns the sharded head off (0, false, off or no, in any case): chunks of 4096 were measured with the head only, and without it the log-prob program failed to compile on TPU at chunks of 2048 and 4096. - TUNIX_LOGPS_SHARDED_HEAD=fused (was local). TRAINER_TUNED_FLAGS=false keeps local. - PACK_PAD_POSITION=start: each packed row puts its padding first. Bench (AI-Hypercomputer/maxtext#5701 driver, Qwen3.5-397B-A17B, 128 TPU v7x, mesh fsdp 32 x context 4 x expert 2, ragged_buffer_factor 2.75, 9 micro-batches per step, mean of steps 1 to 4, no token dropped in any arm): - before these keys, the fused head and pad first: 106.30 s per step (o275x-hd), 2,774 tokens/s/chip; - this cell: 82.671 and 82.670 s per step in two pods (o275x-c9pfxhv-l3), 3,567 tokens/s/chip. Part of the difference is the staged step boundary (AI-Hypercomputer/maxtext#5718), which this recipe already turns on and o275x-hd did not run; per regular micro-batch (the median of micro-batches 3 to 9), which the boundary does not change, about 10.77 -> 9.01 s; - within it, each against a control in the same pod: padding first 83.285 -> 82.043 s per step (-1.242 s) and the fused head 83.253 -> 82.708 s per step (-0.545 s). The recipe lists the measurement of each key. The measured build switched the MaxText changes with environment variables, and it also had one attention change on that neither MaxText PR carries (attention's load-balancing permute hoisted; alone 99.504 -> 99.267 s per step). The keys themselves, on the MaxText change's own code (2d82c5199) without gdn_conv1d_dense_token_stack, the fused head and pad first, measured 85.295 and 85.251 s per step (o275x-prc-l3-prof, o275x-prc-l3-r), within 0.23 s of what the environment-switch arms predicted. Both arms set moe_route=device, which nothing in #5698 or #5758 tags, so moe_route stays out of the block as before; and the stop of the frozen router's gradient, an environment switch in the measured build, is unconditional in #5758, with no key. Numerics (loss gates off, random advantages, three steps, against two reference runs whose gradient norms differ by 1.9e-5 relative): all six gates pass for this cell (nxc9pfxhv) and for the keys on the MaxText change's own code (nxprc). moe_wo_gmm_fused_dswiglu stays out: next to the fused SiLU GMM at these tiles it computed a wrong SwiGLU gradient, and its fix is still being verified on TPU. moe_routing_maps and moe_gmm_v2_dlhs_transpose_rhs exist in AI-Hypercomputer/maxtext#5698, and the seven other new MaxText keys in the MaxText change of AI-Hypercomputer/maxtext#5758 (on top of #5698); none of the nine is defined at this tree's pin (e526ea02), and MaxText refuses a key it does not define when the trainer starts. TRAINER_TUNED_FLAGS=true, the default of mlperf_gbs1024_submission.sh, therefore needs requirements/maxtext_requirements.txt moved to a MaxText with #5758 before an image built from this tree runs it. TRAINER_TUNED_FLAGS=false and the other 397B v7x recipes pass none of the nine. Tests (mlperf_submission_launcher_test.py): the exact-set test expects 23 keys, _MAXTEXT_DEFAULTS has the nine new keys' defaults, and the switch-off and other-recipe tests check that none of the nine is passed; the default launch has base-2 exponentials on, chunks of 4096, the fused head and --pack_pad_position=start; TRAINER_TUNED_FLAGS=false keeps the local head and the padding at the end; an explicit pad position, head or chunk size wins; TUNIX_LOGPS_SHARDED_HEAD=0, false, OFF or No gives chunks of 1024; the other 397B v7x recipes pass no pad position. Fifteen deliberately wrong copies of the recipes each fail at least one of these tests: a key dropped from the block or given the wrong value, one tile left out, wi_tile_dlhs_batch_seq raised to 512, sa_use_base2_exp=false kept, moe_wo_gmm_fused_dswiglu or moe_route=device added, a new key passed outside the block at its default, the local head under the switch, the chunk fallback removed or made case-sensitive, chunks of 1024 by default, pad first outside the switch, pad_align passed, and both recipes as they were before this change.
At d66a820, the MoE, GDN and model files of this branch come from #5758, which is built on all of #5698. But this branch ports only part of #5698 onto atwigg/mlperf. Several files it needs are still the atwigg/mlperf versions, so the trainer does not start: - utils/sharding.py lacks CompoundLogicalAxis, so `import maxtext.layers.moe` raises ImportError; - utils/max_utils.py lacks is_eval (moe.py); - models/weaver.py is missing, and common_types.py lacks DecoderBlockType.WEAVER, which nnx_decoders.py maps for every decoder; - kernels/megablox/pallas_mosaic_tpu_v2_tgmm_kernel.py builds GmmConfigs without transpose_rhs, which the gmm_v2 kernel requires: TypeError when the MoE backward compiles. This brings src/maxtext to the tree it was developed and measured on: #5758 (20c877a, on #5698 + #5718) plus this branch's own commits (the GDN segment-path skip, the GDN forward zero-fill skip, gather_scanned_kernels_once and the cp-as-ep-dense rules), which apply to it cleanly.
TL;DR
prepare_fwd_bwd_input(payload)turns the next micro-batch intofwd_bwd's inputs and starts copying them to the chips;fwd_bwd(prepared)then skips that work.wait_for_device()waits until the engine's queued work has finished and returns the seconds it waited.update(). On TPU only order and data were checked (a bench pair and the reduced rig), not loss or gradients. Not run yet: a TPU A/B with real gradients, and any end-to-end RL run.mlperf_gbs1024_submission.sh startwith the image pinned by digest (staging is on by default;STAGE_NEXT_MICROBATCH=falseturns it off). After this PR merges, re-pin 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878 to the merge commit. Draft until the two TPU runs under "Not yet tested" have run and passed.Terms used below
update(), so the copy overlaps compute.update()at 128 chips (5 at 256 chips, 3 on the reduced rig); 1,152 trajectories (72 prompts x 16 generations) on the bench. Bench numbers are means of steps 1-4; step 0 is warm-up.ragged_buffer_factor2.75 with dropless replay, context on device, log-prob chunk 1024. Not the recipe's own earlier default (fully dropless, chunk 512), which was not measured here.gdn_cp_matmul_precision=high,cast_grads_after_all_reduce=true) and the sharded log-prob head (measured as google/tunix#2873's file). One of the 7,moe_route=device, had no effect on the bench and is left out of the shipped recipe.retry_when_tokens_dropped: a micro-batch that overflows its MoE expert buffers is rerun. Eachfwd_bwdreads the previous micro-batch's overflow flag on the host, andupdate()reads the last one.--pipeline_step_boundary: it sends the next step's first micro-batch without waiting for this step's metrics and logging. On its own it copies nothing to the chips early.What changes
prepare_fwd_bwd_input(payload)(new)src/maxtext/training_engine/maxtext_engine.pyfwd_bwdruns, thenjax.device_puts the traced inputs onto the shardingsfwd_bwdis compiled with. Returns aPreparedFwdBwdInputwithout waiting for the copy.wait_for_device()(new)PreparedFwdBwdInput(new)batch.fwd_bwd(signature and one line)PreparedFwdBwdInput, and skips the mapping for the latter. Nothing after the mapping changes.InflightThrottler.block_until_ready()(new)src/maxtext/training_engine/inflight_throttler.pysrc/maxtext/training_engine/README.mdtests/post_training/unit/Size: code +74/-3, README +10, tests +1,090. The methods are not on
AbstractTrainingEngine, so callers look them up (Tunix usesgetattr).prepare_fwd_bwd_inputis safe between a step's lastfwd_bwdand itsupdate(), where Tunix stages; contracts and a call-order example are under API contracts below.How it works
The trainer chips at one step boundary
gantt title One step boundary on the bench, 128 chips, seconds dateFormat X axisFormat %s todayMarker off section Staging off Last micro-batch of step k :b1, 0, 11 Update :b2, 11, 13 Idle - micro-batch 1 copied in :crit, b3, 13, 21 Micro-batch 1 of step k+1 :b4, 21, 32 section Staging on Last micro-batch of step k :a1, 0, 11 prepare_fwd_bwd_input copies micro-batch 1 :a2, 0, 8 Update :a3, 11, 13 Micro-batch 1 of step k+1 :a4, 13, 24Schematic in whole seconds, bench at 128 chips. Measured: the 7.6-7.7 s idle fell to about 0 s, and regular micro-batches stayed at 10.75 s. The copy bar is drawn as long as the gap it removes; the bench does not time the copy itself.
The calls at one staged step boundary
flowchart TB subgraph S1["Call 1 - last micro-batch of step k"] direction LR A1["Tunix fwd_bwd<br/>with the update deferred"] --> C1["MaxText fwd_bwd<br/>dispatches the micro-batch"] end subgraph S2["Call 2 - stage micro-batch 1 of step k+1"] direction LR A2["Tunix stage_fwd_bwd"] --> C2["NEW prepare_fwd_bwd_input<br/>starts the copy and returns"] end subgraph S3["Call 3 - deferred update"] direction LR A3["Tunix apply_deferred_update"] --> C3["NEW wait_for_device<br/>waits for the last micro-batch"] --> D3["MaxText update<br/>reruns it if it overflowed"] end subgraph S4["Call 4 - step k+1 begins"] direction LR A4["Tunix metrics, checkpoint<br/>and weight sync, then fwd_bwd_staged"] --> C4["MaxText fwd_bwd<br/>on the prepared input, no copy"] end S1 --> S2 --> S3 --> S4Tunix's call order once staging is on, one row per call; each Tunix call reaches MaxText through its TrainerWorker. NEW marks this PR's methods.
Why the stage goes before
update(). With dropless replay (the layout measured here),update()first reads the last micro-batch's overflow flag on the host, so it returns only after that micro-batch has finished. A stage placed afterupdate()can hide only behind the update itself; on the reduced rig (update about 1.6 s) that left 7.43 s idle per boundary. Placed before it, the copy hides behind the whole last micro-batch (13.4 s on the rig, 10.75 s on the bench), and the rig's idle fell to 0.02 s.Results
RL trainer bench (#5701): staging off vs on, same pod
Mesh fsdp 32 x context 4 x expert 2 at 128 chips (fsdp 64 at 256),
ragged_buffer_factor2.75 with dropless replay, log-prob chunk 1024. Staging ran through an earlier form of #5719's bench-driver switch, which callsprepare_fwd_bwd_inputandwait_for_devicedirectly; the two engine methods and thefwd_bwdchange were this PR's, and the rest of the engine was the bench image's.Seconds per optimizer step:
Changes are computed from unrounded means (71.474 and 56.115 s at 256 chips), so they can differ from the rounded columns by 0.01 s.
tok/s/chip, same runs:
fwd_bwdcall took 7.6-7.7 s longer than a regular call at 128 chips (15.1-15.5 s at 256). That call waits for micro-batch 1, including its host-to-device copy. With staging the extra time was gone (-0.04 to -0.05 s at 128, -0.03 to -0.04 s at 256). Regular micro-batches did not change (medians, off vs on: 10.7545 vs 10.754 s at 128 chips, 11.0765 vs 11.075 s at 256).prepare_fwd_bwd_inputreturned in 0.007-0.041 s andwait_for_devicethen waited 10.49-10.72 s: the last micro-batch was still computing while the copy ran. The repeat arm's wait reached 10.725 s. Without the precision keys: 0.009-0.042 s and 10.69-10.91 s. At 256 chips: 0.013-0.095 s and 10.80-11.05 s.Reduced rig with Tunix's real orchestrator
Tunix's
StandardRLProgram, through a subclass that only adds observation hooks, andTrainerWorker, at an earlier revision of the boundary code (26b06c504), with recorded trajectory groups from the first optimizer step of a 397B RL run. No weight sync or checkpoint save. The trainer ran the team cell (ragged_buffer_factor2.75 with dropless replay; about 13.4 s per micro-batch). Each scored arm ran 5 steps, all arms in one pod.--no-stage_before_update)wait_for_deviceinside the deferred update still waited about 6.7-9.0 s: the last micro-batch was still computing after the stage returned. Staged after the update it read 2.83-2.94 s, the floor this wait shows on Pathways with nothing left to compute.Context. google/tunix#2878 tabulates the 128-chip stack under this result, with each stage's numerics check: from the team cell (130.80 s per optimizer step, 2,255 tok/s/chip) to the final stack without staging (106.16 s, mean of 3 pods). 3,000 tok/s/chip is 94% of the pretraining target (3,193 tok/s/chip, from the MaxText pretraining run of this model on 128 v7x chips, rows without padding); on non-padding tokens, about 2,712 (85%).
Correctness
update()after the next one was prepared, and through Tunix's realTrainerWorker; the negative controls are caughtapply_optimizer, is pinned by the bench switch's CPU tests (#5719), not on TPUThe TPU runs cannot settle numerics. The bench's speed runs log loss and gradient norm as 0 by design: its log-prob gates drop every sequence of a random-weight model. The real-gradient A/B turns the gates off and gives every sequence segment a seeded random advantage. The rig runs at learning rate 0 and its loss moved by more than 2% from step to step even with staging off, so its automatic verdict reads FAIL although 74 of 79 checks passed (2 gradient-norm bands missed, 3 loss bands unscorable); every order, data and decision check passed. Until the real-gradient A/B runs, numerics rest on the CPU tests; google/tunix#2878 lists the caller's own tests.
How to run end to end
Run it now (before merge, for example for gate 2): skip steps 1-2. google/tunix#2878 already pins this PR's head (
c79bcfd16) in both MaxText lines ofrequirements/maxtext_requirements.txt. After merge: steps 1-3.atwigg/mlperf.maxtextandmaxtext-vllm-adapter) fromc79bcfd16to this PR's squash-merge commit. MaxText squash-merges and deletes the branch, so pin the commit that lands onatwigg/mlperf.bash build_docker.sh --maxtext --raiden --deepswe; its Dockerfile installs MaxText and the vLLM adapter from the pin). Createraiden_wheels/first: the Dockerfile copies it, and it is not in git (see 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878's build notes). Push the image, and launch the submission recipe with the image pinned by digest:TUNIX_IMAGE(gcr.io/cloud-tpu-multipod-dev/atwigg/trellis:latest, built on 2026-09-27) has neither this PR nor 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878's flags. A launch with it fails at argument or config parsing instead of running the wrong code.enable_checkpointing=falseand the inertmoe_route=device.The switches, with their defaults in
mlperf_gbs1024_submission.sh:PIPELINE_STEP_BOUNDARY1PIPELINE_STEP_BOUNDARY=false, which turns staging off tooSTAGE_NEXT_MICROBATCHPIPELINE_STEP_BOUNDARYSTAGE_NEXT_MICROBATCH=falsekeeps the pipelined boundary without stagingRAGGED_BUFFER_FACTOR2.75withRETRY_WHEN_TOKENS_DROPPED=true-1gives the fully dropless layout, not measured with staging. Write-1, not-1.0: the recipe compares the value as a string.The stage-before-update order is the orchestrator's default (
--stage_before_updateinrun_deepswe_dist.py). The recipe does not pass this flag, so trying the other order (--no-stage_before_update; 7.43 s idle per boundary on the reduced rig) needs an edit to the launcher's orchestrator command.Every other recipe keeps staging off, and
STAGE_NEXT_MICROBATCHwithoutPIPELINE_STEP_BOUNDARYis refused. google/tunix#2878 describes its other switches (TRAINER_TUNED_FLAGS,TUNIX_LOGPS_SHARDED_HEAD).What the orchestrator log should show:
[pipeline] STAGED next step's first microbatch on the trainer before the update in <s>; the last microbatch then still computed for <s>, after optimizer steps 1-6 of the recipe's 10-step run (6 of 9 boundaries), or fewer if rollouts are late. Tunix refuses to stage after the steps that write scored checkpoints (fromval_start_step7). A "still computed" value of about 2.7-2.9 s means the chips were already idle (the Pathways floor, measured on the 2026-10-09 build, not on the recipe's 2026-10-13 build), not that time was left.[pipeline] stage_next_microbatch is set but the trainer cannot stage a microbatch ...: it means the image's MaxText lacks this PR and staging is off.Not yet tested / gates before un-drafting
This PR stays a draft until these have run on TPU and passed. Each needs TPU time that has not been available.
Also pending on the Tunix side, outside this PR's code: google/tunix#2878's same-pod real-gradient A/B of its first commit's settings (the libtpu flag and the six MaxText keys it ships;
moe_routeis now left out) against the team cell, with pass bounds written down before the run.ragged_buffer_factor-1, no replay)update()does not wait for the last micro-batch, so the deferred update'swait_for_devicemay hold it back by one readiness round trip (about 2.7-2.9 s on Pathways, inferred). The submission recipe defaults to the replay layout.maxtext_engine.pyitself and 2 othertraining_engine/files.Commits in this PR
Six linear commits on
60425bfbd(atwigg/mlperf):[MaxTextEngine] Prepare a micro-batch's inputs on the mesh ahead of its fwd_bwdPreparedFwdBwdInput,prepare_fwd_bwd_inputand thefwd_bwdchange, with the one-device tests.[MaxTextEngine] Report how long a caller waited for the devicewait_for_deviceandInflightThrottler.block_until_ready, their tests and the 4-device test.[MaxTextEngine] Test prepared micro-batches under dropless replay and the fused updateTrainerWorkerstaging test, makes the 4-device test re-exec itself; README.[MaxTextEngine] Test the order Tunix uses to stage before a replayed updatemaxtext_engine_stage_before_update_test.py: the staged-before-update order with reruns insideupdate(), and two negative controls.[MaxTextEngine] Drive the staged-before-update order through Tunix's TrainerWorkerTrainerWorker(defer_update,apply_deferred_update).[MaxTextEngine] Document that staging leaves a pending overflow check aloneSuperseded PRs (do not merge)
Each is titled [DO NOT MERGE] and opens with a pointer note.
STAGE_BEFORE_UPDATEswitch in #5701's bench driver, which calls these two methods. An earlier form of it ran the three staging A/Bs in Results; the earlier stages of the stack (130.80 s to 106.16 s) ran on #5701's driver without it.atwigg/mlperfplus this PRa29991d3f)Known limitations and follow-ups
None of these blocks review.
wait_for_deviceon Pathwaysfwd_bwd(compile(None))fwd_bwdonce, at the first staged boundary. AOT kernels (compile(dummy_data)) are unaffected, and the submission recipe compiles them at warm-up. The tests count_compile_for_batchcalls and cannot see the re-trace. Follow-up: compile the lowered program on that path, as the background compiles already do, or warn.prepare_fwd_bwd_inputdoes not block, and none pins the secondswait_for_devicereturns: four blocking variants of the prepare, and three wrong-value variants of the wait, pass the tests they were run against. The only multi-device test uses an all-fsdp mesh (production is fsdp 32 x context 4 x expert 2). Follow-up: add those tests.fwd_bwd(under replay, until its overflow check resolves): about 20 MB per device on the bench (5.15 GB over 256 devices; derived). The in-flight throttler does not count it and no compile can see it. It fit on the bench.jax.device_puton every traced leaf, without the zero-width and non-fully-addressable skips of_place_read_only_inputs.STAGE_NEXT_MICROBATCHis set and the trainer cannot stage.API contracts, dropless replay and the fused update
prepare_fwd_bwd_input(payload) -> PreparedFwdBwdInput_prepare_batch(payload), the mappingfwd_bwdruns. With a mesh and a dict batch itjax.device_puts the traced half onto_batch_data_shardings(the per-leafin_shardingsoffwd_bwd's compile) and leaves the closed-over half (Tunix'salgo_config,pad_id,eos_id) as is. Otherwise it returns the mapped batch, andfwd_bwdplaces it as before.fwd_bwdand itsupdate(): under dropless replay it neither reads nor resolves the last micro-batch's pending overflow flag, and the mapping reads only the payload, so the prepared input is the same before or after the update.wait_for_device() -> floatfwd_bwdloss and the lastupdate's gradient norm; at mostmax_inflight_computations, 2 by default) and returns the seconds it waited.wait_for_next/wait_for_alland their metrics are unchanged. A pending overflow check, and its rerun, stay with the nextfwd_bwdorupdate().jax.block_until_readyraises, so a failed micro-batch or update surfaces here; Tunix reports it as the update's failure.is_ready(): under Pathways a readiness check is a proxy RPC and can read "not ready" for finished work.last_update_tokenis the gradient norm waited on here.retry_when_tokens_dropped, #5570)update()reads that micro-batch's kept inputs, even when the next one was already prepared.maxtext_engine_prepared_input_dropless_test.py;maxtext_engine_stage_before_update_test.py(replays per step[1, 2, 1])retry_dropless_final_micro_batch, #5673)apply_optimizerinfwd_bwd, exactly as for a payload.[0, 1, 0]in the order test)update()per step. With--no-stage_before_updatethe stage lands after the fused update; that is bit-identical too.maxtext_engine_prepared_input_test.py;maxtext_engine_stage_before_update_test.py(both also throughTrainerWorker)The supported call order, for any caller:
Repro: CPU tests and the bench A/B
CPU tests, from the repo root with Tunix importable:
The two
TrainerWorkertests skip unless the installed Tunix hasstage_fwd_bwdandapply_deferred_update(google/tunix#2878).Bench A/B at 128 and 256 chips
This route was checked against the measured arm offline, not re-run on TPU: the measured arms were rendered from the team recipe's launch render, with their flags set explicitly.
Bench: Add an RL trainer bench under experimental/rl_trainer_bench #5701 (
src/maxtext/experimental/rl_trainer_bench) at its commit "rl_trainer_bench: rename e2e identifiers to bench/launch names", with [DO NOT MERGE] perf(rl_trainer_bench): opt-in STAGE_BEFORE_UPDATE step boundary (-7.2% step time on the 397B bench; stacked on #5701) #5719's commit "rl_trainer_bench: STAGE_BEFORE_UPDATE switch for the step boundary" on top. The TPU runs used an earlier form of the same switch.Image: a client image whose MaxText has this PR's training-engine changes and whose Tunix does not fuse the update into the last train call (with perf(orchestrator): fuse final microbatch fwd_bwd+update RPC and defer get_metrics off weight-sync path google/tunix#2719's fused update the switch raises at the first boundary). The measured runs used, on the Pathways build of 2026-10-09:
gcr.io/tpu-vm-gke-testing/sivaibhav/trellis-experimental@sha256:b4647c188b25638e644c7a6fc53fb8e0751741525a277a1bcf07dca92ba9a6acus-docker.pkg.dev/cloud-tpu-v2-images-dev/pathways/gke/datenglin/unsanitized_server@sha256:662f33b0ca6cd236a12534e4c78a8b7d715868669366738a015401925a39cb86us-docker.pkg.dev/cloud-tpu-v2-images-dev/pathways/gke/datenglin/unsanitized_proxy_server@sha256:e714351cdde50d306dd62f127cffabcdaf98dff80ad22e8b4c66c5a1d1df6ddbAt pod start they overlaid files carrying this PR's training-engine changes on the image's MaxText, and installed the head file measured on TPU (google/tunix
a29991d3f).Launch render (Add an RL trainer bench under experimental/rl_trainer_bench #5701's README, step 1):
DRY_RUN=true bash tunix/experimental/examples/recipes/mlperf_gbs1024_submission.sh startfrom 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878, withBATCH_SIZE=72(1,152 trajectories per optimizer step, as measured).Bench render (Add an RL trainer bench under experimental/rl_trainer_bench #5701's README, step 2):
render_bench.shwithSTEPS=5 MICRO_STEPS=9andEXTRA_MAXTEXT_FLAGS="enable_checkpointing=false", as the measured arms set (they also setmoe_route=device, which did nothing on that MaxText). Render twice, identical exceptDRIVER_ENV="STAGE_BEFORE_UPDATE=on"vsoff.Run both arms back to back in one pod.
Read the client log:
completed step: <n>, seconds: <s>, ... fwd_bwd dispatch s: [...]for steps 1-4, andSTAGE_BEFORE_UPDATE step=<k> staged=... stage_s=... device_wait_s=.... A healthy staged arm shows each step's 2nd dispatch entry near the regular call time, anddevice_wait_sfar abovestage_s.256 chips: render with
TRAINER_SLICE=4x8x8(fsdp 64 x context 4 x expert 2) andMICRO_STEPS=5(micro-batches of 128 rows), and addFINAL_PAD_FRACTION=0.625toDRIVER_ENV, as the measured arms did, so a step keeps the 560 real rows of the 128-chip runs.Evidence: the tests in this PR and the recorded CPU runs
maxtext_engine_prepared_input_test.py(one device, toy model, Tunix-shaped input; 9 tests)_compile_for_batchcall (the same count as the payload run); also withwait_for_deviceafter each prepare, queue and metric writes unchanged. A spy onjax.block_until_readysees every queued entry in order. Traced leaves carry exactly_batch_data_shardings; step state is untouched. Through Tunix's realTrainerWorker(skipped without staging).InflightThrottler.block_until_readyneither dequeues nor writes (3 tests). Positive control: a one-token change is seen.maxtext_engine_prepared_input_dropless_test.py(one-layer Qwen3.5 MoE rig; replay and the final program on; 3 tests)maxtext_engine_prepared_input_sharded_test.py(1 test, plus 3 in its child)maxtext_engine_stage_before_update_test.py(dropless rig on thecompile(None)path; 3 steps of 3 micro-batches, each step's last one overflows; 5 tests)fwd_bwd(last), prepare(next),wait_for_device(),update(),fwd_bwd(prepared)equals the base order: replays per step[1, 2, 1], the flag still pending after the wait, one_compile_for_batchcall (the JAX re-trace on this path is not counted). The same with the final dropless program on ([0, 1, 0]), and throughTrainerWorkerwith the deferred update (defer_updatenever reaches the engine). Negative controls: a rerun of the prepared micro-batch instead of its own, and preparing step k+1's 2nd micro-batch instead of its 1st, are both caught.Recorded runs (CPU, jax 0.11.1):
TrainerWorkercode equals 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878's (only docstrings differ): 18 passed, 10 subtests.atwigg/mlperf): 16 passed, 2 skipped (the twoTrainerWorkertests).maxtext_engine_test.py111 passed on base and branch, with the same 1 failure; the other suites (maxtext_engine_*_test.py,router_replay_engine_test.py) 130 passed and 13 failed of 143 on base, 148 passed and 13 failed of 161 on this branch. 12 failures are identical pre-existing ones; the 13th is a libtpu lockfile collision between parallel test workers, which hit a different test on each side.Evidence: TPU run records (predictions, arm order, rig checks)
Bench (predictions as predicted value [allowed band])
Reduced rig
26b06c504; since then it changed only in comments, docstrings, one flag's help text and the keyword order in one constructor, plus an unrelated base merge.fwd_bwdreturns only when its micro-batch is done), so its step-time-saving prediction missed (4.78 s against 9-20 s). Its staged-after control expected Tunix to switch staging off; it did not, because with nothing left to computewait_for_deviceread 2.69-2.87 s at all 8 staged boundaries.Corrections to commit messages (pushed history is not rewritten)
[MaxTextEngine] Prepare a micro-batch's inputs on the mesh ahead of its fwd_bwdfwd_bwdonce (see Known limitations).[MaxTextEngine] Report how long a caller waited for the devicemaxtext_engine_prepared_input_sharded_test.py, with Tunix's real model input.[MaxTextEngine] Test prepared micro-batches under dropless replay and the fused update--no-stage_before_updateorder.[MaxTextEngine] Test the order Tunix uses to stage before a replayed update_compile_for_batchcalls on thecompile(None)path, where the first prepared input re-tracesfwd_bwdonce (one more backend compile than the unstaged order).[MaxTextEngine] Document that staging leaves a pending overflow check aloneragged_buffer_factor2.75 withretry_when_tokens_dropped. In the MLPerf recipes this layout is opt-in (RETRY_WHEN_TOKENS_DROPPEDdefaults to false); only 397B v7x RL trainer: tuned settings, sharded log-prob head, staged step boundary (130.8 -> 98.3 s per optimizer step, trainer bench) google/tunix#2878's submission recipe uses it by default.