iHore/vLLM
8
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Revert "[Docs] Add train-inference consistency agent skill" (#61551) (#62186)

Browse Source 
Signed-off-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
This commit is contained in:
aoshen02
2026-10-17 13:54:30 +00:00
committed by GitHub
co-authored by Claude Opus 5.5
parent 2d8efee9a7
commit 237d99b4af
18 changed files with 0 additions and 3093 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.agents/skills/train-inference-consistency/SKILL.md
-31
View File
@@ -1,31 +0,0 @@
---
name: train-inference-consistency
description: Build and validate train-inference consistency (true on-policy RL), where vLLM rollout and the training forward give bitwise identical logprobs at high throughput. Use when aligning a training stack with vLLM batch-invariant inference, validating forward, backward and quantization contracts, or diagnosing logprob mismatch, layout and performance regressions between training and rollout.
---
# Train-inference consistency skill
Make the training side and the inference side in RL produce bitwise identical logprobs for the same token (diff = 0), and, while keeping 0 diff, bring performance back close to the high-performance baseline.
How to enable batch-invariant inference in vLLM is in [batch invariance](../../../docs/features/batch_invariance.md); this skill covers making the training side match it.
Start here: **[train-inference consistency workflow (C0–C10)](workflow/README.md)**. It answers "what to do next, and what counts as passing"; acceptance details are in [validation.md](workflow/validation.md), and the report template is [report-template.md](workflow/report-template.md).
Find a guide by problem:
| Guide | Questions it answers |
| --- | --- |
| [Fault localization and mismatch bisect](guides/debugging-mismatch.md) | Find the side first, bisect, classify before changing precision, save and compare protocol |
| [Training forward](guides/training-forward.md) | Module inventory and route (mLite / MCore / FSDP2), reproducing the inference forward by calling its kernels, layout and length, framework integration |
| [Backward validation](guides/backward-validation.md) | When to reuse backward, gradient contract, module-level VJP comparison, proxy SFT three-way control |
| [BI kernels](guides/bi-kernels.md) | Which kernel each operator uses in BI mode and what it costs |
| [Precision and quantization](guides/precision-quant.md) | The conversion chain master -> training forward weights -> exported bytes -> vLLM bytes; QAT/STE; known pitfalls |
| [Rollout engine and sampling](guides/rollout-engine.md) | Sampling RNG, prefix cache semantics, MTP constraints, R3 constraints, rollout replay |
| [Performance and profiling](guides/performance.md) | Workload and measurement basis, fair A/B, breakdown and ablation |
| [Checkpoint and resume](guides/checkpoint-resume.md) | Exact resume, derived state after restore, memory budget for colocated training and rollout |
| [Environment provenance](guides/environment-provenance.md) | Which version is actually running now; immutable delivery image |
| [Recipe and acceptance](guides/recipe-and-acceptance.md) | BI=1 / BI=0 arms in one recipe and one image; evaluation contract; paired acceptance |
Process rules: [evidence and experiments](practices/evidence-and-experiments.md), [code and delivery](practices/code-and-delivery.md).
Note: this skill comes from the DeepSeek-V4 and Nemotron-H train-inference consistency work and covers train-inference consistency only. Day-0 model support, per-model case records and cluster-specific notes are not included.
.agents/skills/train-inference-consistency/agents/openai.yaml
-4
View File
@@ -1,4 +0,0 @@
interface:
display_name: "Train-Inference Consistency"
short_description: "Make RL training and vLLM rollout logprobs bitwise identical."
default_prompt: "Use the train-inference consistency workflow to align training with vLLM inference for this model."
.agents/skills/train-inference-consistency/guides/backward-validation.md
-221
View File
@@ -1,221 +0,0 @@
# Backward: reuse and validation
## Reusing backward: when you actually need to write your own
When to reuse backward and when to write glue.
### Conclusion
Backward is determined by the **math formula**, not by the forward kernel. The backward of `y = f(x, W)` needs only the formula and the saved inputs: `dX = dY·W`, `dW = dYᵀ·x`. A vLLM forward kernel does not by itself require a new backward. "No backward operator at the same precision" is not a reason to rewrite: under the high_precision contract the backward runs in BF16/FP32 anyway.
### How to wire it: forward once, backward uses what the forward saved
The training forward is the inference forward. Each module is an autograd Function: its forward calls the inference kernel to produce the visible output and saves what its own backward needs (inputs, visible output, LSE, fixed routing IDs, quantized values and scales); its backward uses only those, through a reused library backward or a closed-form VJP. If the inference kernel does not output something the backward needs, change the kernel to output it.
Delivery requires 0 diff and a training step within budget on the main topologies fixed at C1; C6 checks both together. Nemotron audited this against DS4: attention saves q/k/v, out and LSE; routed MoE saves x, the visible FC1 output and the visible output; linear and norms save their inputs; the router saves logits and fixed IDs. Its open gap is Mamba SSD, because vLLM's visible SSD does not output the chunk state.
### Three cases
| Case | Characteristics | Examples | What to do |
| --- | --- | --- | --- |
| A only rounding differs | Same math; summation order/chunking differs; difference in the last BF16 bit | RMSNorm, Mamba SSD scan, attention, conv1d, unquantized linear | Reuse the training framework's backward (MCore/TE/mamba_ssm); write no code |
| B forward has rounding | FP8/NVFP4 quantization rounds; the true derivative is 0 almost everywhere | Quantized linear, FP8 Q/K/V | First fix the STE proxy (rounding passes through as 1; dX uses dequantized values by default; DS4 and Nemotron currently approximate with TE high_precision master values, and the deviation must be quantified per [layer 2](#layer-2-gradient-contract-comparison)), then reuse the ordinary BF16 GEMM backward |
| C no ready-made backward for the whole segment | Composite of several operators with no library backward for the whole segment | Routed experts: sort -> grouped GEMM -> relu² -> restore -> top-k weighted sum | Write glue that chains ready-made operators; GEMM still uses TE grouped GEMM; compare gradients against an independent reference |
Decision flow: (1) does the forward have rounding/quantization (top-k ID selection does not count; IDs are fixed and not differentiated)? No -> A. Yes -> (2) fix the STE contract -> (3) is there a packaged training backward for the whole segment (TE Linear/GroupedLinear, MCore MoE)? Yes -> B, reuse it and feed master values; no -> C, glue. If A has no same-math training implementation, also write only glue. Look first in mLite `primitive/modules` and `ops`, then in MCore/TE/mamba_ssm.
### How DS4 did it
DS4's only hand-written kernel is the small FP8 quantization Triton kernel in the forward. The backward has no hand-written kernel, only PyTorch closed-form formulas and glue, all in FP32.
| Operator | Case | Backward source | What we wrote |
| --- | --- | --- | --- |
| FP8 dense linear (qkv, shared expert, o-proj) | B | TE `general_gemm`, BF16 master | Nothing |
| Routed experts (grouped) | B+C | TE `general_grouped_gemm`; reuses gate_up, the output of the forward vLLM FC1, without recomputing | Glue: sort, activation derivative, summation |
| Sparse attention | A | cuDNN sparse attention backward, consuming the forward's saved out and LSE | Nothing |
| SwiGLU (with clamp), RMSNorm, q/kv norm, mHC | A | Hand-written closed-form derivative, FP32, `torch.compile` | PyTorch formula |
| Routing score | A (IDs fixed) | FP32 replay of `sqrt(softplus)` + normalization, handed to autograd | PyTorch formula |
| Others | A | Generic bridge: forward via vLLM, backward recomputes the PyTorch graph then `autograd.grad` | One utility function |
DS4 wrote its own for four reasons, none of them a missing same-precision backward kernel: (1) library backwards need intermediates from the library's own forward (TE RMSNorm needs rstd, MCore MoE the token permutation; the vLLM forward has neither); (2) model-specific math is not in the library (`sqrt(softplus)` scoring, SwiGLU with clamp, mHC Sinkhorn); (3) discrete selection must copy the forward (recomputed top-k may pick a different expert on ties, so store the IDs and differentiate only the continuous score); (4) the generic bridge costs too much memory and time (autograd stores all 20 mHC iterations; use a closed-form derivative fused with compile). A custom backward function has no autograd graph inside, so its activation derivative must be closed-form (`dynamic_clamped_swiglu_vjp`; likewise relu² in Nemotron `routed_vjp`).
### Three rules when reusing
1. Intermediate definitions must match: attention backward uses the forward's out and LSE, so LSE base (ln or log2), softmax scale and memory layout must line up. If they do not, recompute the forward with the training kernel inside the backward.
2. Backward need not be bitwise identical: recomputed intermediates differ from the vLLM forward in the last BF16 bit and gradients differ by the same amount, which is allowed. The requirement is that the gradient contract holds and the error is near the noise floor.
3. Reused code must be verified too, per [layer 1](#layer-1-module-level-vjp-comparison): compare dX and dW against an FP32 reference that does not call the backward under test, with the same dY, in multiples of the noise floor.
### Two test-infrastructure false negatives
- A unit test that instantiated QuantFP8 directly lacked production training's vLLM config context (`CompilationConfig(custom_ops=['none','+quant_fp8'])`, Humming, PP4), and the first long-sequence test failed before computing anything. Fix: the test wrapper uses the production config; production operators and tolerances are unchanged.
- An import check without GPU visibility makes vLLM select a non-CUDA platform, which cannot import `reshape_and_cache_flash` (exported only on CUDA). Dependency checks need real device visibility.
- Both are test errors, not train-inference differences.
## Backward validation: is the implementation right, is the approximation reasonable, is training biased
Forward 0 diff shows the logprobs match, not that gradients are right. Validate the backward in three layers:
| Layer | Question answered | Method | Cannot replace |
| --- | --- | --- | --- |
| 1. Module-level VJP comparison | Does the implementation match the agreed gradient contract (missing scale, wrong transpose, broken gradient chain) | compare dX, dW and a one-step update of every module against an independent reference that does not call the backward under test | layers 2 and 3 |
| 2. Gradient contract comparison | Is the contract reasonable: does an approximation (e.g. dX from the master value) bias training | same batch, compare gradients under the two conventions per class and per layer | layer 1 |
| 3. Proxy SFT three-arm comparison | Does the whole training chain have a systematic bias | three SFT paths (under test, gradient reference, native baseline) with the same data and hyperparameters | layers 1 and 2 |
Run 1 → 2 → 3. If layer 1 fails, fix the implementation before any SFT. If layer 3 is abnormal, go back to layers 1 and 2. Passing all three clears the backward; none proves long-term RL stability.
### Write down the gradient contract first
#### The three tiers of a weight
In quantized training one parameter has three values, and the gradient must end up at the one the optimizer manages:
| Tier | Precision | What it is |
| --- | --- | --- |
| main | FP32 | optimizer state; updates accumulate here |
| master | BF16 | the model parameter, converted from main every step; non-quantized layers compute with it directly |
| deployed value | NVFP4 / FP8 etc. | master quantized per the deployment rule; both the training forward and inference use these bytes |
Keep the optimizer's parameter identity: when replacing the forward, do not create parameters the optimizer does not know about. The deployed value is not a trainable parameter.
#### Contract table
A kernel without autograd does not need a surrogate gradient: non-quantized GEMM, norm and attention reuse the existing backward when the math matches; STE is needed only through discrete operations such as quantization rounding. Before implementing, decide per module class:
| Item | What must be decided |
| --- | --- |
| Visible forward | the inference implementation; input/output dtype, layout; quantize/dequantize positions, scale granularity, rounding, clamp, accumulation precision |
| Trainable variables | input, master weight, bias, norm, router; whether scale is fixed, stop-gradient or trainable |
| Surrogate function | whether dX uses the master or dequantized value (see below); identity or clipped STE, and the value at the threshold |
| Discrete state | fixed expert IDs, sparse indices, ordering; continuous score/combine weights that still need gradients |
| backward | reused implementation or custom VJP; cast, accumulation/reduction precision, communication and normalization |
| Save/recompute | quantized values, scale, norm/attention intermediates, routing, RNG; memory cost and recompute rules |
| Reference and acceptance | independent reference implementation, target shape/execution path, metrics and tolerances agreed in advance |
Identity STE sets the surrogate derivative of Q(x) to 1; clipped STE sets 1 inside the range and 0 outside, with the threshold value stated and never switched silently. Declare weight and activation quantization separately; frozen weights do not mean nothing upstream of the activations needs gradients.
#### Linear: which weight dX uses
For `Y = X · Wᵀ + b` with upstream gradient `G`:
```text
dW = Gᵀ · X_ref # returns to master via STE, then the optimizer updates main
db = sum_over_tokens(G)
dX = G · W_ref
```
Two choices for `W_ref`:
- **Dequantized value**: `W_ref = dequant(quant(W_master))`, the value the forward actually used. This is the true derivative of `Y = X · Wqᵀ` with respect to X, and what fake-quant QAT gets from autograd. **Use this by default.**
- **Master value**: `W_ref = W_master` (BF16). Reuses the plain BF16 GEMM backward without dequantizing; the cheapest hookup when the inference kernel takes only packed bytes and has no backward. Cost: dX is no longer the exact forward derivative and ignores the quantization error's effect on the input gradient.
The two give different gradients and must not be swapped silently. With the master value, the layer-2 comparison is mandatory: quantify and record the bias. Likewise for activation quantization, state whether `X_ref` is the quantized or original activation.
Before reusing a linear or grouped GEMM backward, check flatten/transpose, padding and valid tokens, bias reduction, saved inputs, scaling and the accumulation/backprop dtype. Matching output shape does not prove reusability.
### Layer 1: module-level VJP comparison
#### How to build the reference implementation
The reference may use PyTorch ops and autograd but **must not call the backward under test**, or both paths share the error. For discrete operations implement the agreed surrogate backward explicitly; the zero gradient of bare rounding is not a reference.
This bridge takes the output value from the inference implementation and the gradient from an independent surrogate. Use it to build the reference path:
```python
import torch
class VisibleWithSurrogate(torch.autograd.Function):
@staticmethod
def forward(ctx, visible, surrogate):
assert visible.shape == surrogate.shape
assert visible.device == surrogate.device
assert visible.dtype == surrogate.dtype
return visible
@staticmethod
def backward(ctx, grad_output):
return None, grad_output
with torch.no_grad():
visible = inference_linear(x, deployment_weight, scales)
w_ref = dequantize(deployment_weight, scales) # default contract; use master_weight for the master-value contract
surrogate = torch.nn.functional.linear(x, w_ref, bias)
y = VisibleWithSurrogate.apply(visible.detach(), surrogate)
```
- Build the surrogate graph in normal grad mode, outside the custom Function's forward.
- With the dequantized surrogate, `w_ref`'s gradient must reach master via STE: bridge a tensor valued as the dequantized weight whose gradient passes straight to master, then feed it to a plain linear.
- Do not use `visible + surrogate - surrogate.detach()`: floating-point add/subtract and compiler transforms change the result. Do not modify the bridge output in place downstream.
- This bridge is the independent validation reference, not the production wiring; production wiring is in [How to wire it](#how-to-wire-it-forward-once-backward-uses-what-the-forward-saved).
- `visible_forward(visible, native, *inputs)` is another way to write the same bridge: `visible` saves inputs without building a graph, `native` is the training framework's implementation of the same math and reruns the forward on those inputs in backward; precision and speed depend on which `native` is chosen.
#### Comparison steps
1. **Fix the state.** Fix inputs, weights, quantization config, mask/routing. Reference and candidate get independent copies; clear accumulated gradients. Check that the visible forward matches inference, then the saved state the backward needs; at integration, confirm the real forward produces that state.
2. **Same upstream gradient.** Both sides use the same external dY, dtype and values. Cover fixed random, zero, and sparse (non-zero at selected positions) gradients, then samples from real training. `y.sum().backward()` alone misses transpose, permutation and broadcast errors.
3. **Compare item by item.** Compare dX, dW and every bias/norm/router/scale gradient separately: missing, shape, finiteness, non-zero where expected. Do not ignore unused differentiable inputs by default; list parameters the branch genuinely does not use. Report max absolute error, relative L2 `‖g − g_ref‖₂ / max(‖g_ref‖₂, eps)`, norm ratio and cosine; direction alone is not enough. With a zero norm the cosine is undefined; use absolute error.
4. **Set tolerances in advance.** Derive atol/rtol and the relative L2 threshold from dtype, reduction length and reference precision; record the reasoning. No single threshold for all modules, no relaxing after a failure, and no relaxing a same-semantics kernel optimization because "STE is an approximation anyway". Cross-framework bitwise identity is not required.
5. **Coverage.** Small M (decode), large M (training/prefill), different context lengths, unaligned tails, padding/packing/mask, zeros and extreme values; both sides of the quantization threshold, the saturation region, different scales. For MoE cover empty experts, repeated routing, shared experts, top-k ties; for shared parameters check gradient accumulation across use sites. Test under the target dtype/layout, recompute, autocast, CUDA Graph, gradient accumulation and the actual parallel reduction; local first, then integrated.
6. **One-step update.** From the same parameters and optimizer state, run one update each and compare the change in master. If gradients are close but updates differ, check scaling, clipping, optimizer state and parameter binding.
Cross-check smooth non-quantized modules with small, high-precision [gradcheck](https://docs.pytorch.org/docs/stable/generated/torch.autograd.gradcheck.html) or finite differences: fix the discrete branch, avoid non-smooth boundaries, and perturb the weights the forward actually sees. **Do not validate STE with finite differences of the real quantized forward**: that measures the discrete forward, which the STE surrogate derivative deliberately differs from.
#### What to check per module
| Module | Reuse first | Semantics to check |
| --- | --- | --- |
| Dense/QKV/LM head | training linear/grouped VJP, gradient returns to master | dX value convention, cast, bias, gradient accumulation of shared parameters |
| Norm/RoPE | existing training VJP, FP32 statistics if needed | epsilon, RMS or LayerNorm, position offset, rotation layout |
| Activation + quantization | the actual composition of activation derivative and STE | saturation region, whether scale is detached, fusion order |
| Attention/sparse attention | training backward with the same mask/index semantics | causal boundary, softmax scale, LSE definition, GQA, order of selected KV; do not mix auxiliary quantities from different forwards |
| MoE | expert linear/grouped VJP and a differentiable combine | inverse dispatch permutation, gradient accumulation of repeated tokens, empty experts, shared experts, top-k weight normalization |
| Router/sparse selection | fix the discrete selection, keep the continuous score's derivative | no continuous gradient for IDs/indices; do not detach the whole router output |
| Linear attention/SSM | training backward with the same chunked arithmetic | chunk boundary, initial state, state precision, conv state |
| Selected logprob | log-softmax/gather VJP matching the definition | vocab-shard reduction, temperature, token mask, loss normalization |
Recompute must save or deterministically rebuild the scale, routing/indices, RNG and necessary intermediates, and must not switch the surrogate function.
### Layer 2: gradient contract comparison
Layer 1 answers "does the implementation equal the contract", not "is the contract reasonable". When the contract contains an approximation (most often dX from the master instead of the dequantized value), quantify its distance from the exact derivative:
1. On a fixed checkpoint and batch, compute dX (and gradients related to activation quantization) for every quantized layer under both conventions.
2. Report cosine and norm ratio per module class and per layer; inspect outlier layers.
3. Measure each target precision combination separately: changing weight or activation precision changes the bias.
On Nemotron-H (NVFP4 routed experts, FP8 Mamba projections), master-value and dequantized-value dX have cosine ≈0.99 and norm ratio ≈1.00 per class and layer, with both BF16 and FP8 activations, so this approximation does not cause the training score gap. Remeasure for other precisions, a larger learning rate or longer training.
### Layer 3: Proxy SFT three-arm comparison
After the first two layers pass, run a controlled SFT over the whole training chain:
| Path | Visible forward | backward | Purpose |
| --- | --- | --- | --- |
| Under test | aligned inference numerics | production backward | the scheme to deliver |
| Gradient reference | same as under test at the same parameter snapshot | independent surrogate reference (the bridge above works) | isolate backward implementation errors |
| Native baseline | original training numerical path | verified original training backward | measure the recipe's overall quality change |
- First compare gradients and one-step update of under-test vs gradient-reference on a fixed checkpoint and batch, then run curves; once trajectories diverge, visible values differ too. The under-test vs native-baseline gap includes quantization and forward-path differences, not only backward. If the gradient reference's full curve is unaffordable, compare on representative batches and state the missing range.
- Controlled variables: initial weights and optimizer state, data and order, labels/mask, packing, loss normalization, global batch and gradient accumulation, learning rate and schedule, clipping, precision and loss scaling, seed, recompute policy.
- Record train loss, fixed-validation-set NLL, gradient norm, update RMS, clipping ratio, non-finite values, throughput, against optimizer step or cumulative effective tokens. Fix the comparison window and acceptable deviation before the run; a short run is only a smoke test.
- Localization: loss differs at the initial weights -> check forward, data, mask; initial loss matches but one-step update differs -> backward, reduction, scale, clipping, optimizer; diverges only after several steps -> save the first abnormal checkpoint, check recompute, quantization refresh, learning rate, RNG, then replay at the same weights.
- Mark unexplained or unbudgeted differences inconclusive; do not cover them with forward 0 diff.
On Nemotron-H the three paths (aligned trainer with BF16 activations, Megatron QAT, Megatron BF16) ran 40 steps on gsm8k with the same data in the same order: loss 0.3388 / 0.3383 / 0.3374, gradient norm 0.685 / 0.674 / 0.728, curves overlapping. The training loop itself is fine; the RL score gap comes from RL-specific factors.
### Common false differences
Each of these made gradients look mismatched when the comparison method was the problem:
- **Upstream gradient precision differs.** FP32 dY on one side and BF16-rounded dY on the other creates a difference unrelated to the backward. Use the same external dY.
- **Dump timing.** Expert gradients are multiplied by the expert-DP to DP scaling ratio only inside the optimizer step; dumping earlier shows a false multiplicative difference.
- **Unsharded copy under FSDP2.** After forward-only the root module is not resharded, so the parameter fetched by name is the unsharded compute copy, not the real master. Perturbing it or reading its gradient both mislead.
- **"Frozen" is not frozen.** Zeroing the gradient still lets AdamW weight decay change the parameter; a real freeze needs grad None or removal from the optimizer.
- **Test environment differs from production.** A unit test missing production config context, or an import without GPU visibility, selects the wrong implementation or fails to import (see [Two test-infrastructure false negatives](#two-test-infrastructure-false-negatives)).
### What to write in the report
Per module class: gradient contract (dX value choice, STE type, scale handling), reference implementation, upstream gradient construction, shapes/paths covered, metrics and tolerances with reasoning. Layer 2: per-class per-layer cosine/norm ratio and outlier layers. Layer 3: config diff table of the three paths, curves, conclusion within the window. List missing coverage explicitly, not as "the rest is analogous".
For the basic definition of STE see [Bengio et al.](https://arxiv.org/abs/1308.3432); it does not decide the surrogate gradient rule for a specific model.
.agents/skills/train-inference-consistency/guides/bi-kernels.md
-143
View File
@@ -1,143 +0,0 @@
# BI kernels: which kernel each operator uses and what it costs
Read [the numerical contracts](../workflow/README.md) first, then build candidates from this page. Each row gives the kernel used under BI=1, the cost of invariance, known pitfalls and directions already tested dead. Numbers come from Nemotron-H (GB200 single node, 4 GPUs) and DS4-Flash (GB200 DP8/EP8); only conclusions are given. When the user asks to optimize only the explicit Top-K/index overhead, do not expand into whole-model tuning; switching DSL does not relax the contract.
## By operator
| Operator | BI=1 kernel / approach | Cost and upper bound | Pitfalls |
| --- | --- | --- | --- |
| Attention (FA4) | `FlashAttnFixedSplit`: 32 splits × 640 keys, `MAX_SEQ_LEN` 16384; with FP8 Q/KV the KV dequantization path is fixed | The 16384 limit exists for invariance; exceeding it must raise, not silently change the split. MLA matrix absorption changes rounding; mathematical equivalence does not give train-inference bitwise equivalence | On SM100, do not declare support by bypassing the selector's "backend does not support BI" check. Implement and verify the numerical contract in the backend first, then let the selector declare it; BI0 keeps the default fast path |
| Dense GEMM (FP8 / NVFP4 linear) | fixed split-K; under Nemotron BI small-M does not use CuTe GEMM, FP8 bmm pinned to cutlass, W4A16 uses Humming | tcgen05 hand-written skinny GEMM: a 64-row UMMA tile under "TileM must not depend on M" caps tensor core utilization at 50%; large M with narrow N is about 2× cuBLAS; loading the skinny bi_gemm `.so` on DS4 BI=1 does not change throughput | The small-M gain comes from cluster split-K over DSMEM (S=8 at K=4096), removing the producer/MMA warp `try_wait` peek (~0.5 µs), and `st.async` sent directly to the owner slot (~0.5 µs at M=64); stage count, spin `test_wait` and S=16 gave nothing, do not retry |
| mHC (DS4 TileLang) | BI=1 pins `split_tokens=1` | no extra cost | Pitfall only under BI=0: `n_splits` varies with `num_tokens` and is a `@tilelang.jit` specialization parameter, so every unseen prefill size recompiles for 10–14 s, all DP ranks stall, and every BI=0 comparison is polluted. BI=0 must load the `fixed_splitk` patch and check `tl_after_bench=0`. `mhc_post` spreading h over 128 blocks is numerically right but 40% slower; gains can only come from fusion |
| MoE (NVFP4 W4A16, Nemotron) | CuTe-DSL W4A16 modular MoE (takes the router's top-k); one CTA runs the whole K (`_W4A16_K_TILE=256` fixed), no split-K, no atomic | Per-row K accumulation order is the same under both tile sets; whether a single FP32 element is bitwise identical across tcgen05 instruction shapes (M128/M256, N8/N128) is undocumented and must be measured | flashinfer_cutedsl and Humming are both modular; moving R3 routing capture into the kernel (monolithic, #44214 style) saves nothing |
| MoE (FP8 block, DS4) | backend locked to DEEPGEMM (open source, built-in BI mode, the deployment kernel; FlashInfer TRT-LLM MoE ships as precompiled cubins that cannot be modified or verified as BI); `flashinfer_cutlass` rejects the block-fp8 QuantKey, `triton` resolves to BATCHED_TRITON which does not declare BI → a2a locked to deepep low-latency | BI=1 vs BI=0 gap floor is 33%, 83% of it non-overlappable deepep LL communication (86 serial a2a round trips per decode step); BI=1 does 14% less GPU work than BI=0 but takes 28% more wall clock (overlap ratio 1.33 vs 1.97) | DBO measured −39% (split ubatches read expert weights twice per layer, and weight reads are already at 84% of HBM roofline); TP2/DP4 −17%, TP4/DP2 −28%, DP8 is optimal for BI=1 too; deterministic top-k + histogram radix-select together +6.7% (alone +2.3% / +0.1%, super-additive, measure together); the biggest lever is DeepEP LL over MNNVL instead of RoCE (`VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL=1`, `DEEPEP_MAX_NVL_PEERS=8`): gap 33% (start) → 17% (after kernel PRs) → 4.5% (paired, same fabric), finally BI=1 ahead by 3.2%; transport does not touch BI (output bitwise identical across transports) |
| Stable top-k / routing scores | BI grouped top-k, tie-break by request-local index | Deterministic top-k (radix above 4K) 5.9 µs vs a histogram radix-select 56.7 µs | Full output contract below; "select candidates unordered, then sort" is correct only if the candidate set itself meets the contract for cutoff ties; a final stable sort cannot repair an earlier step that randomly dropped equal-score candidates |
| Mamba2 (SSD scan + conv1d) | exact replay ported from vllm#55905 (three extra chunk buffers x/dt/B); a partial chunk is replayed from the FP32 boundary state through the same SSD path | row-gated emit (emit only outputs needed now, update state at chunk boundary) restores Graph, but measured 0.1874 ms/token vs BI0 0.1726, still 8.57% slower; not default | Under BI Mamba2 declares 5 states (conv + temporal×4) but align-mode PC copied only 2 → after a hit, every prompt longer than one block is wrong; now copies the full per-layer declared state list, with a unit test asserting equal counts. The conv window is written tail-aligned by `conv_state.size(-1)` while the kernel reads `[0, width-1)`; when spec decode widens the cache they misalign; write only `[0, conv_kernel_size-1)`. `test_mamba2_ssd_emit` is occasionally flaky because the reference's autotune picks different blocking; pin the reference to the BI config. Only chunk64/8 heads exposes that the cumsum tile changes FP32 prefix sums (#55895). Gating inside the FP32 scan vs rounding first then gated RMSNorm round differently (#56244). Merging #58851 × #57329 directly: 28-shot GSM8K 432 → 4/500, because later groups reused the first group's checkpoint target block; 438 after adding the per-group gather. Under BI Mamba2 rejects the KV connector; real PD disaggregation is not covered |
| RMSNorm / gated RMSNorm | `cuda_rms_norm` (`CudaRMSNorm`), `grouped_gated_rms_norm` custom op, norm/act+quant fusion off, standalone `+quant_fp8` | Training-side TE fused RMSNorm differs from vLLM native RMSNorm by 1 ulp → only 74.4% of in_proj GEMM inputs bitwise identical and 1.7% of FP8 codes flip; so the training side aligns with the vLLM kernel instead of te.RMSNorm | Direct and compiled paths of the same gated RMS kernel differ; the difference disappears behind an opaque custom-op boundary. Review comment: this is a second BI RMSNorm implementation; DS4 reuses `rms_norm_batch_invariant` |
| FP8 activation quantization | standalone `+quant_fp8` op, `use_cuda_fp8_quant` under BI | Audit scale grouping, amax, rounding, clamp and aggregation order separately; the communication mode does not guarantee quant BI | The DS4 vLLM fork's `is_batch_invariant_quant_kernel_enabled()` only checks that the `.so` exists, not the BI switch, so BI=0 runs the unoptimized BI kernel for nothing (nsys 9.9%, 35.7 µs); with the fixed `.so` + `VLLM_BI_QUANT_TOKEN_BLOCKS=128` it is 5.3 µs, DS4 16703 → 18313 tok/s. Before a BI=0 comparison, grep nsys for `vllm::batch_invariant::`, or set `VLLM_BATCH_INVARIANT_KERNEL_LIB=/nonexistent` to fall back to upstream |
| cudagraph | pinned to FULL under BI (on DS4 FULL_AND_PIECEWISE is fastest: 16409 vs FULL 16068 vs PIECEWISE 10423 tok/s, BI=1 bitwise PASS) | FULL has worse TTFT (1792 vs 764 ms): on DS4 FULL lands as FULL_DECODE_ONLY and mixed prefill is always eager | Piecewise breaks BI through the switch "does this step run the graph or fall back to eager"; the criterion is total token count ≤ `max_cudagraph_capture_size`. Graph and eager paths are each BI but not bitwise identical to each other. Unrelated to padding. It reappears under FULL when the capture limit is set below `max_num_seqs × decode_query_len`. A short-prompt probe passing in all three modes does not prove PIECEWISE safe |
| determinism | `torch.use_deterministic_algorithms(True)` + `CUBLAS_WORKSPACE_CONFIG` (rewritten to `:16:8` by vLLM BI init); the mamba_ssm SSD backward uses atomics for dA/dD/ddt_bias, call `set_deterministic_mode(True)` explicitly | affects only backward step 0 | see the section below |
| KV block zeroing (DS4 multi-segment KV) | unrelated to BI, but both arms need it | The DS4 vLLM fork's `KVBlockZeroer` picks a 16 B program block by the gcd of per-segment page sizes, 1–12 ms per step (2.27 ms/step = 7.6% of wall clock); upstream #52058 uses a 3-D grid + masked tail; chunk mode 60 µs | Visible only as `_zero_kv_blocks*` in nsys; for a memset-like kernel taking ms, first check `blk_size` ≤16 |
| Prefix caching (Mamba align mode, BI=0) | precompute multi-group alignment indices once (#58851 fits best), skip align postprocess without spec decode, skip preprocess/precopy when the host determines nothing crosses a block; on prefill export block-boundary checkpoints inside the scan (#57329) | PC on as is: BI0 throughput −4 to −20%, TPOT +13 to 21%, c=1 TTFT roughly doubles (eager launches per step 9→37, 8K prefill split in two); after the fix BI1 is about on par with PC off | The BI path must export checkpoints separately or BI1 is silently wrong; a chunk truncated by the token budget must also export its last checkpoint (under DP+EP one extra prefill step on any rank stalls global decode) |
## Reference: DS4-Flash per-operator choices
The full per-operator table behind the blog's summary table (consistent arm = BI=1). "Same as left" means the training forward calls the same vLLM kernel as rollout.
### Dense, norm and mHC
| Operator | rollout (consistent) | Training | What is fixed / given up; cost |
| --- | --- | --- | --- |
| Embedding | Table lookup | Forward: Same as left<br>Backward: PyTorch | Naturally consistent |
| mHC residual mixing (pre / post / Sinkhorn) | Split TileLang mHC kernels, `split_tokens=1` | Forward: Same as left<br>Backward: Hand-written VJP, Sinkhorn iterations unrolled, FP32 | Gives up the baseline's fused post+pre kernel (split count varies with token count)<br>Cost: Both sides lose one fusion |
| RMSNorm | Batch-invariant RMSNorm (one program per row) | Forward: Same as left<br>Backward: FP32 VJP | Per-row reduction order fixed<br>Cost: Small |
| FP8 linear layers (q_a/kv, q_b, wo_b, shared expert) | DeepGEMM FP8 GEMM, batch-invariant mode; weights in 128×128 blocks, activations quantized dynamically per token in groups of 128 | Forward: Quantizes BF16 master weights on the fly to the same block FP8, same DeepGEMM<br>Backward: Transformer Engine (TE) BF16 GEMM w.r.t. BF16 master weights (STE) | Gives up the baseline's FlashInfer switch at M<32; scale is UE8M0 on SM100 (8-bit, exponent only, i.e. a power of 2), FP32 on SM90<br>Cost: rollout: no faster small-batch kernel; training: FP8 weights repacked every forward (including recomputation) |
| Small linear layers (compressor, indexer weight projection) | FP32-output cuBLAS GEMM / batch-invariant persistent Triton matmul | Forward: Same as left<br>Backward: BF16 | Summation order along K fixed<br>Cost: Small |
### Attention (DSA sparse attention + FP8 KV cache)
| Operator | rollout (consistent) | Training | What is fixed / given up; cost |
| --- | --- | --- | --- |
| q-norm + RoPE + KV write | Fused kernel quantizes K and writes the FP8 cache | Forward: Same kernel writes a private training-side FP8 cache, then reads it back<br>Backward: FP32 RoPE / norm in PyTorch (straight-through) | K takes the same FP8 round trip<br>Cost: Training: one extra cache and dequantization |
| FP8 KV cache | Stored as FP8; prefill and decode dequantize to BF16 before attention | Forward: Same as left<br>Backward: — | Gives up the default FP8-reading decode kernel (blog 4.2)<br>Cost: rollout decode |
| CSA compressor | vLLM compression kernel, state reset per request | Forward: Same kernel's values, functional graph's gradients<br>Backward: autograd | In-request state starts as in prefill<br>Cost: Training: loop over requests |
| DSA indexer scoring | DeepGEMM FP8 MQA logits; prefill reads contiguous KV, decode reads paged KV | Forward: Same as left, per request<br>Backward: None (indexer not trained) | Paged and contiguous reads bitwise identical |
| Top-K selection + index arrangement | Batch-invariant Top-K (64-bit key + radix select), sorted descending, invalid entries last | Forward: Same as left<br>Backward: None | Ties go to the larger in-request index (blog 4.3)<br>Cost: One extra sort per side, recovered by fusion |
| Sparse MLA | `flash_mla_sparse_fwd`; prefill one segment per request, decode same sparse kernel | Forward: Same as left, once per request<br>Backward: cuDNN frontend DSA backward | In-request geometry ignores neighbors; decode emits 1 token per request per call<br>Cost: rollout: multiple prefill launches; training: serial over requests |
| o_proj | SM100: DeepGEMM FP8 einsum; SM90: grouped Triton W8A8 | Forward: Same as left<br>Backward: TE BF16 grouped GEMM + inverse-RoPE VJP | Aligned per architecture; no bitwise identity across architectures<br>Cost: Loop over groups on SM90 |
### MoE
| Operator | rollout (consistent) | Training | What is fixed / given up; cost |
| --- | --- | --- | --- |
| Router scoring + expert selection | Same FP32-output gate GEMM for all batch sizes; sqrt-softplus + bias selects top-k | Forward: Same as left, routing recomputed bitwise<br>Backward: Expert IDs fixed, FP32 gradient w.r.t. routing weights | Gives up the small-batch gate kernel; **R3 is not needed** (blog 4.4)<br>Cost: Small |
| Dispatch | DeepEP low-latency (LL below), fixed per-rank capacity, all ranks segment in sync; FP8 quantization uses elementwise round-to-nearest | Forward: DeepEP normal, plus one dispatch of routing metadata (fingerprint, expert, weight)<br>Backward: DeepEP backward | Experts get bitwise-identical FP8 values and scales either way<br>Cost: rollout: LL communication cannot overlap compute; training: metadata dispatch and hash matching |
| SiLU · mul · quantization | Batch-invariant fused kernel: product rounded once to BF16, then quantized in groups of 128 | Forward: Same as left (mandatory)<br>Backward: FP32 VJP | Rounding position fixed (blog 4.4)<br>Cost: Small after optimization |
| Routed expert GEMM | DeepGEMM masked grouped GEMM, static capacity as scheduling hint | Forward: DeepGEMM contiguous grouped GEMM, per-expert rows aligned to 128<br>Backward: TE BF16 grouped GEMM | Gives up the baseline's FlashInfer TRT-LLM MoE (ships as precompiled cubins; cannot be modified or verified as batch-invariant)<br>Cost: rollout: a main throughput gap; training: alignment padding |
| Combine | DeepEP LL combine, fixed merge order | Forward: FP32 accumulation in top-k slot order on the token's owning rank, then BF16<br>Backward: Gradient scatter | Merge order by slot, not arrival order or rank<br>Cost: 43 layers × 2 serial all-to-all per step |
| Shared expert | FP8 linear layer; auxiliary stream under LL + MNNVL (multi-node NVLink) | Forward: Same as left<br>Backward: BF16 (STE) | Stream changes overlap, not values<br>Cost: A source of recovered speed |
### Output layer and infrastructure
| Item | rollout (consistent) | Training | What is fixed / given up; cost |
| --- | --- | --- | --- |
| lm_head + logprob | Batch-invariant matmul; FP32 batch-invariant log-softmax | Logits per 8192 tokens, vLLM logprob value; backward via PyTorch log-softmax on the same logits (straight-through) | temperature must be a scalar<br>Cost: Training: each chunk's FP32 log-softmax kept for backward |
| Cross-GPU reduction | TP=1, no all-reduce; NCCL channel count and algorithm pinned as a guard | TP=1; CP exchanges boundaries only, no reduction | TP>1 not supported (blog 4.5)<br>Cost: Cannot split with TP |
| CUDA Graph | Kept, with the right mode (blog 4.6) | — | Graph and eager paths bitwise identical |
| Weight sync | Reload copies new values to original addresses, refreshes derived weights | Exports block-FP8 weights | Graph-referenced addresses unchanged<br>Cost: One export and transfer per step |
Details that back the table:
- **Scales and precision.** DeepGEMM scales are powers of 2: UE8M0-packed on SM100 (a Triton kernel emits packed weights), FP32 on SM90 (vLLM weight post-processing). Router gate and compressor projection output FP32; indexer weight projection and lm_head are BF16. Before the first update the checkpoint scales are reused; afterwards they are recomputed from amax. FP8 weights are not cached by default, so every forward including recompute repacks them.
- **FP8 KV layout.** Per token per layer: 448 FP8 NoPE dims, 64 BF16 RoPE dims, one UE8M0 scale per 64 elements, 1 alignment byte = 584 bytes (BF16: 1024). The sparse decode kernel must be selected explicitly; the BI switch alone does not change the default FP8-reading FlashMLA decode. The CSA compressor zeroes state at each request start.
- **Top-K key.** 64-bit key: high 32 bits = order-preserving transform of the score, low 32 bits = request-local index; radix select at 8 bits per round. Deterministic Top-K and the Top-K + index-sort fusion gave +2.3% and +0.1% alone but +6.7% together; the interaction is unexplained, so measure them together. Under DCP>1 the cross-GPU candidate merge uses the opposite tie rule; unsupported. The indexer auxiliary loss does not support CP>1.
- **Expert GEMM.** Contiguous layout aligns each expert's rows to 128. LL quantization switched to elementwise round-to-nearest; training quantizes the contiguous layout under the LL dispatch convention. The SiLU·mul·quant kernel has fixed launch geometry, is mandatory on both sides (training refuses to start without it), and uses different SiLU formulations on SM100 and SM90 (no cross-architecture bitwise promise). Backward: TE BF16 grouped GEMM, second GEMM dgrad/wgrad first, SwiGLU recomputed from saved gate_up in FP32.
- **Layouts tried for performance only.** TP2/DP4 17% slower, TP4/DP2 28% slower; DP8 is best for BI=1.
- **Cross-GPU.** Training dispatches routing metadata once more (per route: 16-element token fingerprint, expert index, FP32 routing weight), matches outputs by bitwise hash, and scatters with a deterministic Triton kernel. NCCL channel count, thread count, protocol and algorithm are pinned and symmetric-memory / custom all-reduce are off; under TP=1 this is a guard, not the mechanism. CP pads each sequence, splits contiguously, exchanges boundary hidden states and all-gathers compressed KV rows.
- **CUDA Graph.** Workspaces are allocated in warmup before capture and locked afterwards; the Top-K row-start buffer is a resident zero buffer. Weight-derived mHC parameters are recomputed in place on reload.
## The full output contract of stable Top-K
What orders equal scores (e.g. the request-local index)? Take it from the training side or current baseline; do not invent a tie-break. Preserve:
- Comparison key for equal scores, positive and negative zero, NaN/Inf (if allowed); a global batch offset cannot replace the request-local index.
- Output indices and their order, not just the same set; swapping equal-score indices changes later floating-point reduction order.
- valid length, mask, invalid sentinel, K larger than the valid count, empty rows, padding, tail blocks.
- prefill/decode index coordinates, cache page mapping, the position of the same semantic token.
Check the candidate for repeated allocation, sort input range, radix/merge passes, workspace, launch/grid/block; prove order preservation before deleting a pass. For any CUB/custom sort or selection, verify the binding and the real lifetime of its temporary storage; a library's stability claim does not replace integration tests. A pre-generated index/workspace saves overhead but must not overwrite in-use content under multiple streams, concurrent requests or Graph replay.
## Index / elementwise fusion
Limit scope to confirmed hotspots such as gather lengths, row start/end, sparse index concatenation and invalid-row zeroing. For every row write down input coordinates, whether the interval is half-open (closed left, open right), output capacity, valid range, and who writes the padding.
Fusion can save a launch or temporary tensor, but can also raise register pressure or hurt memory access. Measure the gain on all real shapes; the number of changed lines says nothing. Shrink the grid and loop stride together so the full output is covered; fill a reused buffer with poison values to catch unwritten tail/empty rows.
Integer indices are not automatically correct: off-by-one, request offset, CP boundary, overflow and uninitialized rows can bias a whole sequence. If fusion reaches into dequant/norm, keep the original dtype conversions and rounding positions. Unit-test Top-K and index post-processing separately, then combined, then module forward, then engine gates. Record whether the optimization also speeds up BI0: dynamic block/grid selection on a shared path is a general optimization and cannot all be counted as BI-specific overhead reduction.
## Split-K and the M of grouped GEMM
A GEMM's M is usually the number of rows processed; masked grouped GEMM computes several experts in groups and uses only each group's valid token rows. `expert_num_tokens.max().item()` pulls the device max token count to the host, commonly as a capacity/scheduling hint; it does not "fix the shape" or "guarantee BI". The readback synchronizes and usually cannot be used in CUDA Graph capture. Before removing it, read the API: is the parameter a correctness bound, a capacity, or a tuning hint? Restore an estimate/safe upper bound the library supports, and cover heavy load and empty experts. Do not claim without source evidence that underestimating M makes the default kernel access memory illegally, and do not set an arbitrarily small value for speed.
One launch per token avoids some batch-dependent reduction/scheduling but means many launches and low utilization. Batched launch with fixed `n_split` may keep the per-row numerical order, but the K range of each split, the accumulation order within it, the accumulator dtype, partial reduce order, tail/padding and actual kernel specialization must still be fixed or verified. "Fixed K/n_split" is a candidate mechanism, not a proof of prefill=decode: different tiles, epilogues, precision or preprocessing can still differ at the same split count. Test identical operator inputs and the full semantic path.
## BI coverage of attention
Attention BI covers QK, online softmax, PV and the merge order of blocks/splits. Beyond the split count, the tile scheduler, Q-stage, cluster/scheduling choice, page layout, mask/sink, score scale and accumulator/rounding can depend on batch/length. Passing small batches does not cover long context, prefill chunk and decode, Graph buckets, mixed batches, page boundaries, prefix cache hit/miss. Graph+prefix enabled does not mean one test executed all these branches; record the coverage.
## Observation in Graph mode
Save the resolved Graph mode, capture sizes, prefill/decode/mixed step types and evidence of kernel calls. A successful capture guarantees neither eager-equal numerics nor that every step uses the same kind of graph. A Python hook usually runs at capture, not replay; a capture dump cannot stand for all replays.
To observe replay, use controlled device-side outputs or static diagnostic buffers, and state whether instrumentation changes the capture. Eager hooks can localize, but the final recheck must return to the original Graph mode. Even if the compiled path with Graph off matches production Graph on fixed samples, eager may still differ; keep artifacts for all three. A hook not firing on the compiled path does not mean the layer did not execute.
Turning the graph off changes speed and therefore online batching. First fix each step's tokens, requests/neighbors, padding and cache state, then compare the supported modes (eager/FULL/PIECEWISE). Use a 2×2 experiment to rule out the dispatch/split plan the switch changes at the same time. A short filler may put a mixed step into a capture bucket while a long filler goes eager. Prefill/decode overlap differs between the full and a small model; a small model passing does not cover untriggered production mixed steps. For intermittent failures keep multi-round samples and the trigger rate; do not pick only successful rounds.
<a id="determinism"></a>
## determinism takes effect only on the main thread
In PyTorch 2.13 (nv26.07) Inductor config overrides are thread-local (ContextVar). `torch.use_deterministic_algorithms(True)` sets `torch._inductor.config.deterministic` on the main thread; on autograd's CUDA device thread it stays False. A `torch.compile`d norm VJP first compiles on the autograd thread during backward, sees `deterministic=False`, and the generated looped-reduction kernel enables `dynamic_scale_rblock`: besides `R0_BLOCK=1024` it adds candidate 512, and the first launch picks by timing. The two reduce in different orders and differ at bit level. Only at the end of the dynamo compile is that thread set to deterministic; later compiles have one fixed config. Result: only backward step 0 is not bitwise reproducible across runs; forward logprob and the RL closed loop are unaffected.
Apply both fixes: `torch.compile(..., options={"deterministic": True})` at the compile site (environment-independent), and `TORCHINDUCTOR_DETERMINISTIC=1` in the recipe environment (covers any first compile on the autograd thread, including the vLLM forward function compiled inside recompute). Also: verl's full_determinism actually sets `warn_only=True`, and the mamba_ssm SSD backward must enable deterministic mode explicitly.
## How to break down the cost from BI0 to BI1
Draw the actual call chain first: what the default fast kernel, the BI branch and the extra preprocessing/sorting/postprocessing each are. Do not infer "split-K was changed" or "a temporary Triton was used" from function names; compare source, binding, compiled artifacts and real launches.
You may copy the BI0 implementation into a standalone candidate and add the deterministic/BI conditions one at a time for a bisectable A/B. It need not live in the same file, but the default path must stay unchanged. If the BI0 algorithm cannot meet the required tie/reduction contract, switching algorithms is allowed; state the cost.
Fairness between the arms comes before any kernel conclusion: same fabric (on DS4 two environment variables closed most of the remaining gap), same node, same image, BI=0 not running the BI kernel, no JIT stall, and the KV zeroing fix in both arms. Nemotron's BI1 vs BI0 E2E throughput loss is the median over 5 pairs, with single pairs reported separately.
## Minimal record for each candidate
```text
baseline / candidate source+binary hashes
one intended change; BI-specific or common
same-input bitwise + BI composition + prefill/decode gates
BI0 unchanged + real shape/tie/mask/tail + Graph/workspace reuse
paired latency, launch/allocation counts, repetitions
pass / fail / inconclusive; retained or rejected; remaining scope
```
On failure, save the minimal counterexample, then withdraw from the default path. When the user requires all micro gates to pass before E2E, do not skip them; performance numbers do not cover a correctness failure. Do not quietly remove an existing stable-sort patch to reproduce a long-length mismatch; that tests a regression, not a defect of the current best code.
.agents/skills/train-inference-consistency/guides/checkpoint-resume.md
-304
View File
@@ -1,304 +0,0 @@
# Checkpoint, resume and memory budget
Use this guide to get checkpoint save, load and resume right in long RL runs, and to budget host (CPU) memory peaks and growth. For per-phase GPU memory peaks see [Locating GPU OOM and staging peaks by lifecycle](#locating-gpu-oom-and-staging-peaks-by-lifecycle).
## Find by symptom
| Symptom | Where to go |
| --- | --- |
| OOM while loading a checkpoint | [Load](#4-load-enough-capacity-does-not-mean-enough-for-the-transient-peak) |
| Train-inference mismatch on the first step after restore | [13](#13-stale-quantization-scale-on-the-first-step-after-restore) |
| Resume is set but training starts from step 0 | [14](#14-resume_modeauto-starts-from-step-0) |
| Host RSS grows every step | [3.3](#33-locating-slowly-grows-every-step) |
| Planning host memory for a checkpoint in the hundreds of GB | [4](#4-capacity-planning) |
| vLLM fails to start when colocated with training (KV budget too small) | [5](#5-gpu-memory-set-gpu_memory_utilization-per-mode-when-training-and-rollout-are-colocated) |
| GPU OOM during init / export / wake | [Locating GPU OOM](#locating-gpu-oom-and-staging-peaks-by-lifecycle) |
## Terms
- **pinned memory:** page-locked host memory for direct GPU DMA (`cudaHostAlloc` / `cudaHostRegister`), cached by PyTorch's caching host allocator.
- **cgroup:** limit `memory.max`, usage `memory.current`, split in `memory.stat` into anon, file (page cache) and shmem.
- **RSS and PSS:** RSS counts shared pages fully in every process; PSS splits them, so only summed PSS approximates the physical total.
- **DCP / Megatron dist_ckpt:** PyTorch and Megatron distributed checkpoint formats; each rank writes its shard, `.metadata` last.
- **staging / D2H:** the host buffer GPU tensors are copied into before writing to disk.
- **Ray object store (plasma):** Ray's shared memory for passing objects, mapped into each worker and charged via PSS.
- **old-logprob:** the logprob training recomputes for rollout tokens with the pre-update policy.
- **Ray OOM threshold:** `RAY_memory_usage_threshold`, the usage fraction at which Ray kills workers. Default 0.95.
## Key rules
- Check file, memory and derived state separately; complete files can still mismatch on the first step after restore.
- Capture the first exception (errno, rank, file, offset), not the last message.
- Completion marker last; clear the old directory before overwriting; all ranks check and fail together; a missing file raises instead of falling back to a fresh optimizer.
- Under EP4 + FSDP2, ranks' expert shards in DCP share key and offset, so restore reads another rank's experts. Put the EP rank in the key (port the reference fix), add a 4-GPU bitwise read-back unit test, and raise on EP size mismatch.
- Rotation deletes only tracked, complete, idle directories; in long runs keep one checkpoint every 20 steps outside rotation.
- Set rollout `gpu_memory_utilization` per training mode ([5](#5-gpu-memory-set-gpu_memory_utilization-per-mode-when-training-and-rollout-are-colocated)).
- Save fixes pass [the acceptance in 3.2](#32-reproduce-and-discriminate). Before a long run on a new save path or layout, run to the first save and do one kill→resume.
- Resume is proven by matching per-rank digests in a new process, a byte-identical first export, an aligned first old-logprob and one non-zero update; `resume_mode=auto` alone proves nothing.
- Read host memory as anon/file/shmem; `memory.current` includes page cache, and Ray kills on it.
- Pinned buffers round up to a power of 2 (Nemotron: 146.9 GiB pinned for 88.2 GiB of optimizer state, +67%); allocate large long-lived pinned blocks at exact size.
- The save step is the peak; a higher Ray threshold or smaller microbatch does not fix headroom.
- For per-step growth, use the per-process PSS floor and a small CPU-only repro before changing code.
- Verify per save path that the host cache is released after save.
## Checkpoint and resume
### 1. Three independent contracts
Classify the problem before calling it "the checkpoint is broken":
| Contract | Typical failure |
| --- | --- |
| File: serialization and publish | Disk write error; an incomplete directory treated as complete |
| Memory: transient peaks of save and load | Killed by Ray after the write finishes; cgroup OOM during load |
| Derived state: the first export after restore | Files complete, but the first step after restore mismatches ([stale quantization scale](#13-stale-quantization-scale-on-the-first-step-after-restore)) |
The third class is a derived-state lifecycle bug, not file corruption.
### 2. Which state exact resume needs
| State | Requirement |
| --- | --- |
| Model parameters or LoRA adapter shards | Saved per rank; strict load that rejects missing/unexpected keys ([training forward](training-forward.md)) |
| Optimizer state, including the CPU-offloaded part | Same step as the model. Nemotron chunked optimizer offload: 88.2 GiB per rank, 94.7 GB per rank on disk; DS4.1 saves LoRA parameters and Adam `exp_avg` / `exp_avg_sq` in FP32 |
| LR scheduler | Check the saved value, not the source default (DS4.1 checks its constant, no-warmup schedule per rank) |
| RNG | Four streams: CPU, CUDA, Python, NumPy |
| Counters and data position | global step, optimizer step, data cursor; DS4.1 also saves the cumulative count of effective-signal steps |
| Run contract | Hashes of code, data and recipe, compared item by item on restore; reject a change in total steps or data by default. To extend training, declare a new phase explicitly; do not edit the old contract file or add a generic ignore switch |
| Completion marker | `.metadata` / tracker (e.g. `latest_checkpointed_iteration.txt`), written last, see 3.4 |
Not in the checkpoint, rebuilt after restore:
- **Parameter-derived caches:** quantization scale, packed/transposed/fused weight, version markers (see 5).
- **Records outside training:** the metrics jsonl is overwritten on resume.
### 3. Save
#### 3.1 Capture the first error, do not fix the last error message
Keep the writer callback's original exception / errno, rank, path, offset, requested / returned bytes and exception chain, which DCP and Ray wrappers may drop; also filesystem/mount, free space, quota and cgroup OOM events, not just the final `ActorDiedError`. `unexpected pos` can come from disk space, quota, permissions or I/O failure; the first exception tells them apart.
#### 3.2 Reproduce and discriminate
A fast repro keeps the real tensors, serializer and concurrency, skips rollout/backward and repeats only save, for a bounded number of rounds sized to the failure probability; stop at the first failure and keep the scene. DS4 model-only save failed on the 3rd attempt on one path and the 6th on another; Nemotron with host cache release failed on the 2nd-4th save.
**Acceptance for a save fix: N consecutive saves + kill→resume→save again.** Also run it before the first long run of a new save path, parallel layout or node count; running to the first save point is enough.
Order of discrimination:
1. Same buffer/serializer, target filesystem versus node-local disk as the control; read the errno of the first failure.
2. Control concurrency, pinned/pageable and memory peak; if the serializer changed along with the filesystem, the result cannot be attributed to one factor.
3. Check growth of threads, processes and file handles, separately from the write error.
4. Only then read kernel / filesystem client code; a matching historical issue is supporting evidence, not proof of a fix.
#### 3.3 Reliable publish
Serialize on node-local disk and close the file; copy to a **temporary name on the target filesystem**, handle short copies, flush/fsync as durability requires, verify integrity, then rename within the same filesystem. Publish metadata/tracker/completion marker only after all ranks finish. A move across filesystems is not an atomic rename.
#### 3.4 Commit protocol: write the completion marker last
| Rule | If you do not |
| --- | --- |
| Write `.metadata` / tracker last, after every rank's shards and auxiliary files (optimizer, RNG, scheduler) succeed and are confirmed collectively | A rank fails, the marker already exists, and restore reads incomplete state |
| Before overwriting a step directory, delete all old files first (uncommit) | Old optimizer / RNG silently pair with the new model |
| Before restore, all ranks check file completeness together and fail together | The rank missing a file exits; the others hang until the NCCL timeout |
| Ranks with no shard to write still join the same collective | Collective order diverges across ranks and hangs |
| A requested load with a missing file raises an error, not just a log line | Training silently continues with a fresh optimizer |
Acceptance: a save failing midway (Nemotron hit a real ENOSPC) makes all ranks exit together without advancing the tracker, and training resumes directly; new code loads an old-code checkpoint. Never kill a job mid-write; verify the checkpoint is complete first.
#### 3.5 Verify a written checkpoint
Check the expected shard set, non-empty sizes, leftover temporary files, checksums where needed, and do one real full restore. A correct save digest is not a restore test (DS4.1 checkpoints 20/24 had correct save digests but were never restored in a new process, and the report says so).
### 4. Load: enough capacity does not mean enough for the transient peak
Load may copy the full optimizer into anonymous memory, and FSDP CPU offload may create a second set of pinned storage. Accept by loading model + optimizer + RNG/extra at real scale, not one small shard.
Read a checkpoint with the image that wrote it: DS4.1 host PyTorch 2.11 cannot deserialize the DTensor `_FlatLayout` written by the image's 2.13; it reads only inside the image with `torch.load(weights_only=True, mmap=True, map_location="cpu")`. Do not modify the checkpoint or turn off `weights_only` to make it load.
### 5. Derived state after restore
DS4 chain: initial HF load saves FP8 source scales → DCP restores the trained BF16 parameters → the first export still uses the old registry → the actor re-quantizes by parameter version, but rollout uses the old scales. The first resumed step mismatches; once the next optimizer hook clears the cache it matches. The prediction "it recovers on the next step" plus a before/after-fix comparison on the same checkpoint locates the cause.
Fix it in the lifecycle that owns the state: after load/update, invalidate or rebuild parameter-derived caches (scale, packed/transposed/fused weight, version markers), reusing an existing hook without side effects such as a real optimizer update. Do not hide it by waiting a step, disabling resume or clearing all caches by hand. DS4's minimal patch calls the existing post-update hook ([Megatron-LM #197](https://github.com/ISEEKYAN/Megatron-LM/pull/197)). List every parameter-derived cache and its invalidation code, add the missing calls and one regression test.
### 6. How to prove resume is correct
1. **Before launch.** For node-local checkpoints, every rank's node in the new allocation has its shard; the tracker / latest pointer points to a complete checkpoint; the launcher's resolved `resume_from_path` and global step are as expected. Confirm the restored step before the first training log line; `resume_mode=auto` is not evidence.
2. **New-process restore.** All worker PIDs changed; per-rank adapter / optimizer / scheduler / RNG digests match the save-time digests item by item.
3. **First export and old-logprob.** Exported files are byte-identical to those of the saving step; also check the exporter scale / policy fingerprint and the full valid logprob. Checking only that the second step returns to 0 misses the bug in 5.
4. **One non-zero update.** A zero-advantage step only exercises the flow (two DS4.1 resume segments had zero gradient on all ranks at step 2, proving nothing about restored adapter values).
5. **Gate coverage.** step0 cold load, first export/old-logprob after restoring a non-zero step, one later update, address / state validity under Graph and sleep/wake, new code reading an old checkpoint.
6. **What acceptance compares.** Restored per-rank digests equal save-time digests; first exported files are byte-identical to the saving step's; per-step `probs_diff` after restore is 0 under BI=1 (Nemotron: 0 on the first step after kill→`resume_mode=auto`). The resumed trajectory need not match an uninterrupted run bitwise (rollout sampling is not bitwise across processes). There is no "metrics are close" threshold; reward curves are reference only, on the same step range and basis.
### 7. Rotation, disk and time cost
| Item | Record |
| --- | --- |
| Size per checkpoint and save time | DS4 full32: 99 files, about 3.2 TiB, save about 269 s. Nemotron-H 30B-A3B, 4×GB200: 353 GB (EP4 + FSDP2) or 429 GB (PP4 dist_opt), save 72-152 s, about 2-3% of step time when saving every 5 steps. Layout-specific; state these for your model in the plan |
| Rotation | The mLite engine ignores `max_actor_ckpt_to_keep` |
| Checkpoint used for eval | Rotation may have deleted the step to evaluate; a checkpoint from another chain cannot substitute |
| Resume on a different node after node failure | Copy checkpoint, dequantization cache, pinned image and dependency cache to the new node before launch; disable automatic requeue; change only the node selection; write a transfer record |
| Paths for two-arm experiments | BI=0 / BI=1 output and checkpoint paths are separate; renaming the run breaks a resume chain in progress ([recipe and acceptance](recipe-and-acceptance.md#rules)) |
## Memory budget: host and GPU
Host memory fails in two ways in long RL runs: phase **peaks** (save / load / sleep) hitting the limit, and slow per-step **growth**; they are measured differently (1-4). GPU memory is a third class: with training and rollout colocated, rollout `gpu_memory_utilization` is set per training mode (5).
### 1. Who holds host memory
| Source | Mechanism |
| --- | --- |
| Pinned buffer of optimizer offload | The caching host allocator rounds up to a power of 2, so the footprint is tens of percent above the nominal state; summed across ranks it pushes residency before save close to the limit |
| Parameter shards of FSDP CPU offload | Each rank's shard stays resident on CPU; doubles if load creates a second pinned copy |
| D2H staging during save | Stays in the pinned cache pool; whether `torch.accelerator.empty_host_cache()` releases it must be verified per save path |
| CPU backup of serving sleep | L1 sleep keeps one full backup per rank; the pinned cache is not returned after wake |
| Ray object store | plasma mappings count into each process's PSS as shared pages; node shmem does not grow with it (3) |
`param_offload` trades CPU pressure for GPU memory; optimizer offload is separate. Change one offload switch at a time and measure the phase peak before and after.
### 2. Limits and who does the killing
- **cgroup.** The job limit `memory.max` is usually below physical memory; budget against it. A cgroup OOM shows an `oom_kill` event and process SIGKILL, with no Python stack.
- **Ray memory monitor.** Kills workers at `RAY_memory_usage_threshold`=0.95 by default; log line `exceeded threshold of <bytes>B (95.0% ...)`. Its usage figure includes page cache. Raising the threshold does not fix headroom: Nemotron at 0.985 was still killed during save.
- Distinguish GPU allocation exception, host cgroup OOM and Ray kill. A CUDA allocator allocation-failure / retry warning is not a terminal OOM and the step may complete, but it does show memory pressure.
### 3. How to measure
#### 3.1 One record per phase
Log one line at load, save, sleep, wakeup, reset KV cache and generation start / end, with monotonic time, node / rank / PID and policy version. Per node, record cgroup `memory.current` / peak / events and `memory.stat` anon/file/shmem; sample process RSS/PSS, threads and FDs as needed, not with high-frequency scans of all processes.
`generate_sequences` includes wake/init. An error inside it does not prove token generation started, let alone a backward/optimizer OOM.
#### 3.2 Pitfalls in the readings
| Reading | Problem |
| --- | --- |
| `memory.current` | Includes page cache, which masks per-step growth; but it is what Ray and the save peak see |
| `free`, `psutil.virtual_memory().used` | System-visible memory, not process RSS or job cgroup usage |
| `scontrol` FreeMem | Excludes reclaimable cache |
| `ru_maxrss`, Slurm MaxRSS | Cumulative peak, not current RSS |
| VmPin/VmLck | Cannot alone infer CUDA pinned |
| Sum of per-process RSS | Counts shared mappings repeatedly; use PSS |
Pinned memory from cudaHostAlloc shows up as shmem in cgroup `memory.stat` (measured on Nemotron).
#### 3.3 Locating "slowly grows every step"
1. Use node `MemTotal−MemAvailable` and cgroup anon/file/shmem to find which class grows.
2. Sample PSS from `/proc/<pid>/smaps_rollup` per process once a minute, recording only comm / argv0; take each step's minimum as the floor and find processes that rise monotonically. Within-phase peaks swing with offload and are useless here.
3. For a rising process read `smaps`. On someone else's job or a running job, only read `/proc`; do not attach or send signals.
4. Reproduce the candidate mechanism in a small CPU-only process, then decide where to change code.
A floor that drops after resume (Nemotron run 1: about 455 back to about 433 GiB) shows the growth accumulates over the process lifetime.
### 4. Capacity planning
Budget the save-step peak as the sum of what is alive at the same time, against cgroup limit × Ray threshold:
```text
peak ≈ resident pinned optimizer offload (actual footprint after allocator rounding)
+ save D2H staging (and the part of it left in the cache pool)
+ page cache of the newly written checkpoint files
+ CPU backup of serving sleep and unreleased pinned cache
+ resident amount of the other processes (rollout, driver, Ray workers, etc.; take the floor per 3.3)
limit = cgroup memory.max × RAY_memory_usage_threshold
```
Use measured values wherever possible.
- **The save-step peak is closest to the limit.** The long-run plan states, per training mode, the measured cgroup peak at the save step and in the kill→resume segment, and which load is still unmeasured.
- **Estimate from headers before running.** Safetensors headers alone give the expanded size per dtype (full DS4.1 checkpoint: about 1421.66 GiB as BF16, including auxiliary content). It is not peak RSS; add serializer copies and mapped source pages.
- **A slope only hints at headroom risk.** Continue from a verified checkpoint and observe long enough; do not extrapolate an OOM step linearly from a short timeline.
- **On OOM, do not shrink the microbatch or raise the guard first.** Find the phase and which objects are alive together.
### 5. GPU memory: set gpu_memory_utilization per mode when training and rollout are colocated
The training process keeps part of its GPU memory during rollout. Nemotron single node 4×GB200:
| Training mode | Held per GPU by training during rollout | rollout `gpu_memory_utilization` |
| --- | --- | --- |
| EP4 + FSDP2 (mLite; BI=1 arm default `ACTOR_OPTIMIZER=fsdp2`) | Not measured separately | 0.65 (verified) |
| dist_opt (BI=1 arm optional `ACTOR_OPTIMIZER=dist_opt`; numbers from PP4 dist_opt measurements) | 54-64 GiB (FP32 quantization master resident on GPU) | 0.7, 0.6 and 0.65 fail to start; measured 0.55 |
| EP4 + Megatron (mbridge) dist_opt + chunked optimizer offload (BI=0 arm, QAT + R3) | Not measured separately | 0.45 (current recipe value) |
FlashInfer autotune warmup needs about 8 GiB more, another direct cause of failing to start. In a two-arm comparison, differences such as gpu_mem 0.45/0.65 are confounders and are listed first ([recipe and acceptance](recipe-and-acceptance.md)).
Rules:
- Write the value into the recipe per training mode, not as one global default; the recipe's hardware comment lists only verified topologies.
- Check the resolved value inside the container and the actual KV budget in the vLLM log, not the outer shell variable.
- `param_offload=True` saves GPU memory but moves parameters, changing the `data_ptr` of the quantization master so the pre-export check errors; as with DS4, enable only optimizer offload.
- What sleep level 1/2 release, and why wake may still find no GPU memory: [rollout engine §5](rollout-engine.md#5-lifecycle-and-gpu-memory-ledger) and [Locating GPU OOM](#locating-gpu-oom-and-staging-peaks-by-lifecycle).
## Locating GPU OOM and staging peaks by lifecycle
Use this when the shard estimate fits but init, export or wake still OOMs. Find which step allocates and what is still alive, and change that boundary; do not lower the batch or change sharding first, and never blame update for a failure before update.
### 1. The minimum to record
```text
phase / process / rank / topology
requested bytes / allocated / reserved / device free
host RSS / container or Ray limit
last completed phase / first failing call
```
Mark missing fields as not obtained; do not guess. If output just stops with no error, check CPU/GPU/I/O, log timestamps and phase (reading weights, compiling, collective, generating); slow is not a hang.
### 2. The memory peak is set by "the objects alive at the same time"
Parameters divided by GPU count estimates only part of the steady state and misses these peaks:
| Phase | Objects and boundaries to look for | Minimal verification |
| --- | --- | --- |
| Init / FSDP wrap | Full materialization before sharding; largest single parameter; CPU-resident/tied parameters | Whether meta → shard → load takes effect; per-rank peak |
| forward / backward | Full module from all-gather, activation, workspace | Measure logprob and actor update separately; frozen parameters may be gathered too |
| Export / merge / quantize | A single full tensor, fused groups, lazy generator references | The actual export set; release timing after yield/exception |
| compile / graph capture | autotune workspace and graph pool | Find the allocation in the log; eager only isolates, recheck the production graph separately |
| KV planning / wake | KV budget, trainer residency, allocator free cache | Resolved GMU/TP and allocator metrics, not the configured value |
| checkpoint / host | Full CPU copy, optimizer gather, shared mappings, Ray guard | Per-process and node usage; RSS does not sum to the physical total |
A small model with the same special layers, attention period and scale layout verifies the mechanism; only the full model verifies capacity, such as the all-gather peak of a huge single parameter.
### 3. Reserved is high, so why can wake still have no GPU memory
After update's temporaries are freed, PyTorch keeps the blocks as reserved-but-unallocated. Another allocator such as CuMem cannot use them.
`empty_cache()` releases free cache only, not live tensors, and not reserved memory held by active blocks or a graph's private pool. `gc.collect()` helps only when collecting a reference cycle drops a tensor's last owner; high reserved alone does not prove GC is needed.
Compare allocated/reserved/device-free and the same allocation before and after cleanup. If allocated stays high, find the live owner; if reserved-unallocated is small, empty-cache cannot explain a workspace OOM of tens of GiB. Clean up where the resource is handed back (end of update or before wake), chosen by the real call path and a retest.
### 4. How to tell a fix from a config downgrade
- Lifecycle error: fix materialization/release timing, the completion contract, or the wrong gather set.
- A legitimate peak exceeds capacity: pick a topology that carries it or shrink the declared configuration, and record the support boundary.
- Configuration not propagated: confirm the resolved value inside the container; editing the shell variable again is not a new comparison.
- Killed by Ray/container policy: account for actual headroom before judging the limit; do not raise the guard by default.
- Long duration: time by phase and optimize only large shares; GPU utilization is not evidence of completion.
Do not add a resident full-tensor dump to chase an OOM; use the phase record above to decide the next step. For the full-tensor boundary of export quantization see [precision and quantization](precision-quant.md#41-fsdp-path-verl).
### Evidence boundary of historical experiments
- A configured path is not an executed path: record a marker/count and treat zero hits as an invalid verification.
- Before comparing tensors, align request, token position, rank/owner and valid region; equal byte hashes alone do not prove the same objects.
- Verify a root cause with a single-variable intervention and a negative control that actually triggers the failure.
- Keep failed, invalid and superseded records with reasons. A small model locates the mechanism but does not extrapolate to full-model capacity or performance; the profiler explains structure, latency comes from independent timing.
- Bind every result to the actual patch stack and imported-file hashes; overlaying only new files keeps the old overlay, so restore and apply by manifest ([code and delivery](../practices/code-and-delivery.md)).
## Pitfalls
### 13. Stale quantization scale on the first step after restore
- **Symptom:** After DS4 restores from a non-zero step, the first rollout mismatches and the next step matches again by itself.
- **Cause and Fix:** See [5](#5-derived-state-after-restore).
- **How to verify:** On the same checkpoint, compare the first export's scale fingerprint and the old-logprob before and after the fix.
### 14. resume_mode=auto starts from step 0
- **Symptom:** A job with `resume_mode=auto` retrained from step 0; the log said `latest_checkpointed_iteration.txt` was not found.
- **Cause:** The checkpoint is on node-local NVMe and the newly allocated nodes have no tracker / shards. Same class: shards on the NVMe of the original eight nodes with the global `data.pt` on only one, so the launcher must check the exact node set.
- **Fix and verification:** Check before launch per item 1 of [6](#6-how-to-prove-resume-is-correct).
.agents/skills/train-inference-consistency/guides/debugging-mismatch.md
-321
View File
@@ -1,321 +0,0 @@
# Fault localization and mismatch bisect
Use this when training and rollout do not match, or a stage of a multi-node run died. It finds the first divergence with the fewest observations and produces a reviewable root cause, without dumping the whole model and without changing precision first. Performance attribution is in [Performance](performance.md); OOM and staging peaks are in [Checkpoint, resume and memory budget](checkpoint-resume.md).
## Symptom → first step
| Symptom | Do first |
| --- | --- |
| Not bitwise identical already at θ0 | [01 Isolate the side first](#isolate-the-side-first): four control groups; whichever two are equal rule out that side |
| Only some sequences are wrong | 01 same-prefix sibling: two samples of the same prompt must be bitwise identical over the shared prefix, otherwise it is the engine side |
| Wrong starting from some token position | [02 Bisect](#bisect-and-first-divergence): single-request prefill, bisect along one axis of layer/module/position |
| The difference grows slowly with training (θ0 is basically right) | [01 Isolate the side first](#isolate-the-side-first): save a small batch of prompt/response, score repeatedly at the same version to get each side's noise floor, then compare early and late adapters |
| Same operator input, still different | [Numerical contracts](../workflow/README.md), [BI kernel appendix table](bi-kernels.md) |
| Multi-node will not start / NCCL errors | [05 Communication fault ladder](#gb200-first-tell-apart-topology-implementation-mode-and-build) |
| Not sure what to save to be able to reproduce | [06 Save and compare protocol](#save-and-compare-protocol) |
## Key rules
1. Find the earliest divergence, not the largest layer; all later differences propagate from earlier ones.
2. Isolate the side before bisecting; classify before changing precision: input mismatch → weight mismatch → discrete path mismatch → kernel numerics → precision contract. Precision comes last.
3. Align probes by semantics (LSE in ln or log2, before or after norm, before or after the residual add). Misaligned probes produce false divergences.
4. Stop bisecting at a discrete decision (top-k, index); fix the routing, then continue.
5. A single round does not count; hidden coupling (one switch also changed scheduling) cannot be attributed to the switch; the target path must actually be exercised, otherwise the result is VACUOUS.
6. A two-node test must prove cross-node transfer happened; successful initialization does not mean bandwidth is reasonable; one `max_abs=0` does not prove consistency for arbitrary batches.
7. Looks like precision but is something else: determinism only takes effect on the main thread (switch scope).
## Isolate the side first
Applies when the logprob difference between FSDP/LoRA training and vLLM rollout is abnormal or grows with training.
Let μ be the selected-token logprob saved by rollout and π the real trainer old-logprob. `μ != π` does not by itself show vLLM is wrong. The goal is not unconditional bit-exactness: rule out correctness errors, measure the noise floor, and explain growth beyond it. For quantization mechanisms see the [precision and quantization guide](precision-quant.md).
### Per-sequence observation patterns
Record valid/differing token counts, first divergence position, longest run of differences, prompt/response length, UID/group, and packing/rank mapping. Distinguish: a few sequences differ from the first response token; all drift slightly; differences only past a chunk/page; differences only after an update. The pattern chooses experiments; it does not prove root cause.
Confirm token shift, mask and EOS. ±1/±2 offset comparisons are fine for diagnosis, but a fix cannot be declared through a misaligned comparison. Check that no score from another sample or version was referenced.
### Existing rescoring tool
A verified prefill rescoring tool takes the full sample tokens as input with zero new tokens and the current adapter, returns scores from `prompt_len-1`, checks the response token IDs, rejects missing scores, and clears the prefix/KV cache before scoring.
### The cheapest control: do not dump all layers yet
Save a small batch of real `prompt_ids + response_ids`, the prompt length and the valid response mask. From then on only score: do not regenerate, and do not rewrite the sequence through decode/re-tokenize. Record the base revision, adapter digest/version, config, code SHA, hardware and kernel backend. Compare training and serving under the same weights, with parameters unchanged during scoring.
Score the same version repeatedly on each side for its noise floor, then compare early and late adapters. Rising ordinary mismatch (samples differ every round) is a signal, not proof that same-input error grows with training.
| Path | Method | What it can answer |
| --- | --- | --- |
| Prefill | vLLM full input token IDs + `prompt_logprobs`, slice out the response | Difference between training teacher-forcing and the serving full-sequence forward |
| Decode trace | Normal prompt, force fixed response IDs, keep the real logits/logprob | Additional difference from KV/recurrent state, step-by-step generation, and decode kernels |
`prompt_logprobs=0` means no extra top-k; target-token scores are still returned. If zero generation is not allowed, generate 1 token, discard it, and read only input scores. The first input token has no conditional score; the first response token is predicted by the logits at the last prompt position, so do not shift by one. Check trace switch, runner support and MTP limits against the actual vLLM version. Do not quietly turn off production MTP for the probe; an explicit no-MTP control is allowed but does not cover the MTP path. Prefill consistency does not imply decode or MTP consistency.
### Metric measurement basis
- Prefer raw logprob. When comparing the sampling distribution, unify temperature, penalty, top-k/top-p and similar processing. temperature=0 in the request does not prove which kind of logprob is returned; check the backend semantics.
- `logprob_mae` (token MAE): `sum(mask * abs(lp_train-lp_serve)) / sum(mask)` over valid tokens; also report p95/max and position.
- `probability_mae`: `mean(abs(exp(lp_train)-exp(lp_serve)))` over valid tokens. verl's `training/rollout_probs_diff_mean` is this, not logprob MAE.
- Per-token logprob difference in VERL: set `actor_rollout_ref.rollout.calculate_log_probs=true`; after the trainer computes `old_log_probs`, call `verl.utils.debug.metrics.calculate_log_prob_diff(old_log_probs, rollout_log_probs, response_mask.bool())`. It returns the 1-D absolute difference over valid tokens, not a mean; compute mean and quantiles yourself and keep mask indices for localization. The default debug metric applies `exp` first, so it yields a probability difference. Checked against VERL main at commit 8e4a572 (`verl/utils/debug/metrics.py`); on another version or fork, check return shape, mask and `exp` first, and record the VERL revision.
- Relative L2 and cosine are for tensors; do not mix them with logprob MAE/KL. Flag zero vectors separately. Report full-vocab log-softmax, probability total variation and argmax separately from the token metrics. An overall shift of raw logits is not a change in the probability distribution. There is no cross-model universal threshold.
- Schulman k3: with `d=lp_train-lp_rollout`, take statistics of `exp(d)-1-d`. It is a KL estimate only when samples come from the current rollout distribution; on a fixed old trajectory it is a diagnostic statistic, not the current policy KL. Check for overflow; do not silently clip and then claim the difference is small.
### Build four control groups
1. Inference alone: frozen tokens, repeated scoring, batch1/mixed batch, decode/full-prefill.
2. Training alone: A/A and A/C packing replay under the same weights.
3. Cross-stack: μ/π for tokens with the same semantics, plus full logits where needed.
4. Score layer: run the scorers separately on exactly the same logits; check dtype, temperature, logsumexp and the reduction over vocabulary shards.
If training A/C differ and inference alone is stable, look at the training layout first, but keep the inference verification scope. If both sides are stable alone but differ across stacks, check weights, arithmetic semantics and the scoring path first.
### Same-prefix sibling
Same policy, same prompt, same response prefix, same score processing rules: conditional probabilities must be equal. Equal sibling μ with different π is evidence that the training side is sensitive to layout/state. For processed logprobs the penalty history and processing config must also match.
Compare only the shared prefix, never past the generation fork. Use a clean sibling and repeated execution as negative controls. "One side looks more stable" is not proof of correctness over the full scope.
### Localizing CUDA async errors
An async illegal access may be reported after the faulty kernel. Save the first fault, every rank's logs, input/version and execution stage; localize with scoped synchronization, CUDA core dump/cuda-gdb, or a sanitizer. Once CUDA state is corrupted, end that experiment and collect no more numerical evidence from it. Disabling async globally or switching attention backend does not prove the stability problem solved.
### Validity of controls
Disabling the cache may execute a prefill that never ran before. Disabling Graph also changes speed and scheduling; to blame Graph itself, fix each step's composition offline and compare. A negative control must verify the target branch's execution count: zero difference on an unexercised path counts as not covered. For experiment design see [04 Ablation method](performance.md#ablation-method-make-ab-really-differ-in-only-one-variable).
After isolating the side, bisect along one axis ([02](#bisect-and-first-divergence)); for classification and handling order see [03](#classification-and-handling-order).
## Bisect and first divergence
Separate input/weight errors, operator implementation differences and quantization error with the smallest experiment, without a model-specific debug framework. Metrics: [Metric measurement basis](#metric-measurement-basis); classification: [03](#classification-and-handling-order).
### Choose the search axis first
- Version regression: good/bad revision on the same reproducer, compatible environment, fixed input.
- Config regression: paired single-factor intervention; a 2×2 matrix when factors couple.
- Data trigger: request subset, length, neighbors, boundary remainder.
- Compute path: layer → attention/MoE/norm → internal operator → metadata/launch/instruction order.
Do not change image, weights, kernel and parallel layout at once. For a dirty good/bad, save the full patch and run files; the same git SHA does not mean the same program. Build failure, dependency incompatibility and failure to reproduce are not a bad numerical revision; mark them skip/inconclusive.
### Single-request prefill: narrow step by step, do not dump the whole model directly
1. **Read existing training logs first.** Check the metric definition, policy version and sample source. Samples from different rounds only hint at an anomaly; they cannot prove same-input mismatch grows with training.
2. **Fixed-input scoring.** Save token IDs and position/mask from a single render. Both ends use same-source weights and the same-version adapter, and each scores repeatedly to measure its own noise. Rule out weight, input and score-semantics mismatches first ([01](#isolate-the-side-first)).
3. **Minimal forward.** Single request, one prefill, one output token. Keep the minimal prefix or layer set covering the divergent path and its upstream dependencies, with identical config on both ends and no training update. Confirm no cache hit skipped the prefill, and separate extra calls from warmup/draft/decode. An eager probe does not stand for graph/MTP acceptance.
4. **Find the earliest divergence.** Look first at the embedding, block outputs, and logits at the last valid prompt position. Align layout, valid tokens, shards, padding and deferred residuals by semantics, not by module names or tuple indices. Expand only the earliest block that exceeds its own noise.
5. **Same-input operator replay.** Give both implementations the same logical input and weights to separate this operator's difference from upstream propagation. If needed, replace only one operator's output and compare the final distribution; recurrent state must also satisfy the contract.
6. **Return to the original failing path.** After the fix, re-verify with the real adapter, input and run config. Passing with reduced layers/eager does not replace verifying the full model, decode, or CUDA Graph/MTP.
### Per-module dump: find the earliest divergence, not the largest layer
1. Build a semantic mapping table: layer, function, input/output role, module path on both sides, logical shape. FSDP wrapper names and fused vs separate Q/K/V need not match; combine before comparing.
2. Capture block boundaries and final logits first, on a few tokens/samples; refine the earliest significantly divergent block into norm, projection, router, expert, residual etc.
3. Align valid tokens, position, mask, TP/FSDP shards, expert ID, dimension order and padding. Never compare a physical shard with the full tensor.
4. A hook takes a snapshot, not an overwritable reference, and never changes the computation. For FSDP, sample at boundaries already materialized in this forward; do not all-gather the whole model or materialize a huge frozen PLE on GPU to observe. Export bounded data per layer/needed shard and release GPU staging promptly.
5. Record forward/recompute counts so backward recomputation does not overwrite the normal forward. Graph replay does not re-execute Python hooks, so an eager dump speaks only for eager. For graph-only problems use in-graph observation or snapshots of existing buffers; do not clear the graph or recapture to cover the problem.
6. If the operator input already differs, the operator is not yet proven buggy; give both ends the same logical input and check whether it amplifies the difference. When a fused kernel's inside is not visible, mark the observation boundary instead of guessing.
Minimal artifacts: manifest (run conditions) + bounded tensor dump + comparison table with module/role, shape/dtype, fraction of finite values, max abs, relative L2, cosine, first anomalous position. Save a block's detailed intermediate tensors only after the summary has localized it; do not write all tokens of the whole model by default. Artifact format: [06](#save-and-compare-protocol).
### Align probes by semantics
Build a cross-stack boundary table: embedding, residual stream, norm input/output, Q/K/V, compressed rows/index, attention output, router logits/indices, expert input/output, combine, lm_head and scorer. Record tensor meaning, dtype, global token/shard coordinates and valid mask, not just module-name pairs. EP dispatch may reorder, duplicate or pad tokens; align expert tensors by semantic token/expert/shard key, not local offset. After expert GEMMs match, still compare separately the communication reduction, the shared + routed expert combination, and its rounding after compilation.
If the saved rollout's first response token already differs, recompute it once from the original prompt alone. This rules out later-trajectory divergence as the only cause, but attributing it to prefill, attention or a kernel still needs controls on the same policy, position/mask/scorer and per-layer inputs.
Compare original-dtype digests per layer or at a few checkpoints first; then dump only the neighboring boundaries and target samples. Keep one clean sibling. Snapshot async buffers before reuse without changing the production stream order in a way that hides the problem.
If the difference does not move under any knob, suspect a too-coarse boundary first. In an external TPU MiMo case, attention core and o_proj were captured together and the difference stayed at `7.8125e-03` across CP 8/2, segment alignment 128/2048 and FP32 accumulation on/off; split apart, the core was 0 and the whole difference was in o_proj. Give adjacent or fused GEMMs separate boundaries by default.
### Minimal but credible observations
- Copy tensors before buffer reuse; save only necessary boundaries and tokens, with a byte cap. Model-specific mappings and artifacts stay in the model working directory.
- The manifest records code/image version, checkpoint and adapter identity, input digest, position/mask, dtype/quantization, score semantics, graph/MTP, rank/layout and intervention options. Same config does not mean same weights.
- One directory per call; write the manifest after completion. Refuse comparison on missing boundaries, repeated calls, or incompatible version/shape; do not silently take the intersection or truncate. Non-finite values fail explicitly; metrics on zero norm are marked undefined.
- Compare max abs, MAE, relative L2, cosine; for final full-vocab logits also log-softmax, probability total variation and argmax.
### Recommended logs: identity + metric measurement basis + evidence boundary
Reuse existing logs and JSON artifacts; no extra logging framework. One manifest per run, one detailed report per comparison. The terminal prints summaries, not full tensors. Mark unknown fields unknown; do not backfill them as verified.
```text
[mismatch] run=<id> step=<n> policy=<digest> input=<sha> phase=<prefill|decode> rank=<n> repeat=<n>
scope=<fixed-input|independent-rollout> intervention=<none|name>
metric=<name> value=<number> reference=<run/repeat/boundary>
self_repeat=<noise> report=<relative-path> classification=<defect|numerical|quantization|unknown>
```
`metric` uses the names in [Metric measurement basis](#metric-measurement-basis) (`logprob_mae`, `probability_mae`, KL diagnostic statistic, tensor relative L2, vocab TV), never interchangeably.
Time load, rollout, score and update separately so cold start is not counted as training. A distributed job's outer console log is only a forwarded copy: when updates stop, check TaskRunner/worker raw logs, checkpoint times and completed steps first. One forward/backward stack does not prove a hang; do not restart a healthy job because log forwarding stopped.
### Limits of bisect
A floating-point difference may vanish in later rounding, masking or routing, so "this layer differs" is not monotonic in depth and O(log L) does not always find the first difference. Coarse-scan for a suspect interval, then bisect/fine-scan, verify the boundaries before it, and confirm responsibility with the same operator input.
A cut point must carry the complete state: residual, MHC mix, SSM/conv state, KV, positions, routing and necessary metadata. A fused upper-layer post and this layer's pre have no natural independent cut point. Prefix replacement that drops hidden state creates new differences.
Only a same-input operator A/B separates upstream propagation from this operator's error. Injecting reference activations, fixing the routing or replacing the kernel can confirm causality, but they are diagnostic interventions; accept the production fix only after restoring the original path.
### Form a reviewable root cause
The result must contain: the minimal input, the last identical/earliest different boundary, the same-input control, the source mechanism, the trigger condition, and the minimal change that makes the failure disappear. A config difference is only a lead: in Nemotron, running both reduction block sizes with the same kernel gave consistent results for both, which overturned the guess that the block size difference was the cause.
If after three rounds only the final score is known to differ, add observation points or re-check the sample mapping instead of rerunning in full. Keep the failing samples and the ruled-out hypotheses so the next round does not repeat them.
### Conclusion and stop conditions
The report answers only: what was observed, what the controls proved, what is not yet proven, and one minimal next step. Correctness errors must be fixed. Normal numerical differences are judged by noise and task impact. Quantization error needs same-input evidence; low-precision inference alone is not grounds to blame it. When improvement is needed, evaluate in order: supported high-precision rollout, QAT matching serving, real low-precision training. Do not cover wrong weights or adapters by changing precision.
**Replacing with the reference output is diagnosis, not a fix. Lower local error does not mean lower training mismatch; per-operator ablation gains do not add up.** Stop once this round's evidence goal is reached; do not change the training config or start new experiments along the way.
## Classification and handling order
After isolating the side ([01](#isolate-the-side-first)) and bisecting ([02](#bisect-and-first-divergence)), classify the difference, then decide what to touch. Precision comes last. Model-specific probes, runtime patches and measurements stay in the model working directory.
### Classification table
| Class | Evidence | Handling |
| --- | --- | --- |
| Correctness error | Missing weights, wrong mask/position/layout, wrong scale/block, wrong quantization range, operator semantics violate the contract | Must be fixed and the original failing path re-verified; for a diagnostic task, deliver the reproducer and evidence first, without widening modify permission |
| Normal floating-point difference | Explained by same-version repeatability/reduction/kernel noise, no contract error found | Report the noise floor and task impact; bit-exact is not the only pass condition |
| Error from correct quantization | Mapping and quantization rules are correct, and a same-input control localizes the divergence to the quantization path | Precision ablation in the order below |
| Unlocalized | Insufficient evidence or confounded variables | Keep unknown, continue with minimal probes, do not declare "it is just quantization" |
Precision ablation order: supported and affordable high-precision rollout → QAT matching the serving rules → real low-precision training (most backend requirements, last). A natively quantized model does not support BF16 just by changing the dtype. Fake quant must match the serving weight/activation rules but does not reproduce its GEMM. QAT is not a workaround for a wrong scale. QAT contract and route selection: [precision guide](precision-quant.md#1-choose-the-precision-route-first).
Keep a single checkpoint source; open a separate high-precision adapter branch when needed.
### Handling order
Input mismatch → weight mismatch → discrete path mismatch (top-k, index) → kernel numerics → precision contract. Rule out each item before moving to the next.
Acceptance: after the fix, the same failing probe passes, the original path is not broken, and error growth on fixed samples is explained/controlled; then look at training process metrics and the fixed task evaluation. Set tolerances from the model and the baseline; do not copy another model's 0.002. Statistics on a fixed old trajectory are diagnostics, not a KL estimate of the current policy.
This procedure does not authorize stopping jobs, switching dtype, adding GPUs or changing the production recipe. When the next diagnostic gate needs new permissions/resources, report the smallest next probe and its cost.
### Delivery
Follow [06 Save and compare protocol](#save-and-compare-protocol): manifest, bounded dump, comparison table, semantic boundary of the earliest divergence, classification and confidence, one next step; state which paths were actually tested. No separate framework, no cross-model universal threshold.
A case that looks like precision but is a switch-scope correctness error: [determinism takes effect only on the main thread](bi-kernels.md#determinism-takes-effect-only-on-the-main-thread).
## GB200: first tell apart topology, implementation, mode and build
Based on experiments on GB200 four-GPU nodes; future DeepEP/NVSHMEM releases may change the support matrix. Check the source of the target fork/commit and the installed artifacts before relying on it.
### Do not use one name for four things
| What to ask | Example record |
| --- | --- |
| Which implementation | AllToAll, some version of DeepEP, HybridEP; note the fork/revision |
| Which path of that implementation | The actual class, API and kernel behind high-throughput / low-latency / v2 |
| How it communicates | NVLink within a node; RDMA across nodes; or the MNNVL path actually enabled |
| What it was compiled into | 4-peer/8-peer macro, target GPU, NVSHMEM dynamic dependency |
`NVL4 + RDMA fallback` describes the communication topology. **It does not state the DeepEP mode**: do not answer low-latency or high-throughput unless the call site proves it. How `DeepEP v2`, `HybridEP` and HT/LL relate depends on the version; do not infer containment from the names.
MNNVL being off does not rule out cross-node EP: the verified path may be NVLink within a node and RDMA NICs across nodes. Nor does that path work on every GB200; devices, network, peer build, underlying libraries and the called API all need verification.
### Get the world size right first
Write the trainer and rollout process groups separately. Formulas for dense world and expert world, and grouping pitfalls such as "two DP4/EP4 are not one EP8 group", are in [training-side layout and boundaries](training-forward.md). The runtime decides the actual grouping; do not change the GPU count to fit a config string.
### Same-fabric fairness
Before any communication-bound A/B, prove both arms use the same transport. In one GB200 NVL case the BI=0 arm's NCCL AllGather/ReduceScatter went `via P2P/MNNVL` (same clique, cliqueSize 8) while the BI=1 arm's DeepEP low-latency a2a was pinned to RoCE (`VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL=0`, `NVSHMEM_HCA_PREFIX=rocep`); a gap blamed on non-overlappable communication latency was mostly the wrong NIC, fixed by two environment variables.
Check steps:
1. `nvidia-smi -q | grep -A4 Fabric`: the ClusterUUID/CliqueId of the two nodes must match.
2. Run `NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,NET` once per arm and check whether the transport is P2P/MNNVL or NET/IB.
3. grep the arm environment files for DeepEP/NVSHMEM pins. A hard `export` overrides the value passed by the container `-e`; write it as `${VAR:-default}`.
Mechanism: `VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL=1` → `deep_ep.Buffer(allow_mnnvl=True)`, DeepEP no longer exports `NVSHMEM_DISABLE_MNNVL=1`, and NVSHMEM 3.4.5 uses MNNVL for peers in the same clique automatically. `DEEPEP_MAX_NVL_PEERS` is not a runtime switch (it is DeepEP's compile-time `NUM_MAX_NVL_PEERS`); the peer-4/peer-8 build is selected by switching the overlay. MNNVL is not hybrid-only: the DeepEP v1 LL kernel branches per destination on `nvshmemi_get_p2p_ptr` (direct NVLink write vs IBGDA), and MNNVL only makes cross-tray peers P2P-mappable; "hybrid mode" is a v2/elastic concept. MNNVL was on by default before DeepEP #190 (2025-06-06) and turned off afterwards to save NVSHMEM memory.
## Save and compare protocol
Large tensors and checkpoints stay in `results/<run-id>/`; source and small summaries go into version control. Each framework dumps differently, so write these per project following the key points below: logprob comparison, per-layer trace, replay controls, a fixed-token oracle (logprob/argmax of a fixed token set before and after a weight update), GSM8K record comparison before and after an update, a `.meta.json` provenance sidecar for the results directory, and a pre-experiment snapshot commit + clean worktree of uncommitted changes with a check of the container's actual import path and file hashes.
The manifest records artifact path, size, hash, generating command, code/image/policy version, number of valid samples and hook execution count. Confirm the write succeeded before cleaning up compute processes. Do not auto-delete existing artifacts.
### Archive experiment material by step
Build the step index before the experiment and fill in status and evidence after each run, not at the end. The layout below is logical; link existing material by explicit path instead of moving or copying it:
```text
<workspace>/
steps/
INDEX.md # current step, evidence entry of each step, gates not passed
01-baseline/
INDEX.md # tuning comparison table, candidates and review records
<run-id>/
config/ # resolved config, environment variables, version list
scripts/ # actual server/bench commands or pinned-version entry points
results/ # raw results/logs, or the list of persisted artifacts
summary.md # conclusion, limits, failure cause and next experiment
05-trainer-forward/ # classified by C0-C10 gate number, created as needed
10-final-acceptance/ # references the baseline ID and this candidate version
results/<run-id>/ # large tensors/checkpoints, not in Git
```
- **Every run is traceable.** Save the actual launch command, full effective config, code/image/checkpoint identity, nodes and run time, evidence of the backend that actually ran (a CLI flag is not evidence), exit status and raw results path. A branch name is not provenance: record branch, snapshot commit, base SHA, patch, tree hash and image digest in a `.meta.json` sidecar. Replace secrets with references. Large files go to persistent storage with size and hash, never only a node-local path.
- **Baseline tuning: one row per run**: run ID, candidate config (MoE/communication), actual backend, load and concurrency, repeat rounds, valid input/output tokens, cache hits, throughput/latency, exit status, evidence link; N/A where inapplicable. Keep failed, unfinished, diagnostic and excluded runs with the reason. The summary lists the included run IDs; do not pick the best round or keep only the final mean.
- **A candidate is not a reviewed baseline.** The index marks the recommended candidate and the human-confirmed baseline ID separately; the baseline links to exact config/image/code, result set and review record, or reads "pending review". Later comparisons reference that ID; do not swap the denominator for a same-named script's new content or another backend's old result.
- **Reruns keep history.** A new run gets a new ID and overwrites nothing. Mark an invalidated conclusion with the reason and the replacing run. Continuing an experiment keeps separate attempt/segment and exit records. Reuse evidence across steps by reference, not by copy.
- **On handoff, read the index first.** The step owner updates its record. Reporting, recovery and acceptance check raw evidence through the index; a missing file, mismatched identity or chat-only conclusion is an evidence gap, not a pass. Index historical material; do not delete or bulk-move it.
### Replayable raw data
Keep the real framework batch: prompt/response IDs, attention/response masks, positions, lengths, UID/group/sibling, rollout logprobs, necessary reward and routing records, plus the packing mapping, parallel config and how to restore that step's rollout weights. The raw batch is the replay input; exported per-token logprobs are for comparison only.
Step N's rollout usually comes from the policy before update N, so the post-update checkpoint may not reproduce it. Under partial rollout or async staleness, segments of one sequence may come from different policies (re-prefill after interruption continues under the new policy); record the policy version per segment, or a single-checkpoint replay turns expected policy difference into false inconsistency. Freeze weights, mutable buffers and RNG/dropout state, and state whether the scorer uses eval or training mode.
### Compare per-token logprobs
Each side (rollout and training old-logprob, or two replays) exports `{request_id: (logprobs, response_mask)}` as original-dtype tensors with `torch.save`. Do not convert to JSON or rebuild from printed logs. Key points:
- **Align by logical request, not physical row.** Remove padding and restore the pre-packing UID order. Both request sets must match; a missing request raises an error rather than comparing the intersection.
- **Compare only valid response tokens.** Positions outside the mask are excluded. Every sequence needs at least one valid token; an empty batch is an invalid experiment.
- **Bitwise comparison looks at bytes, not values.** Value comparison treats `+0.0` and `-0.0` as equal and two NaNs as unequal. Compare integer views: `a.view(torch.int32)` (float32) or `a.view(torch.int16)` (bf16). Do not cast first and then claim bitwise identity in the original dtype.
- **Same contract on both sides.** Same policy version, same raw/processed basis (for processed, record temperature, truncation, penalty), same input. Compare input token hashes before scores; differences from different inputs are not numerical differences.
- **Report** valid token count, byte mismatch count, first mismatching (request, position), and max/mean absolute error. Matching sampled-token logprobs do not mean matching full logits.
```python
def compare_logprobs(a, b):
assert a.keys() == b.keys(), "request sets differ"
n = bad = 0; first = None; max_err = 0.0
for uid in sorted(a):
(x, mx), (y, my) = a[uid], b[uid]
assert torch.equal(mx, my), f"{uid}: masks differ"
x, y = x[mx], y[my]
ix = x.view(torch.int16 if x.dtype == torch.bfloat16 else torch.int32)
iy = y.view(torch.int16 if y.dtype == torch.bfloat16 else torch.int32)
diff = (ix != iy).nonzero()
if len(diff) and first is None:
first = (uid, int(diff[0]))
n += len(x); bad += len(diff)
max_err = max(max_err, (x.double() - y.double()).abs().max().item())
return {"tokens": n, "mismatched": bad, "first": first, "max_abs": max_err}
```
### Per-layer trace: find the boundary of the first divergence
When logprobs mismatch, record tensors with forward hooks at the same semantic boundaries on both sides (for example each layer's attention output, MoE routing, residual), find the first differing boundary in execution order, then do a bounded dump there. Key points:
- **Key by semantics, not rank.** Each record's key is `(request_id, boundary name, logical slice)`, for example `("step7-p2-s3", "layer2.attn.out", "tokens=0:128")`. Stitch TP/CP/EP shards back by global coordinates and remove padding before recording. Naming and order must match on both sides. A missing probe invalidates the whole experiment; it cannot be skipped.
- **Hashes first, tensors on demand.** Take SHA256 of the contiguous tensor's original-dtype bytes (`hashlib.sha256(t.contiguous().view(torch.uint8).cpu().numpy()).hexdigest()`), with shape, dtype and semantic execution sequence number. A hash shows where it differs, not by how much; save raw tensors only at the first differing boundary.
- **Clone before the buffer is reused.** Inference-side workspace and CUDA Graph output buffers are overwritten by the next call; `clone()` in the hook before hashing or saving.
- **The first difference is not the root cause.** To blame an operator, also prove its input is identical on both sides (record its input hash too).
### Replay controls
With fixed rollout and fixed weights, change one variable (packing order, batch composition, parallel layout, kernel switch), replay old-logprob, and use the comparison above to judge whether that variable affects the result. Key points:
- **One process per variant**, with a synchronous entry point that blocks until the replay finishes. Do not use background submission such as `sbatch`: on timeout or failure the job cannot be reclaimed. Variants that change compilation, libraries or initialization environment need a new process.
- **Each variant writes to a new directory** without overwriting old results; record the actual command, environment, exit code and elapsed time. Preview the command and budget before executing; on failure, stop at that variant.
- **Prove the replay is frozen**: record the weight fingerprint, confirm no generation and no optimizer update, and confirm the packing actually in effect, not just the command-line arguments.
- For the experiment record use [the report template](../workflow/report-template.md).
.agents/skills/train-inference-consistency/guides/environment-provenance.md
-106
View File
@@ -1,106 +0,0 @@
# Environment, image and runtime version provenance
Once the environment is set up (the image pins the native runtime, uv manages training dependencies, source arrives as a traceable snapshot or a minimal overlay), prove that it is what is actually running. Code existing does not mean it is executing.
## Find by symptom
| Symptom / question | Where to go |
| --- | --- |
| Not sure which version of source / wheel / .so / image is actually running | [Run manifest](#use-one-run-manifest-to-answer-what-exactly-is-it-right-now) |
| How many commits "latest" actually differs from the target branch | ["Latest" must come with a target branch](#latest-must-come-with-a-target-branch) |
| Overlay, or change the source and rebuild | [Boundary between overlay and native wheel](#boundary-between-overlay-and-native-wheel) |
| Recompiles on every start / not sure where the cache should go | [Compile cache](#compile-cache-where-it-lives-decides-whether-it-can-load) |
| Not sure which path DeepEP actually runs / how it was built | [Four things about the communication implementation](#communication-implementation-do-not-use-one-name-for-four-things) |
## Key rules
1. Pin the image digest, CPU architecture and Python version first, then resolve training dependencies. Image tag, PyPI version and upstream PR merge status do not replace one minimal inference run in the image.
2. `--system-site-packages` does not stop uv from installing a same-named package that then imports first; check the install plan, the actual import path and the GPU operators together.
3. Each of the six manifest layers (source, installed packages, native libraries, JIT kernels, container, run) has evidence nothing else replaces: `latest` is not the digest, `pip show` in another venv is not the actual import path, the Python source HEAD is not the loaded .so.
4. Audit inside the worker, interpreter and mount namespace that actually run; the login container shell and the Ray worker may differ.
5. One image tag may point to different digests on different dates/nodes; ARM64 and AMD64 tags may be independent builds.
6. Modify in the host worktree, execute in the container: same-commit worktree plus per-file hash check, never an overlay stack detached from the formal source.
7. Write version pins such as "the Megatron fork is only compatible with Megatron-Bridge v0.6.2" into the run manifest.
8. How the image was produced (container commit exported to sqsh, or Dockerfile rebuild), where it was pushed, and the three-repo SHAs in the tag are part of the run identity.
## Environment identity: code existing does not mean it is executing
### Use one run manifest to answer "what exactly is it right now"
Record recipe, trainer, Megatron/MCore, vLLM, DeepEP, DeepGEMM, FlashAttention/FlashInfer and NVSHMEM separately, each with at least:
| Layer | Evidence to save | What cannot substitute for it |
| --- | --- | --- |
| Source | remote, target branch, HEAD, dirty patch hash, submodule gitlink/actual HEAD | `latest`, the repo directory name |
| Installed package | Actual interpreter, import path, package version, wheel/build revision | `pip show` of another venv |
| Native library | Path/hash of the `.so` actually loaded, build inputs, target architecture, dynamic dependencies | Python source HEAD, extension file name |
| JIT kernel | DSL/source hash, compiler version, specialization/config, cache origin | `.so` unchanged |
| Container | OCI manifest/platform digest or `.sqsh` SHA256, parent image, build record | A mutable Docker tag, image creation time |
| Run | Resolved config, worker environment, model/weight digest, the real kernel | An export in the head node shell |
Write `unknown` for unknown items; do not infer the whole source tree from a wheel version string. Audit inside **the worker, interpreter and mount namespace that actually run**. If dynamic linking is suspect, check the process's actual mappings/loader result, not the static prediction of `ldd`.
ARM64 and AMD64 tags may be independent builds or members of one multi-arch manifest. Compare the target platform's manifest/config/layers; a similar tag or update time does not show it is the same image. Record generic scheduling switches such as `CUDA_DEVICE_MAX_CONNECTIONS=1` per the task convention; do not add them quietly and then credit BI.
A cross-repo compatibility pin is a manifest row too: state the pin (the Megatron fork works only with Megatron-Bridge v0.6.2) and its basis, and re-verify when either side is upgraded. Do not assemble incompatible releases just so every item reads "latest".
#### Communication implementation: do not use one name for four things
See [debugging mismatch](debugging-mismatch.md#do-not-use-one-name-for-four-things).
### "Latest" must come with a target branch
Pin each repo's target branch per the [delivery target](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image). The reference implementation's PR is a design reference, not automatically the PR base. If the user names another fork/branch, record the override. fetch is not checkout; submodule update follows the parent repo's gitlink, not each submodule's own remote.
```bash
git -C /explicit/repo status --short --branch
git -C /explicit/repo remote -v
git -C /explicit/repo fetch origin
git -C /explicit/repo rev-list --left-right --count A...B
git -C /explicit/repo log --left-right --oneline A...B
git -C /explicit/repo merge-base A B
git -C /explicit/repo diff --stat A B
git -C /explicit/repo submodule status --recursive
```
The left/right counts of `A...B` say how many commits each side has; the merge base and diff say what differs. Diverged branches have no single before/after relation. Commit time and short-hash magnitude do not decide ancestry. If a shallow clone lacks the ancestor, fetch history before claiming the lines are unrelated.
Keep the recipe's declared compatible combination by default. If the user explicitly asks to update a development branch, record the deviation from the pin and check the actual changes and the ABI.
### Boundary between overlay and native wheel
During development, fix the container code-sync and build/test method first and record the tool versions in the run manifest. Modify in the host worktree, execute in the container; do not maintain an overlay stack detached from the formal source:
1. Confirm the config, model registration and import identity of the native image/checkpoint; pin the image digest, the full source commit and the real install directory. Do not guess a missing implementation from the model name or tag.
2. The sync worktree's `HEAD` equals the image's full commit, and the changes under test are a diff on that baseline. They may be ported from the development branch, but a different HEAD cannot bypass the check.
3. In a container not yet running the workload, sync each changed in-package Python file individually: before copying, check the original file's hash against `HEAD`; after copying, check the hash of the file under test; save the sync manifest. For new files, check no original exists on either side. Never overwrite the whole package or modify a running container.
4. Keep the image's version, generated files and unchanged extensions. Recompile native/binding changes and verify the artifacts; for JIT changes check compile inputs and cache; if the ABI/toolchain changes, rebuild the image.
5. After startup, re-check import/library paths in the actual worker and test. Finally build the unified image from the formal PR's clean commit and rerun acceptance; the user does not redo the development sync. Details: [Implementation and delivery](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image).
Overlays are for development/diagnosis only, not the delivery format.
- An overlay covers specific files in the container with external copies; a wheel is an installed package that may carry compiled machine code.
- Cover only the necessary files at their exact target paths; do not mount the whole repo over the installed package and shadow its extensions. Verify the target's original hash first, apply once at startup, stop on failure, and save the overlay manifest, final hashes and import paths.
- A Python-only change can still change the binding/JIT/shape contract; it is not automatically exempt from recompiling or testing.
- Changes to CUDA/C++, bindings, compile macros, submodule kernels or the CUDA/Torch ABI need checking and recompiling. Before acting on a "missing extension", rule out the wrong venv, shadowing by a whole-repo mount, search paths and dynamic dependencies.
- If the image already has the correct artifact, verify it and use it. Do not recompile a DeepEP already proven to work, and do not change only the source pin in the recipe expecting the old `.so` to change.
- Do not fix a FlashAttention extension load failure by `unset`ting a variable. Accept a config fix only once it is shown the variable selected an unneeded/incompatible path and the run passes with the user's target path restored.
- Keep observability overlays separate from numerical/performance changes and off by default; when cleaning up the main PR, back up still-useful diagnostic code before deleting it.
### Delivery gate for immutable images
1. Commit production changes to their PRs; freeze the clean source commit and all native/JIT build inputs and dependencies; record the base image digest, not just the tag. A dirty patch may identify a diagnostic candidate but cannot be an uncommitted dependency of the final run.
2. Build the final image and check import/dependencies/hashes in a fresh container **with no development source mounts and no inherited JIT cache**, keeping the device, model and data mounts production needs.
3. Run the agreed correctness, Graph/cache and performance gates in that image; if observability was on, accept again with the formal config and no probes.
4. Publish exactly the accepted artifact and verify the registry digest/architecture; if conversion or repackaging changes the content, verify the equivalent closure and run a final smoke.
5. Sync recipe, PR description and image receipt; "image push succeeded" does not cover items that failed. Follow this order when the user asks to push after testing.
The receipt states how the image was produced: Dockerfile rebuild from a clean commit, or changes installed in a container then committed and exported as `.sqsh` (the only option when the coordinator has no docker permission; fine for development, but formal delivery returns to the Dockerfile); which registry it went to (Nemotron: public Docker Hub); a tag carrying the vLLM/Megatron/verl SHAs, not just the date. Record both the `.sqsh` SHA256 and the registry digest under the same tag.
Adding a new project source layer to an old OCI base does not bring in another experiment image's vLLM/FA4 changes. A deleted file missing from the top-level tar may still exist in lower layers; deleting a base file needs a correct whiteout or a check in the final rootfs.
Upgrading FlashInfer/NVSHMEM must not silently upgrade Torch/CUDA. Pin the target interpreter and constraints, compare packages/dynamic dependencies before and after, and confirm the APIs used and build compatibility. A version number is not evidence of an HCA fix or of model correctness.
## Compile cache: where it lives decides whether it can load
Put the Triton/Inductor/vLLM compile caches on node-local disk mounted explicitly into the container; a `noexec` mount (such as `/dev/shm` in some containers) cannot load the compiled `.so`, and a per-container tmpfs recompiles on every start. Isolate caches by image, source commit and GPU architecture, and give each rank on a node its own copy, since an existing cache directory does not mean a cache hit.
.agents/skills/train-inference-consistency/guides/performance.md
-160
View File
@@ -1,160 +0,0 @@
# Performance and profiling
Record performance gates and correctness gates independently: low throughput does not mean the numerics are wrong, and correct numerics do not mean performance meets the target. BI kernel choices and costs are in [BI kernels](bi-kernels.md); the statistical protocol for paired acceptance is in [Pairing and computation](recipe-and-acceptance.md#pairing-and-computation).
## 1. Workload and measurement basis
### Three throughput measurement bases
- **Per-request tok/s**: 1/TPOT; decode latency.
- **Engine total tok/s**: completed tokens of all concurrent requests / common timing window. With multiple engines use the full common window and total completed tokens, not a sum of per-engine throughput; with unsynchronized node clocks, do not compute the window from timestamps.
- **RL one-step gen time (makespan)**: one batch of rollouts from the first request in to the last request out. This is what RL cares about; long tails and scheduling can move it opposite to the other two.
State which basis you report; compare ratios only on the same basis.
### Translate the rollout target into actual kernel workload
"batch=128, n=8, input 2K, output 8K" must at least specify:
- Is 128 prompts, expanded requests, or per-engine concurrency? 128 prompts with 8 siblings each is 1024 requests.
- Is 2K/8K the maximum or a fixed length? Record input token mean/max/percentiles, total output, and completed requests.
- `max_num_seqs` limits simultaneously active sequences; `max_num_batched_tokens` is the token budget per scheduling pass; neither equals the total request count.
- Requests per DP engine, prefill chunk size, and the active decode batch over time; a global 128 is not each rank's kernel M.
- Prefix cache cold, warm, or naturally hit; the n siblings of a prompt affect hits and prefill.
- Natural EOS vs ignore EOS / fixed length, sampling parameters, seed, request order, maximum model length, scheduling policy.
### replay bench: equal-workload replay of one RL step
Use the same request count and prompt×n as one RL step, take each request's output length from one real rollout, replay with ignore EOS at fixed lengths, and measure makespan. Measure the repeat noise of one config first, then compare arms; draw no conclusion from a difference smaller than the noise. With the same length table, token count, request count and long tail are identical, so the difference belongs to the execution path. Fixed output suits equal-workload comparison, not quality.
In natural generation, different numerical paths or scheduling can change token counts even with the same weights and seed; record output token receipts and add a frozen-token replay if needed. A short 2K+64 profile can locate startup/local hotspots but **cannot** represent the long-term decode, cache/page pressure, scheduling tail or full rollout speed of 2K+8K. Use different experiment IDs for the short profile, the full stress test and the training E2E.
### Metrics are cumulative values
When one engine runs multiple stages (concurrency or parameter sweep), `vllm:iteration_tokens_total_*`, `vllm:spec_decode_num_*` and `vllm:num_preemptions_total` in `/metrics` are cumulative since process start. Always compute per-stage cost as `after - before`, and note in the script that the first stage is right only because before is all zero, which makes the method look correct. Cross-check the prefill amount with an independent client-side basis (prompt_tokens/cached_tokens in the request records); if the two disagree, the counter basis is wrong.
When the verification method itself fails (a regex that never matches, a hardcoded version number, an unsubtracted cumulative value), the numbers look plausible but measure something else. Reconcile every new script's first numbers against an independent basis.
## 2. Fair A/B
### Do not mix the two kinds of A/B
- **Causal comparison of the BI switch**: same software/model/scheduling/cache, only BI changes. If the switch selects a different kernel automatically, record the actual branch and add no other optimizations.
- **Each side's best available config**: BI=0 uses the strong default / verified best path for that platform and model; BI=1 uses the fastest path that meets the agreed consistency. Backends may differ; disclose every difference.
Do not force BI=0 onto FA2, a specific MoE, non-default attention, or Graph off to shrink the gap. Do not call an experiment that locally disables switches inside BI=1 "native BI=1". A replaced path is a separately named candidate that passes the full contract gate. Give common optimizations (cutting shared launches, enabling a supported network path) to both sides and report the common gain and the reduction in BI overhead separately. BI=1 faster than BI=0 is not automatically fake, but it triggers an audit of config, workload, cache and timing.
### What to fix
Weights, input, generated tokens/lengths, output cap, concurrency, nodes, software, warmup and timing boundaries. If output receipts and workload differ, report only each side's own throughput. Separate the fresh-process first rollout from warm steady state.
**Same fabric**: before comparing communication bottlenecks, confirm both arms use the same transport; see [Same-fabric fairness](debugging-mismatch.md#same-fabric-fairness). **Same node**: small gains need same-node pairing, crossed order and multiple rounds; counterbalance across node pools and do not pick the fastest cross-pool combination. **JIT pollution**: a JIT kernel specialized by token count recompiles for every new prefill size under BI=0; verify zero compile lines during the bench. **BI=0 silently running the BI kernel**: a fork whose gate only checks whether the `.so` exists makes BI=0 run the BI kernel; grep `vllm::batch_invariant::` in nsys.
### Report speed switches and numerical switches separately
One switch (PC, cudagraph, BI) changes both the numerical path and throughput. Before reporting "probs_diff with BI on/off" or "throughput with BI on/off", list the confounders that change with it, rule out or fix each, then compare; see [Paired acceptance](recipe-and-acceptance.md#paired-acceptance).
### Computation and reporting
```text
output tokens/ms = completed output tokens / (wall_seconds * 1000)
BI throughput deficit = 1 - throughput_BI1 / throughput_BI0
BI time overhead = time_BI1 / time_BI0 - 1 # matched work only
```
Latency reduction `1-new/old` and throughput gain `old/new-1` differ; 16.35% lower throughput is not 16.35% more time. Report numerator, denominator, request / completed count, raw receipts, config differences, sample count and variation; take the median over pairs and report single pairs separately. Gains of A and B do not add; measure the combination, and do not claim it beats the best single item from a few observations. List skipping the old-logprob forward separately, only when the RL algorithm allows it; it is not a gain of the BI kernel. Every ratio must be recomputable from explicit artifacts.
## 3. Breakdown and ablation
### Two stages: remove structural defects first, then fine-tune
Once proxy inference consistency passes, check feasibility: sweep token row counts on real operator shapes, count launches and host gaps, look for per-token launches, synchronous readbacks, repeated copies and abnormal workspace, and confirm the target Graph captures/replays. After backward is wired in, measure training forward and forward+backward (by token count) on the main topologies fixed at C1 against the native training path on the same build, and check them against the C1 budget, so a scheme fast at decode but slow in training/prefill is not chosen. After full consistency, fine-tune by E2E profile -> micro benchmark (screen high-share shapes, then sweep all target shapes and boundaries; all must be correct; weight performance by real call frequency) -> multi-round paired E2E on the same workload. New or replacement kernels do not use Triton; use CuTe DSL, CUDA/CUTLASS or an existing implementation; historical Triton data is only a comparison. Final delivery still needs full acceptance, see [validation](../workflow/validation.md). Fix a structural problem first and re-check the numerics; do not turn this stage into endless tuning.
### Break down the overhead first
Time model load/JIT, rollout prefill/decode, old-logprob, backward/update and offload/wake separately. Kernel times may overlap and do not add up to wall time. Keep the nsys window, rank/GPU normalization, Graph node tracing and sampling interval consistent.
Profile the effective config and actual kernels, not just the BI switch. In DS4 the same quantization `.so` loads under both BI0 and BI1, but different MoE layouts select different schedulers. For an NVFP4 MoE model, account for dense linear, expert GEMM, expert routing/indexing, all-to-all and attention separately; one vendor name does not describe the path (one Nemotron run mixed FlashInfer FP8 linear, Humming NVFP4 dense GEMM, Humming indexed experts and FlashInfer NVLink one-sided all-to-all). A run that fails on a backend guard (a BI=0 trial expected Humming but got FlashInfer CuTe DSL) calls for re-auditing the backend and build path and does not enter the performance denominator.
### Collect enough attribution information in one profile
1. Verify the actual serving config and input/output statistics. For inference-only use vLLM serve or an existing client, not training; use the mp/Ray executor actually needed.
2. Add low-overhead, off-by-default NVTX at semantic module boundaries: attention, indexer/Top-K, index preparation, MLA, MHC, shared/routed MoE, norm, linear, sampling, communication. Do not apply the DS4 module list to dense models.
3. Use nsys first for the timeline and hotspots, with CUDA Graph node tracing on; graph launches without node details do not support per-kernel comparison.
4. Count calls, GPU time, CPU launch/sync, allocation and shape per path; aggregate by rank, phase and module.
5. Pick representative prefill, steady decode, mixed and tail windows, then ncu a few kernels/shapes for bandwidth, occupancy, memory access and compute; ncu-perturbed wall time is not throughput.
6. Measure throughput on the same full workload without profiler or heavy dumps, and record whether profile and timing used the same code and config.
A kernel name is not a module identity: copy/fill/cast, `index_select`, generic elementwise, GEMM, and different shapes of one template appear in many layers. NVTX/launch correlation and shape are needed to tell "the indexer's gather" from other gathers. Parent/child NVTX inclusive times do not add; the sum of concurrent kernel GPU durations is not wall time; label multi-GPU sum, per-GPU mean and critical-path time separately.
When a wrapper/overlay also rewrites the kernel selector, confirm who launched a kernel with **three layers of evidence**: per-worker class, object count and import source at initialization; input shape, module name and dispatch at the real forward callsite; launch correlation and GPU kernel in the final profiler. Overridden linear objects or a Humming shape do not prove what launched at replay, and overlay env vars set where the worker does not inherit the entry point falsely suggest the candidate is enabled. Graph may hide ordinary Python CPU op events, so a `cpu_op` count of 0 alone does not prove the callsite did not execute. When evidence conflicts, run a minimal probe that keeps the numerical path; do not fall back to an unverified default operator for throughput.
For a Roofline bound, use the peak compute of the actual dtype and the bandwidth of the memory level studied; L2 / shared-memory / non-Tensor-Core kernels cannot borrow the HBM / Tensor Core bound, and do not infer the metric from chart coordinates.
### Pin "where is it slow" to a process and a call stack
The frontend and control plane are in scope. Split the makespan difference by process: py-spy the engine, worker and frontend once each and see which function holds the main-loop time; use asyncio loop-lag to see whether the frontend is saturated. With R3, most of the frontend main loop is in `output_handler`; without R3 the frontend is almost idle, which moved suspicion from the GPU to the frontend object count. Optimize only parts with a large share; GPU utilization is not evidence of completion.
### Single-operator replay and synthetic micro benchmark
A synthetic shape matrix proves algorithm boundaries; real rollout replay supplies score distribution, ties, masks, non-contiguous layouts, cold/hot buffers and call frequency. Weight by sampled shapes/frequencies and keep high percentiles and rare boundaries; report decode small M and prefill large M separately. Time with and without workspace allocation; if the original allocates per call and the candidate preallocates, do not compare only the kernel body. A micro A/B fixes input and resources and checks the candidate bitwise against the baseline; change one thing at a time, covering at least the batch sizes actually used (1/4/32/128), different neighbor compositions, prefill/decode, Graph, mask/tie/boundaries; mark shapes not run as not covered. When Triton is forbidden, a "temporary benchmark" does not bypass that.
### Ablation method: make A/B really differ in only one variable
These rules apply to intermittent numerical leaks (the batch-invariance kind) and to performance attribution; the examples come from the DS4 batch-invariance investigation. Process rules are in [Experiment discipline](../practices/evidence-and-experiments.md#experiment-discipline).
1. **A single round of 0-diff does not count.** For a leak with a hit rate below one half, a single all-green round has probability above one half. Re-check a statistical conclusion with a deterministic reproducer, especially "turned clean"; 0-diff is necessary but not sufficient.
2. **For "turning off X fixes it", first check whether X's switch also changes another branch.** Example: with cudagraph mode NONE, the budget split function also took a different branch, so the eager experiment moved two variables. Hidden coupling cannot be attributed to the switch.
3. **A switch that changes speed also changes the scheduling composition.** Eager is several times slower than graphs, so batch composition changes. Before a statistical comparison, pin each step's batch composition with offline manual steps; otherwise "clean when off" may just mean it was not hit. Turning off the cache is the same: a prefill that never ran before now runs.
4. **The target path must actually be exercised.** A long-prompt filler looks "heavier", but the step exceeds `max_cudagraph_capture_size` and the graph does not run, so identical results across settings are an illusion. A negative control records the target branch's execution count; zero hits is marked VACUOUS and does not count as verification.
Before localizing, draw a 2×2 (or higher-dimensional) matrix to separate the variables, then choose which cells to run.
### Ablations must be able to identify interactions
With optimizations A and B, measure `base / A / B / A+B`; with network/Graph/cache switches, add the corresponding controls instead of running the full combination mechanically. Reuse A+B data only under identical weights, input, software and timing. The first round after restoring checkpoint N is usually step N+1, and adjacent names do not prove the same policy: fix the checkpoint content digest and export fingerprint, and do not mix BI0/BI1 or R3/no-R3 branches. For rollout-only timing, exit cleanly after completion without optimizer/save, and record that all requests completed. Top-K and index post-processing can usually be ablated orthogonally, but shared buffers, layout or Graph create interactions, so re-accept the combination.
### Segment-by-segment isolation: add the total difference back up
Cut a long path into arms at boundaries, each enabled up to one segment, and measure what each segment adds:
| Arm | Enabled up to | What it answers |
| --- | --- | --- |
| All off | No capture | Baseline |
| A0 | Capture only (hook on, nothing lands on the data plane) | Cost of the capture kernel/sync point |
| A | Capture + worker data plane + packing, not sent to the scheduler | Worker-side processing cost |
| E | worker→engine IPC on, scheduler discards | Inter-process transfer cost |
| B | No capture, return path only (fake data) | Return path cost, independent of capture |
| All on | Current path | Total difference |
Arm differences should roughly sum to all-on minus all-off; otherwise there is an interaction or a dirty cut. This localized the Nemotron-H R3 return path: A0 and A are almost level with baseline, E adds a little, B is close to all on, so about 80% of the overhead is on the per-step engine→frontend return path, not the capture hook.
Measure same-config repeat noise first and ignore arm differences below it.
### Optimization priorities and correctness gates
1. Remove idle/repeated work: launch grids mismatched with the real workload, repeated moves, invalid tokens.
2. Reduce host scheduling/sync and kernel launches; restore a usable Graph with the output contract unchanged.
3. Fuse boundaries such as gather/dequant/index post-processing, keeping the original rounding, traversal order and padding/sentinel.
4. Only then change the kernel algorithm, chunking or collectives; these may change the reduction order and need new equivalence evidence.
Example: a quantization kernel launches 1024 blocks per expert while decode has few valid rows. Shrink the grid and the stride together; under Graph capture use a host-known capacity bound, not the device token count. If dequant/gather are element-wise independent, re-chunking may keep bit-level equivalence; fusing into sparse attention must keep the scale rounding and KV traversal/reduction order.
Each candidate first runs a same-input kernel comparison (real dtype/shape, non-contiguous or masked layout, empty experts, tail, extreme values, ties, reused dirty buffers, Graph replay), then the fixed-token engine gates, then actual rollout/training. Formatting checks do not replace compiling and executing; rerun measurements on the final code, since a JIT refactor that looks behavior-preserving may fail to compile. A candidate with no gain or a regression keeps its conclusion and is withdrawn from the default path.
### Delivery
Each item lists the actual kernel/path, correctness before and after, micro latency, rollout and full-step results, independent sample count, Graph/topology scope, raw receipts and unsupported configs. State whether quality evaluation and long-term stability are done; kernel speedup does not replace an end-to-end conclusion.
## 4. Lever results table
Measured switches: directions and dead ends not to retry.
| Model / arm | Winners | Dead ends (measured, do not retry) | Structural floor |
| --- | --- | --- | --- |
| DS4-Flash GB200 DP8/EP8 BI=0 | FlashInfer trtllm MoE (`--moe-backend auto` + new FlashInfer); `cudagraph_mode=FULL_AND_PIECEWISE` | deepep_low_latency, `--api-server-count>1`, NCCL_ALGO/PROTO, max-num-batched-tokens 4k/8k, linear backend fitrt/cudnn (cutlass crashes), DBO (forces the V1 runner, which does not accept PIECEWISE) | MoE decode is weight-bound (gemm1 near HBM roofline); DBO rereads weights and loses; overlap capped at the shared experts' few µs |
| DS4-Flash BI=1 | Deterministic top-k + fusion, two PRs together (each alone has almost no effect); DeepEP LL over MNNVL (`VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL=1` + nvl8 overlay) is the biggest lever; after that the shared experts aux stream helps only under MNNVL | DBO, TP2/DP4, TP4/DP2, the skinny bi_gemm, async-scheduling (already default), splitting the shared MLP to fill the combine window | MoE backend locked to DEEPGEMM (the only BI fp8 MoE) -> a2a locked to deepep LL; per-layer dispatch+combine is serial; cross-layer overlap breaks BI; finally the TTFT floor |
| Nemotron-H BI=0 | PC off (when R3 forces PC on, upstream #58851 is best); R3 finish-only + per-step packed IPC (overlay), makespan returns close to the R3-off floor; then packing per-step logprobs into a list for the frontend (overlay) unsaturates the frontend and BI=0 is 16% faster than BI=1 | Raising the zero-copy threshold (no effect); `max_num_seqs` 256 (slower); PR 57968 keys-only as it stands | PC on (Mamba align mode) has a fixed cost: extra eager launches per decode step, 8K prefill split into two segments |
| Nemotron-H 256K training update (both arms) | BI=1 per-GPU per-micro-batch token cap 64K -> 128K: update only about 5% faster per token (13.1 -> 12.4 s / million tokens), actor reserved 99 -> 168 GiB (184 GiB per GPU, cap 170); small gain, thin margin | Less recompute (some layers, only Mamba layers, none): both arms OOM at 32K-64K per GPU. Per-GPU token cap for BI=0 (no difference across 32K-192K): with CP>1 the recipe disables dynamic batching, each micro-batch is one sequence, and the parameter has no effect | Update is 14-18% of the step, rollout about 80% (sandbox execution, slowest trajectories). Compute the share before tuning; confirm in the resolved config that a parameter takes effect |
BI kernel choice: [BI kernels](bi-kernels.md).
.agents/skills/train-inference-consistency/guides/precision-quant.md
-490
View File
@@ -1,490 +0,0 @@
# Precision, quantization and QAT
The weight conversion chain from optimizer master to vLLM bytes, QAT and low-precision training, and known pitfalls. Numerical contracts without quantization: [numerical-contracts](../workflow/README.md).
## Find by symptom
| Symptom / question | Where to go |
| --- | --- |
| Weights unchanged, probs_diff grows after export | [P1](#p1-scale-rule-differs-from-the-checkpoint-requantization-flips-codes), [P10](#p10-both-sender-and-receiver-quantize), [P11](#p11-receiver-picks-quantization-targets-by-name-suffix-and-misses-fused-moe) |
| NVFP4/FP8 scale or packed bytes differ from the checkpoint | [Conversion chain §2](#2-inventory-the-checkpoints-full-quantization-rules-first), [P1-P5](#quantization-rules-and-scale) |
| NVFP4 rollout: QAT or BF16 actor? | [QAT §1](#1-choose-the-precision-route-first), [§4](#4-how-much-mismatch-qat-can-remove) |
| QAT is on but the mismatch did not change | [P7](#p7-qat-switch-is-on-target-modules-were-not-replaced), [P8](#p8-training-grid-differs-from-the-deployment-grid) |
| QAT loss does not drop, deployed weights byte-for-byte unchanged | [QAT §6](#6-why-qat-does-not-train), [P24](#p24-master-starts-on-the-quantization-grid) |
| Whether to add activation fake quantization | [QAT §4.1](#41-activation-fake-quant-aligned-per-layer-worse-end-to-end) |
| Same operator input, inference module output differs | [P13-P15](#training-side-numerics) |
| Fixed logprob offset under FSDP | [P16](#p16-fsdp-bf16-policy-overrides-the-fp32-parameter-contract), [P20](#p20-dividing-by-temperature-in-bf16) |
| Some LoRA/adapter gradients are always zero | [P17](#p17-fake-quant-cuts-off-the-gradient) |
| Adding a new quantization format | [Conversion chain §3](#3-what-each-format-outputs), [§5](#5-minimal-adaptation-for-a-new-quantization-format) |
| Implementing or changing backward/STE | [Backward validation](backward-validation.md) |
## Key rules
1. "NVFP4/FP8" is only a format name; a different block scale selection, global scale convention, rounding or static/dynamic activation gives different bytes. Inventory the checkpoint's full rules first.
2. Use the rules of a mainstream implementation (TE, ModelOpt, NeMo-RL) that matches the checkpoint; do not invent your own. Verify from the real BF16 release, not from dequantized weights.
3. Training QAT and the exporter use the same quantization function on the same grid (range, axis, scale, rounding).
4. Compute scale on the full logical tensor. A fusion group shares scale (VERL NVFP4: QKV, gate/up fused global scale). Quantizing each shard separately and concatenating is forbidden.
5. Each tensor family has exactly one conversion owner and is quantized exactly once; post-load reordering does not count as quantization.
6. Activation scales (`input_scale`, `input_amax`, `k_scale`, `v_scale`) are calibration values. Pass them through unchanged; they cannot be derived from weights. Read runtime scales from the actual runner.
7. High-precision training + low-precision rollout is not QAT. Neither fake-QAT nor real low-precision GEMM implies a low-precision optimizer or low-precision gradients.
8. Weight QAT lowers Nemotron-H probs_diff mean from about 0.025 to about 0.010; it reaches 0 only when the training forward uses the same kernels as inference. Whether activation fake-quant helps depends on upstream mismatch; measure first.
9. STE backward has its own acceptance: an independent reference, the same dY, a tolerance agreed in advance. A forward 0 diff does not prove the gradient is right.
10. Check that deployed bytes = `requantize(master)`.
11. Start the master from the real BF16 release, not a dequantized checkpoint, whose values sit on grid points so NVFP4 deployed values only change after thousands of steps at small lr. θ0 bytes then differ from the official checkpoint (Nemotron-H: NVFP4 99.776% identical, FP8 24/46 tensors differ); quantify the cost with proxy SFT and have a human confirm it ([QAT §6.2](#62-fix)).
## Terms
- **NVFP4**: FP4 E2M1 values, one FP8 E4M3 block scale per 16 elements, plus one FP32 global scale.
- **block scale / global scale**: per-block scale (`weight_scale`, `scale`) / per-tensor or per-fusion-group scale (`weight_scale_2`, `weight_global_scale`, `scale2`).
- **amax**: maximum absolute value of a tensor (or block), used to compute scale.
- **4over6 / MSE**: TE NVFP4 block scale selection settings; the Nemotron-H checkpoint uses them.
- **E4M3 / E2M1 / UE8M0**: FP8 (4 exponent, 3 mantissa bits) / FP4 (2 exponent, 1 mantissa bit) / unsigned 8-bit exponent (power-of-2 scale).
- **W4A16 / W4A4 / W8A8**: weight bit width / activation bit width.
- **QDQ**: quantize then immediately dequantize, giving a high-precision value on the grid.
- **fake-quant / fake-QAT**: the training forward uses the QDQ value while GEMM stays high precision; not low-precision training.
- **STE**: straight-through estimator; treats the derivative of rounding as 1 (or with a clipping mask).
- **master**: high-precision parameters managed by the optimizer.
- **conversion owner**: the single party that converts a tensor family to the inference format.
- **fusion group**: tensors fused in serving (QKV, gate/up) that must share scale.
- **R3**: MoE route replay; training uses the expert selection recorded by rollout.
- **BI=1 / BI=0**: the `VLLM_BATCH_INVARIANT` switch; 1 is the true on-policy arm where training calls the inference kernels, 0 is the regular baseline.
- **probs_diff**: verl's `training/rollout_probs_diff_*`, the token probability difference between rollout and the training forward.
- **TIS / `rollout_is`**: rollout importance sampling correction.
- **θ0**: policy weights before the first update.
- **oracle**: the reference implementation or value that decides pass or fail.
- **VJP**: the backward function that computes input and parameter gradients from the upstream gradient.
## Conversion chain: from master weights to the bytes the vLLM kernel uses
### 1. Four stages
```text
A. master weights FP32/BF16 optimizer master (plus the FP32 exceptions the model declares)
↓ forward materialization
B. weights the training forward actually sees
BF16; under QAT it is DQ(Q(W)), the dequantized value on the deployment grid
↓ export (gather → HF mapping → quantize/pack)
C. inference-format bytes
packed NVFP4 + block scale + global scale / FP8 + scale /
INT4 packed + scale / layers still BF16 in mixed precision
↓ loader → post-load
D. tensors loaded by vLLM and used by the kernel (including post-load reordering and derived scales)
```
| Stage | Invariant | How to verify |
| --- | --- | --- |
| A | Record dtype separately for parameters and buffers; an FP32 master does not mean forward uses FP32 | Read optimizer parameter identity and dtype; see [§6](#6-dtype-contract-six-boundaries) |
| B | Under QAT, `Q` is the same function on the same grid as the export in C | Unit test: weights of B = dequantize(bytes of C), bitwise |
| C | Scale computed on the full logical tensor; fusion group shares the global scale; excluded layers keep their dtype; recomputing on the same master is idempotent | Compare packed bytes, block scale, global scale per tensor; see [§2](#2-inventory-the-checkpoints-full-quantization-rules-first) |
| D | Each family quantized once; post-load only reorders, does not recompute scale; calibration values arrive unchanged | Per tensor, compare exported bytes with the bytes vLLM uses after loading |
Record the real format and scale of weights, activations and KV (packed UE8M0, FP32 scale, MXFP4/FP8 scope); activation and packing rounding may differ between SM90 and SM100. Before measuring quantization error, confirm the scale comes from the current weights, not an old tensor.
### 2. Inventory the checkpoint's full quantization rules first
1. From `hf_quant_config.json` and the `quantization_config` in `config.json`, list each layer's format, which layers carry `input_scale`, and whether the KV cache carries k/v scale. Read layer by layer; do not infer from the model name ([P12](#p12-inferring-a-layers-precision-from-the-models-overall-name)). Nemotron-H Lightning 30B-A3B: `MIXED_PRECISION`; 46 FP8 layers (`mixer.in_proj`/`out_proj` of 23 Mamba layers); 5935 `W4A16_NVFP4` layers (routed/shared experts, lm_head); FP8 KV cache; 46 fp32 scalar `input_scale`, FP8 layers only.
2. Encode the BF16 release with the checkpoint's own scales; 100% identical packed bytes confirm the release matches the checkpoint.
3. From the **real BF16 release**, requantize with the candidate rule and report per tensor the match ratio of packed bytes, block scale and global scale. Dequantized weights already sit on the grid and pass under any rule ([P2](#p2-verifying-a-rule-with-dequantized-weights-gives-a-false-positive)). Measured match rates: [P1](#p1-scale-rule-differs-from-the-checkpoint-requantization-flips-codes), [P3](#p3-fp8-per-tensor-scale-convention-is-not-amax448).
Decisions:
- Call the rules of a mainstream implementation (TE, ModelOpt, NeMo-RL) directly; do not invent a scale policy. Deviate, and report it, only when the mainstream implementation cannot do it. Nemotron's self-invented "enlarge scale only on overflow" was replaced by an existing TE rule.
- θ0 deployed bytes = the real BF16 release requantized by the rule, which differs from the official checkpoint: on Nemotron-H, NVFP4 bytes are 99.776% identical and FP8 24/46 tensors differ; the proxy SFT cost is about 0.03 nat ([QAT §6.2](#62-fix)). This policy needs human confirmation.
- The BI=1 training forward, BI=0 QAT and the exporter call the same `requantize`; θ0 bytes of the two arms are identical. `input_scale`, `k_scale`, `v_scale` are calibration values passed through from the checkpoint.
- When a single FP8 checkpoint is the source of truth, the training base is a deterministic dequantized derivative, and requantizing need not restore the checkpoint byte for byte. Measure separately: (1) with zero updates, re-exported bytes equal θ0; (2) the precision difference from the original checkpoint.
### 3. What each format outputs
| Format | Artifacts | Must check |
| --- | --- | --- |
| block FP8 | FP8 weight + `weight_scale_inv`, scale per block | Commonly 128×128; FP32/UE8M0 branch, backend, exclusion scope |
| per-tensor FP8 | FP8 weight + `weight_scale` | Scale from the amax of the full weight, not of a shard; convention may not be amax/448 |
| MXFP8 | E4M3 weight + UE8M0 `weight_scale_inv` | 1×32 block direction, scale encoding, pad and high-precision exceptions |
| NVFP4 W4A16 (compressed-tensors, existing `fp4` sender) | `.weight_packed` + `.weight_scale` + `.weight_global_scale` | 1×16, two-level scale, shared outer scale for gate/up, activation contract |
| NVFP4 W4A4 (ModelOpt, Inkling) | packed `uint8` under the original parameter name + FP8 `scale` + FP32 `scale2` + BF16 `input_amax` + `original_shape` | Naming not interchangeable with compressed-tensors; needs a separate adapter |
| INT4 compressed-tensors | `weight_packed` + `weight_scale` + `weight_shape` | group_size, symmetric quantization, pack layout, list of tensors actually quantized |
| DeepSeek V4.1 attention FP8 | F8_E4M3 weight + F8_E8M0 `.scale`, 32×32 block | Not rowwise 32-element MXFP8 |
`weight_scale_inv` is only a loader protocol name; interpret the value per implementation.
### 4. Who converts: conversion owner
A tensor family has one conversion owner. Once the sender has packed, the receiver must not quantize again. dtype reinterpret, layout staging and post-load repack are not quantization.
| Path | Training | Export | Rollout load | owner |
| --- | --- | --- | --- | --- |
| FP8 rollout-only | BF16 | Full BF16 | `quant_weights()` converts to block-FP8 | rollout receiver |
| VERL FSDP QAT NVFP4 | BF16 + fake quant | After gathering the full tensor, `QATQuantizer` outputs packed NVFP4/scale | compressed-tensors | FSDP sender |
| VERL Megatron QAT NVFP4 | BF16 + ModelOpt fake quant | `QATWeightExporter` | ModelOpt | Megatron sender |
| Megatron FP8 E2E | TE FP8 | Bridge-specific FP8 export planning | vLLM FP8 | Bridge / rollout, per the pinned recipe |
| Nemotron-H BI=0 (verl Megatron) | BF16 + ModelOpt weight STE | Requantized by the checkpoint's rules into ModelOpt `MIXED_PRECISION` | ModelOpt | sender (`mixed_precision_export=True`) |
| Pre-quantized checkpoint passthrough | No requantization | Original tensor/scale | Per the checkpoint layout | checkpoint producer |
- Full merge: convert logically complete HF tensors one at a time by their owner; no whole-model all-gather or per-step disk write.
- QAT: the training forward simulates quantization; export generates the actual packed weights and scales.
#### 4.1 FSDP path (VERL)
- With `actor_rollout_ref.actor.fsdp_config.qat.enable=true`, quantizable `nn.Linear` modules become `QATLinear` after the PEFT wrapper and before the FSDP wrap. Forward does NVFP4 W4A16/W4A4 fake quant and ends in `F.linear`; after the optimizer step, `invalidate_all_scales` clears the weight scale cache. [Implementation][qat-linear]
- Export: `FSDPEngine.get_per_tensor_param()` calls `full_tensor()` on each `DTensor`, then `QATQuantizer.quantize_with_fusion()`:
```text
FSDP local shards → full_tensor() all-gather → full HF weight
→ full-weight amax / global_scale → QKV or gate/up fused_global_scale
→ per-row, per-group block scale → packed NVFP4 weight + scales
```
The cost is the all-gather and GPU staging. The quantizer caches to CPU per decoder layer and moves back to GPU, so peak GPU memory is not necessarily one tensor. [Export][qat-export]
- `get_per_tensor_param_shard()` outputs only the local BF16 shard, with no quantizer and no scale. Shard-local block quantization is an unsupported new protocol and must at least solve:
1. aligning shard and block boundaries, or handling cross-shard blocks;
2. a collective for the full-weight amax/global scale;
3. a consistent fused global scale for QKV and gate/up;
4. W4A4 explicitly carrying the activation-derived `input_global_scale`, which cannot be rebuilt from weight shards.
Until these contracts and the oracle exist, quantizing each shard and concatenating is forbidden.
- Non-QAT: export BF16 in HF coordinates; BF16 rollout does not quantize, and `rollout.quantization=fp8` is quantized by the receiver. LoRA merge exception: [P6](#p6-lora-mergetrue-bypasses-the-fsdp-qat-export).
#### 4.2 Megatron path
- VERL: ModelOpt installs fake-quant modules; `MegatronEngine.get_per_tensor_param()` exports HF weights through Megatron-Bridge, and `QATWeightExporter` reads the module quantizer's amax/block metadata to produce NVFP4 and scales. [ModelOpt exporter][modelopt-source]
- VERL FP8 E2E: TE handles FP8 forward/backward in training; when the Bridge declares `export_weight_dtype=fp8`, VERL does not cache generic conversion tasks. The rollout format is decided by `rollout.quantization` and the serving checkpoint contract; FP8 training does not mean exported tensors are in vLLM layout.
- Megatron capability cannot be inferred as FSDP capability.
- Distinct artifacts: the offline converter prepares the initial checkpoint (may write to disk); the online exporter streams in memory; post-load reorders without recomputing scale; a training checkpoint is not a servable quantized HF checkpoint.
#### 4.3 Receiver quantization (rollout-only)
`rollout.quantization=fp8`: the trainer exports BF16; `load_quanted_weights()` in the vLLM worker calls `quant_weights()`, which uses `scaled_fp8_blockwise()` to produce the FP8 weight and `weight_scale_inv`, then replays the quant method's post-load. There are also a DeepSeek V4-specific entry and an Ascend MXFP8 branch with different scale naming. [receiver][receiver-source] When the checkpoint is already FP8/MXFP4, confirm whether it is requantization or passthrough. vLLM cannot quantize BF16 online into ModelOpt `MIXED_PRECISION`.
### 5. Minimal adaptation for a new quantization format
1. From the checkpoint config and the actual loader, list key, shape, dtype, block size, packing and fusion group of weight/scale, and the unquantized tensor families.
2. Trace export -> loader -> post-load and mark the single owner of each family.
3. If a converter is missing, add the minimal conversion at the existing owner's entry; verify on one tensor / fusion group, then cold-load a small model. Changing the dtype is not a conversion.
4. Test zero merge and a non-zero update; get quantize-after-full-gather right before shard optimizations.
One evidence table is enough: tensor family -> source -> export format -> serving format -> owner -> unit/small-model verification.
#### Minimal acceptance
- [ ] Record training dtype, export dtype/layout, rollout dtype/layout and owner.
- [ ] Fixed-token logprob of the zero-adapter path and the base are within the noise baseline of the same backend.
- [ ] After a non-zero update, tokenwise logprob/KL between trainer and rollout stays within bounds.
- [ ] name/shape/dtype/count/bytes of packed weight, scale and ignored families all close.
- [ ] Deployed bytes = `requantize(master)`.
- [ ] FSDP QAT proves scale is computed after `full_tensor()`; with shard export, prove it is not quantized, or provide the four shard-local contracts of §4.1 and a cross-rank oracle.
- [ ] Each family is quantized exactly once.
### 6. dtype contract: six boundaries
For each model, record `checkpoint → load/bridge → wrapper/sharding → forward → backward/optimizer → export/serving` separately.
| Boundary | Megatron | VERL/FSDP | Must check |
| --- | --- | --- | --- |
| Source weights and conversion | Bridge may reorder, split, cast; model may override default conversion | HF loader, dequantized derived checkpoint, init dtype | Same revision; logical value/scale/layout match, not only target dtype |
| Parameter storage | `params_dtype` plus model exceptions; FP32 master/optimizer | HF keep-precision list, PEFT master, FSDP shard | An FP32 master does not mean forward uses FP32 |
| forward parameters | Float16Module converts selectively | FSDP2 `MixedPrecisionPolicy.param_dtype` | Measure inside forward, not only before the wrap |
| Operator computation and output | TE/FLA/TileLang/fused each define product, accumulation and output precision | HF/PEFT kernel and autocast | FP32 output does not prove FP32 product; BF16 output does not prove BF16 accumulation |
| Gradient and optimizer | TP/EP/SP/CP collective, DDP | FSDP reduce dtype, gradient accumulation | FP32 reduction does not restore values rounded in forward |
| Export and serving | Undo the parallel layout, convert adapter/quantization format | gather/export HF adapter | Compare values on the same logical layout |
FP32 protection takes several forms: mark-to-skip before conversion, reset after declaration (which does not restore values already lost), bypassing the default BF16 downcast, and operator-level selection. It is not "all special parameters are raised in precision".
FSDP: HF `_keep_in_fp32_modules_strict` is a load contract and does not become a per-parameter cast exception in FSDP ([P16](#p16-fsdp-bf16-policy-overrides-the-fp32-parameter-contract)). Respect the model's existing contract first, and verify any change of computation precision against a same-model reference and actual serving. `param_dtype=None`, `param_dtype=fp32` and FP32 autocast are different things.
### 7. Code entry points
FSDP [engine][fsdp-source], [module selection][qat-core], [fake quant][qat-linear], [NVFP4 export][qat-export]; Megatron [ModelOpt exporter][modelopt-source]; rollout [receiver][receiver-source]. Historical PR pointers (not the current state): early return on LoRA merge [#5115](https://github.com/verl-project/verl/pull/5115)/[#5411](https://github.com/verl-project/verl/pull/5411)/[#7014](https://github.com/verl-project/verl/pull/7014); fake quant and export quantization [#7149](https://github.com/verl-project/verl/pull/7149)/[#7224](https://github.com/verl-project/verl/pull/7224)/[#7577](https://github.com/verl-project/verl/pull/7577).
[fsdp-source]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/workers/engine/fsdp/transformer_impl.py
[qat-core]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/qat/core.py
[qat-linear]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/qat/linear.py
[qat-export]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/qat/quantizer.py
[modelopt-source]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/modelopt/qat_weight_exporter.py
[receiver-source]: https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/vllm/vllm_quant_utils.py
## QAT, low-precision training and STE
### 1. Choose the precision route first
Rule out input, weight, adapter, layout and operator correctness errors first, then choose precision:
1. When GPU memory allows and rollout can use BF16, use BF16 for both training and inference; check FP32 exception parameters one by one.
2. When rollout must be low precision (such as NVFP4), use weight QAT or a BF16 actor for training. On Nemotron-H the weight QAT arm is stable. Whether `rollout_is=token` stabilizes the BF16 actor is **not tested**.
3. Consider real low-precision training last: backward/kernel support and maintenance cost more.
Same dtype is not same computation, and QAT is not bit-exact. QAT and low-precision training shrink the **weight value** gap; the kernel and reduction-order gap goes only when the training forward calls the inference kernels (the BI=1 arm of the [BI switch recipe](recipe-and-acceptance.md)). A natively low-precision model does not necessarily run with `dtype=bf16`; do a capacity pre-check first.
### 2. What the four kinds of path are
| Path | Training forward | backward / optimizer | Is it QAT |
| --- | --- | --- | --- |
| High-precision training + low-precision rollout | BF16 weights | High precision | No; quantized only at export |
| fake-QAT | `DQ(Q(W))`, GEMM is still BF16 | STE sends the gradient back to the high-precision master | Yes |
| Real low-precision GEMM (FP8/MXFP8/NVFP4) | TE changes the compute kernel | High-precision master/exception tensors remain | Different from fake-QAT |
| Frozen quantized base + high-precision LoRA | Dequantized base + adapter | Only the adapter is trained | No |
- Neither fake-QAT nor real low-precision training makes the optimizer/gradients low precision. W4A16 quantizes only weights; W4A4 also quantizes inputs.
- Real low precision: NVFP4 forward is not FP4 backward; configurations include high-precision backward and backward on the dequantized operand.
- Only weight-only formats can compute scale from the full weight; formats with activation quantization also need activation scale/calibration state, which weights cannot provide. QAT simulation and export follow the same format rules, but scales are not necessarily constant during training.
- When LoRA is applied separately, do not fake-quant the merged base+adapter yourself.
- Open-source frameworks (verl [#6992](https://github.com/verl-project/verl/pull/6992), [#7971](https://github.com/verl-project/verl/pull/7971), NeMo-RL) essentially only do W4A16; activation inconsistency only appears at W4A8.
### 3. Implementing fake-QAT
Forward form (Nemotron-H):
```text
w_visible = w + (dequant(requantize(w)) - w).detach()
```
- `requantize` is the same function as the exporter; replace the library's built-in quantizer when its rules differ.
- Requantization runs on every forward, and again under full recompute; Nemotron-H step time is about +60%.
- A temporary `QuantModuleRegistry` replacement and the old class-attribute patch give the same probs_diff mean on the same quick workload (0.0095/0.0091/0.0100 vs 0.0096/0.0092/0.0101); module hit check: [P7](#p7-qat-switch-is-on-target-modules-were-not-replaced).
### 4. How much mismatch QAT can remove
Nemotron-H (`rollout_probs_diff` mean, NVFP4 mixed-precision deployment, single node 4×GB200):
| Approach | Training forward | probs_diff mean |
| --- | --- | --- |
| BF16 training + quantized rollout | BF16 weights | About 0.025; no rollout correction |
| Weight QAT (STE) | Dequantized deployed weights, BF16 GEMM | About 0.010 (pearson 0.986→0.997); three 3-step smoke runs of one recipe, 0.0091-0.0102 |
| Weight QAT + activation fake-quant | Same, plus static FP8 activations | About 0.0107, about 10% worse than weights only |
| true on-policy (BI=1) | Deployed bytes + the same kernels as vLLM | max = mean = 0 at every step for 100 steps |
The QAT arm's probs_diff max is still 0.73-1.0; a small mean does not rule out large per-token deviations. Other models (different metrics, not thresholds): DeepSeek-V4 BF16 train + FP8 rollout has a first-step logprob drift of about 0.023 and still reports rising reward/eval ([DeepSeek][dsblog]); Kimi K3 frozen quantized base + LoRA has a train/rollout KL noise floor of about 0.002, a Schulman KL estimate, not a mean absolute logprob diff ([K3][k3blog]).
#### 4.1 Activation fake-quant: aligned per layer, worse end to end
Adding vLLM's static FP8 activation fake-quant to training aligns single-layer output to the BF16 noise floor, but end-to-end probs_diff gets about 10% worse: FP8 rounding is discontinuous, so any element-wise upstream mismatch flips codes and amplifies the difference. To judge, add relative noise δ only to the training-side input and sweep the output error; the crossover was about δ≈1%, and activation quantization hurts when upstream mismatch exceeds it. This is one measurement on one model, not a threshold; for a new model, measure upstream mismatch and repeat the sweep. Nemotron-H did not merge it.
Rule: the weight-side range/axis/scale/rounding must match serving ([rule 3](#key-rules)); matching the activation side depends on upstream mismatch and is not unconditional. Mainstream frameworks (verl, NeMo-RL, slime) use the same **dynamic** activation rule on both sides and do not reuse an external checkpoint's static activation scales; the NeMo-RL Nemotron recipe quantizes only routed experts and leaves Mamba projections, attention and lm_head in BF16. On DeepSeek V4.1, KV quantization simulation adds only about 1.93% improvement on top of R3. Training-side simulation of the FP8 KV cache is not covered.
#### 4.3 When QAT is worth it
Evaluate unified FP8/QAT only when evidence shows the train-inference quantization difference is the main limit; "no score gain" is not a reason to switch precision first. For DSv4.1 rank2 LoRA, the post-merge quantized delta is only 5.1-7.7% of the pre-quantization delta (a norm ratio, not information retention), which only supports treating step size/quantization as controlled variables. The long-running Qwen Flash-Next experiment is same-source BF16/BF16 and does not validate the FP8 rollout route.
### 5. STE contract
The STE surrogate (identity or clipped), whether dX uses the master or dequantized value, scale handling and verification are in [Backward validation](backward-validation.md#write-down-the-gradient-contract-first). By default dX uses the dequantized value; when using the master value, quantify the bias.
### 6. Why QAT does not train
Train-inference consistency is the strictest QAT: forward uses quantized weights, the high-precision master accumulates updates, STE passes the gradient back. Whether it trains depends only on the master's distance to the nearest decision boundary. In W4A16 (the FP4 code is decoded to BF16 in a register and goes through the BF16 tensor core; the gain is 1/4 the weight bandwidth), W4A4 and W8A8 alike, the computed values are "code × scale".
#### 6.1 Dead zone
- Each NVFP4 element has only 8 values, 0/0.5/1/1.5/2/3/4/6 × scale, with RNE rounding. From a dequantized checkpoint the master sits on grid points, half a spacing from every decision boundary, and must travel that half spacing before the code changes.
- At lr 1e-6 one step moves the master about 0.01%; a few hundred steps move it far less than half a spacing, so deployed values stay byte-for-byte unchanged. RL gradients are more random than SFT, closer to a random walk, and need more steps.
- Result: NVFP4-covered parameters (on Nemotron-H, routed/shared experts and lm_head, the bulk) have zero deployed change, and validation NLL is clearly worse than the native BF16 baseline; freezing these weights explains about half the gap.
- Recomputing block scale does not flip codes earlier: from grid points requant is idempotent, the E4M3 scale only changes code after block amax moves 3-6%, and a tiny global scale change only rescales the whole tensor.
- Mainstream QAT (NeMo-RL QARL, slime/Kimi INT4, Jet-RL) starts from the original BF16 weights, which lie between grid points, so any small drift flips a proportional share of codes and there is no dead zone. INT4 with 1×32 groups from a dequantized checkpoint stalls the same way.
- FP8 is barely affected: floating-point relative spacing is 6-12%, small elements have small absolute spacing, and some elements are always near a boundary. DS4 also starts from dequantized FP8 and is affected only mildly.
#### 6.2 Fix
Preferred: use the original BF16 weights (the real BF16 release) as the master and requantize every step by the same rule (Nemotron-H: TE 4over6, E4M3 upper bound 256). State two points:
- **θ0 is not equal to the official checkpoint bytes.** Requantizing from the real BF16 release, 99.776% of Nemotron-H's NVFP4 bytes are identical and 24 of 46 FP8 tensors (ModelOpt calibrated scale) differ. "The official checkpoint was generated by this rule" holds only on dequantized weights, where MSE and MAE cannot tell rules apart on grid points ([P2](#p2-verifying-a-rule-with-dequantized-weights-gives-a-false-positive)). Proxy SFT gives the cost: 300-step val NLL A-bf16 3.4685 vs C-dequant 3.4641 / C-bf16 3.4373, about 0.03 nat, present at θ0 and not growing during training. This policy needs human confirmation.
- **Three acceptance items** (all pass on Nemotron-H): (1) with zero updates, re-exported bytes equal θ0 (not the official checkpoint); (2) fast-tier RL θ0/θ1/θ2 is 0 diff and deployed bytes start changing; (3) proxy SFT gap to the native BF16 baseline is within budget.
Alternative, keeping θ0 bytes: add a fixed small jitter to move the master off grid points (toward the neighboring cell, at most half a step; after requant check per block that bytes are unchanged, and roll back changed blocks). Not recommended: stochastic-rounding requantization (about 1.4× deployment error, needs a self-written quantizer, no mainstream precedent), raising lr (under Adam equivalent to QuRL UAQ, only shortens the dead zone, 10× lr is outside the usual RL range), and frozen base + LoRA (QeRL, gives up full-parameter updates). Whichever you choose, state that from a dequantized checkpoint at lr 1e-6, NVFP4 deployed values do not change for roughly the first 10³-10⁴ steps. See [P24](#p24-master-starts-on-the-quantization-grid).
[dsblog]: https://www.lmsys.org/blog/2026-04-25-deepseek-v4/
[k3blog]: https://www.lmsys.org/blog/2026-07-27-kimi-k3-day0-support/
## Known pitfalls
Ordered along the conversion chain: quantization rules -> export and receive -> training-side numerics -> experiment measurement basis. P24 sits under training-side numerics by content.
### Quantization rules and scale
#### P1 scale rule differs from the checkpoint, requantization flips codes
- Symptom: weights unchanged, rollout shifts after a requantized export.
- Cause: recomputing scale on the dequantized master with the ModelOpt default global = amax/(6·448) differs from the checkpoint's TE 4over6 rule and flips about 40% of Nemotron-H's FP4 codes; recomputing with TE 4over6 is idempotent and byte-for-byte unchanged.
- Fix: use a mainstream rule that matches the checkpoint; training QAT and the exporter call the same `requantize`.
- How to verify: requantizing the same master twice is byte-for-byte identical; requantize from the real BF16 release and compare per tensor with the checkpoint. Nemotron-H:
| Rule | Match with the checkpoint |
| --- | --- |
| TE NVFP4 4over6 + MSE, E4M3 upper bound 256, global = amax/1536 | packed 99.776%, global 100%; no setting closes the remaining gap |
| ModelOpt default global = amax/(6·448) | packed about 50% |
| Standard FP8 per-tensor amax/448 | See P3 |
- Source: procedure in [Conversion chain §2](#2-inventory-the-checkpoints-full-quantization-rules-first).
#### P2 Verifying a rule with dequantized weights gives a false positive
- Symptom: several candidate rules all report a 100% match.
- Cause: dequantized weights already sit on the grid and do not change under any rule.
- Fix: verify from the real BF16 release.
- How to verify: at least one known-wrong rule fails the test.
#### P3 FP8 per-tensor scale convention is not amax/448
- Symptom: FP8 layers differ from the checkpoint after requantization.
- Cause: Nemotron-H's official Mamba projection scale is about 1.5×amax/448; 24 of 46 FP8 tensors differ from amax/448.
- Fix: derive the convention from the checkpoint, or pass the checkpoint's scale through.
- How to verify: report the scale match rate per tensor.
#### P4 Computing scale separately on each shard
- Symptom: packed weights concatenated after an FSDP/TP sharded export differ from quantizing the full tensor.
- Cause: blocks cross shards; global scale and fused global scale need the full tensor.
- Fix: `full_tensor()`/gather, then quantize; shard-local quantization must first meet the four contracts of [Conversion chain §4.1](#41-fsdp-path-verl).
- How to verify: evidence shows scale is computed after the gather.
#### P5 gate/up not quantized as a pair
- Symptom: wrong values after serving fuses gate/up, or the exporter reports a missing pair.
- Cause: the two HF tensors are fused in serving and need a shared outer scale.
- Fix: quantize as a pair; raise an error when a complete pair is missing, and test pairing across safetensors shards.
- How to verify: test pair completeness, cross-shard gate/up and a TE bitwise reference on matching hardware; a test existing does not mean it ran and passed.
#### P6 LoRA merge=true bypasses the FSDP QAT export
- Symptom: QAT and LoRA merge are on, and the export is not the QAT quantizer's packed weights.
- Cause: VERL FSDP returns `_merged_lora_per_tensor_param()` directly for `merge=true`, skipping `QATQuantizer`; with `rollout.quantization=fp8` the receiver may still requantize. The `_qat_enabled` check still follows the `merge=false` branch.
- Fix: the first merge fallback does gather -> merge -> quantize -> yield -> release one full tensor at a time.
- How to verify: check actual module replacement, naming and export scope; both switches being on does not make the combination usable.
- Source: [VERL engine](https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/workers/engine/fsdp/transformer_impl.py).
#### P7 QAT switch is on, target modules were not replaced
- Symptom: the QAT arm's probs_diff equals the BF16 arm's.
- Cause: VERL `QATLinear` enters through `nn.Linear` and does not cover three-dimensional expert Parameters or arbitrary PEFT wrappers; the framework's built-in QAT changes only the GPT layer spec, while Mamba hybrid models read the hybrid stack spec.
- Fix: attach the quantizer after building the model by the actual spec; use the exporter's `requantize` as the quantizer.
- How to verify: every rank prints the count and list of replaced modules and fails closed (139 for Nemotron-H, the same on all four ranks).
- Source: [Conversion chain §4.1](#41-fsdp-path-verl), [BI=0 integration pitfalls](recipe-and-acceptance.md#bi0-arm-integration-pitfalls).
#### P8 Training grid differs from the deployment grid
- Symptom: QAT is declared, the mismatch remains.
- Cause: the converter and the trainer use different group sizes (for example 32 vs 128). An external TPU case is bitwise 0 on the FP8 per-channel grid, but the native weights are 128×128 block, and the author notes the result may not hold on a different grid.
- Fix: align quantized module scope, block/group size, axis, scale dtype/rounding, global scale scope, gate/up fusion, activation scale, TP/ETP partition and high-precision exceptions. FP4 has 1 mantissa bit, so error comes mainly from block partitioning; block size and direction must match the native weights. MiMo-V2.6 does QDQ on the experts after every update under the rollout MXFP4 GEMM's constraints.
- How to verify: training-side QDQ output = dequantize(deployed bytes), bitwise.
- Source: unified NVFP4 constraints in the [technical blog](https://www.lmsys.org/blog/2026-07-29-mxfp8-nvfp4-rl/) (its bit-exact quantizer test verifies the format, not whole-model bit-exactness).
#### P9 Activation scale derived from weights
- Symptom: W4A4 export raises, or the exported activation scale differs from serving.
- Cause: activation scale depends on actual input. VERL W4A4 export raises when an initialized `input_global_scale` is missing; Inkling's `input_amax` cannot be derived from BF16 weights.
- Fix: reuse unchanged from the QAT calibration state or rollout checkpoint; pass through `input_scale`, `k_scale`, `v_scale`.
- How to verify: manifest activation scales are byte-for-byte identical to the checkpoint.
- Source: [VERL quantizer](https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/qat/quantizer.py).
### Export and receive
#### P10 Both sender and receiver quantize
- Symptom: deviation larger than one quantization; or a packed tensor treated as BF16 and quantized again.
- Cause: a tensor family has two conversion owners; the Bridge already outputs packed tensors but the generic packer does not recognize them.
- Fix: one owner per family; pass packed tensors through.
- How to verify: the evidence table lists the owner per family; count/bytes close.
- Source: [Conversion chain §4](#4-who-converts-conversion-owner).
#### P11 Receiver picks quantization targets by name suffix and misses fused MoE
- Symptom: under `rollout.quantization=fp8` some experts are still BF16.
- Cause: `gate_up_proj`/`down_proj` do not end in `.weight` and are skipped; fixed by VERL [#7443](https://github.com/verl-project/verl/pull/7443).
- Fix: check coverage using checkpoint names together with runtime modules/dtype; review the list for each new model.
- How to verify: list the serving dtype of every family.
- Source: [receiver](https://github.com/verl-project/verl/blob/753aed3e1c286ba6825a74342b28669e72c083ea/verl/utils/vllm/vllm_quant_utils.py).
#### P12 Inferring a layer's precision from the model's overall name
- Symptom: diagnosis goes in the wrong direction.
- Cause: in the Nemotron-H "NVFP4" checkpoint, `backbone.layers.16.mixer.out_proj` has `quant_algo` FP8.
- Fix: read `hf_quant_config.json` layer by layer.
- How to verify: the diagnosis states the actual `quant_algo` of the first differing layer.
### Training-side numerics
#### P13 scale taken from the wrong object
- Symptom: the unit-scale unit test passes; real attention output and old-logprob fail.
- Cause: the training adapter quantizes Q with `k_scale`. On the real runner `q_scale=1.0`, `k_scale=0.03173828125`, `v_scale=0.003326416015625`.
- Fix: read scale from the runner/module actually executed; do not infer a default from a missing checkpoint field or another tensor.
- How to verify: replay with real non-unit scales and frozen Q/K/V; after the Nemotron-H fix, 5,832,704/5,832,704 BF16 outputs are bitwise identical.
#### P14 Quantization modules with the same name take different dispatch
- Symptom: calling `ops.scaled_fp8_quant` + `flashinfer_scaled_fp8_mm` directly matches training bitwise, yet the inference linear module differs in 1,361 BF16 elements with the same scale.
- Cause: the module selects the **compiled** `forward_native`, which has 44 more differing FP8 elements than uncompiled native and the CUDA custom op; those two agree. Which compile transformation is responsible is not confirmed.
- Fix: under Nemotron-H BI=1, static per-tensor `QuantFP8` switches to the device custom op (GEMM and format unchanged).
- How to verify: with the same operator input, weights and scale, compare FP8 bytes, GEMM output and module output in turn, calling the native/CUDA/actual module paths separately. When the direct GEMM matches training and the module does not, check quantization dispatch inside the module first; do not change GEMM.
#### P15 Quantizer is not in the `nn.Module` tree, the patch misses layers
- Symptom: the patch "ran", the target layer still has 1,361 differing elements.
- Cause: traversing `model.modules()` finds only 6 registered modules and misses quantizers on plain objects such as `linear.quant_method.kernel.quant_fp8`.
- Fix: traverse the registered tree, the quant-method kernel objects and that version's legacy fields together.
- How to verify: record per worker the hit count (52 after the fix), target layer names, actual `_forward_method` and quantized output bytes; fail if expected targets are zero or a key layer is missing. Private field paths are only a versioned case.
#### P16 FSDP BF16 policy overrides the FP32 parameter contract
- Symptom: Hy4 four-layer fixed-token logprob MAE 0.022982/0.023052.
- Cause: HC and attention sink, FP32 after HF load, are downcast by the uniform BF16 FSDP policy; a later `.float()` in the operator cannot restore them. The router's FP32 conversion is also overridden by the outer BF16 autocast, which changes the expert set.
- Fix: opt in with `wrap_policy.preserve_original_param_dtype=True` (`param_dtype=None`, BF16 autocast kept). This is a whole-model policy, not a per-parameter whitelist; measure communication/memory impact; it does not raise a BF16 base to FP32. Disable autocast inside the router forward.
- How to verify: with original dtype kept, singleton, two ranks and unwrapped are exactly identical; the router has a regression test.
#### P17 fake-quant cuts off the gradient
- Symptom: in DSv4.1 LoRA training, 31 of 160 B matrices are zero, all in `self_attn.kv_proj`.
- Cause: `_fake_quant_fp8_block` goes through an FP8 cast on the autograd path, so small gradients underflow to 0; `_fake_quant_fp4_block` has no STE (integer comparison and LUT gather, so gradient flows only through scale, and the ue8m0 branch cuts even that).
- Fix: keep the quantized forward and use identity for the input gradient: `output.detach() + (x - x.detach())`.
- How to verify: a 1e-6 upstream gradient is preserved exactly; after the fix 160/160 B matrices are non-zero. The loss still rises, so this does not prove training improved.
#### P18 Training forward and export do not quantize at the same place
- Symptom: the model already has FP4/FP8 STE, the mismatch remains.
- Cause: DSv4.1's existing STE acts on indexer/compressed/window KV activations, not as weight fake-quant on `W+ΔW`; the training forward is unmerged base + LoRA, while export is an FP32 merge followed by block32 FP8/E8M0 conversion.
- Fix: confirm whether training sees the post-export rounding before deciding on weight QAT.
- How to verify: sample the actual dtype inside forward, not only the code path.
#### P19 KV cache aligned only on the dtype label
- Symptom: attention results differ with the same weights.
- Cause: in the old vLLM image, KV is FP8/block64 for the non-RoPE part and BF16 for the RoPE part, and the training side handles values differently.
- Fix: set `kv_cache_quantization="fp8_ds_mla"` explicitly in training, use the same rounding for forward with and without gradients, and keep the gradient with STE. Re-verify after the upstream KV layout changes.
- How to verify: compare attention output with the same weights.
#### P20 Dividing by temperature in BF16
- Symptom: logprob has a fixed offset.
- Cause: the old VERL FSDP path divides by temperature in BF16/FP16, then raises precision.
- Fix: `logits.float()` first, then divide by temperature, even when temperature is 1.
- How to verify: both scorers match on the same logits.
#### P21 Derived scale after resume still belongs to the old weights
- Symptom: the first rollout after resume is wrong.
- Cause: BF16 parameters were updated, but the FP8 source-scale registry still trusts the initial weights.
- Fix: invalidate derived caches in the weight restore lifecycle.
- How to verify: first-step probs_diff after resume is on the same measurement basis as before (Nemotron-H kill->resume acceptance: still 0 after resume).
- Source: [checkpoint guide §13](checkpoint-resume.md#13-stale-quantization-scale-on-the-first-step-after-restore).
#### P24 Master starts on the quantization grid
- Symptom: QAT loss barely drops and NVFP4 deployed weights are byte-for-byte unchanged while the master moves.
- Cause: a master initialized from a dequantized checkpoint sits on grid points ([QAT §6.1](#61-dead-zone)).
- Fix: master from the real BF16 release, requantized every step by the same rule, cost quantified with proxy SFT; without the original weights, a fixed within-cell jitter. Not stochastic rounding, larger lr, or LoRA.
- How to verify: the three acceptance items of [QAT §6.2](#62-fix).
### Experiment measurement basis
#### P22 The quantization oracle itself is wrong
- Symptom: the GPU block32 quantization gate fails (zero block expected 2.27373675e-13, actual 1.16415322e-10).
- Cause: the oracle assumes Triton's epsilon position; the CUDA UE8M0 path also clamps scale to 1e-10 before taking the power of 2.
- Fix: fix the oracle, not the kernel. CUDA and Triton are not thereby bitwise identical on tiny/zero blocks.
- How to verify: check the oracle's epsilon/floor position against the kernel source before judging a failure.
#### P23 Attributing every logprob difference to quantization
- Symptom: the difference persists after changing precision.
- Cause: routing selection, attention/kernel and tokenization/mask all produce differences. R3 constrains routing and does not make FP8 GEMM equal BF16; TIS is a statistical correction and does not fix wrong weights.
- Fix: rule out protocol errors first, then record the quantization/numerical noise floor.
- How to verify: after ruling out protocol items (routing, mask) one by one, measure the quantization noise floor on the same input.
.agents/skills/train-inference-consistency/guides/recipe-and-acceptance.md
-113
View File
@@ -1,113 +0,0 @@
# Recipe, evaluation and paired acceptance
Which recipe items must match the reference implementation mathematically (not by name), what puts two evaluations on the same measurement basis, and how to pair two arms so a conclusion can be drawn. The reference recipe for train-inference consistency is [verl#7688](https://github.com/verl-project/verl/pull/7688) (DS4 exact 0-diff true on-policy). Which reference a model aligned with, and the values it ended up using, are per-model facts and are not recorded here.
## Find by question
| Question | Where to go |
| --- | --- |
| Does a 5% gap between two arms count as a gap | [Paired acceptance](#paired-acceptance): noise floor, median over multiple pairs, confounders |
| The validation score only takes a few discrete values | [Evaluation contract](#evaluation-contract): under greedy the score is a multiple of 1/number of problems |
| How to design the BI=0 baseline arm | [How the two arms are defined](#how-the-two-arms-are-defined) |
## Key rules
1. Align the math, not the names; read the final CLI/resolved config, not the README or the model name.
2. Tabulate every deviation from the reference: item, reference value, our value, why, which conclusion it affects.
3. One recipe, one image, a switch selects the arm (following DS4 `EXACT_ALIGNMENT`, default BI=1; BI=0 goes through Bridge/mcore + R3 + QAT); deviations live only in the switch.
4. If any of the six evaluation items differs, it is not the same score: dataset version, prompt template, sampling parameters, number of repeats, scorer, length cap.
5. Rising training reward, probs_diff of 0, and a held-out score gain are three separate conclusions; none substitutes for another.
6. Before running, fix what is compared, how it is computed, the pass line and the noise floor; plant a known error to confirm the criterion can return FAIL.
7. One artifact answers one question: run and report performance, quality and consistency separately.
8. List confounders first, then compare; each is either aligned or quantified separately ([paired acceptance](#paired-acceptance)).
9. Keep failure records as they are and append new results; keep reruns and changed questions apart.
## One recipe, one image: deliver true on-policy and the regular baseline with the BI switch
Following DS4's EXACT_ALIGNMENT: one delivery image, one recipe, and `VLLM_BATCH_INVARIANT`, defined by the recipe, selects the whole arm. The arms may differ in training engine, precision and route replay; the switch is only the single entry point. The pitfalls below come from Nemotron-H Lightning NVFP4 (Mamba hybrid MoE); a new model must re-check them item by item, not copy the config.
### How the two arms are defined
BI=1 is true on-policy: the training forward calls the vLLM kernels and reads the deployed bytes, rollout takes the vLLM BI path, and the deployed bytes are sent down unchanged. BI=0 is the regular baseline: training takes the default MCore path (verl Megatron engine/Megatron-Bridge, without the alignment implementation of [C6](../workflow/README.md#gate-table)), and rollout uses vLLM's non-BI kernels with the configuration audited per [C1](../workflow/README.md#gate-table). For BI=0, decide the items below per model and record who approved each:
| Decision item | When needed | Nemotron's choice (example) |
| --- | --- | --- |
| Training precision: bf16 or weight QAT (STE) | QAT exists only when deployment uses quantized weights | QAT by default, bf16 optional |
| MoE route replay (R3) | MoE only | Always on |
| Requantize from master at export | Quantized deployment only | Requantize by the checkpoint's rules at every export |
| Rollout importance sampling correction (TIS etc.) | Arms whose train-inference mismatch is not 0 | Not enabled |
| rollout MoE backend | When the backend picked by `auto` does not run | Explicit `flashinfer_cutedsl` (`auto` picks Marlin, which reports a PTX version error), matching the BI0 arm of the BI=1 bench |
Do not build BI=0 as "the BI=1 actor with non-BI kernels swapped in". That requires changing the alignment implementation, and the result is neither the regular baseline nor true on-policy; a human rejected this design when it was tried.
### Rules
- **Same image, same recipe.** If the recipe lacks a switch, add it on the delivery branch first, then build one image shared by both arms; no separate image, recipe or source mount for a single experiment. Delete the old standalone baseline recipe.
- **Keep output and checkpoint paths separate.** Suffix only BI=0; BI=1 paths stay as before. With shared paths, `resume_mode=auto` resumes from the other arm's checkpoint. Renaming the run breaks an in-progress resume chain (it cannot find the first-segment checkpoint), so a running chain does not change images.
- **The BI=1 command is byte-for-byte unchanged.** The recipe must offer a mode that only prints the fully resolved training command without starting a job (Nemotron uses `DRY_RUN=1`) and a mode that only runs hydra compose. For each mode × optimizer, print the command from the recipe before and after the change and `diff` (or compare sha256): BI=1 must be identical; BI=0, against the original baseline recipe, may differ only in names, paths and newly added explicit switches. Every compose combination must return rc 0.
- **The header comment states only what was measured.** Reviews have caught untrue claims such as "the BI=0 actor uses vLLM's default kernels". Known rollout config differences between arms (such as R3 forcing prefix caching on) go into the comparison report.
- Delete untested switches (such as unverified mxfp8 and route replay variants) instead of leaving them for users to trip over.
### BI=0 arm integration pitfalls
| Problem | Consequence | Handling |
| --- | --- | --- |
| THD packing does not pass `total_tokens` into `PackedSeqParams` | Mamba gets no `seq_idx` and scans the whole pack as one sequence; logprob is silently corrupted | Pass it when CP=1; raise an error for unverified CP>1 |
| Actor needs the BF16 release, rollout needs the quantized checkpoint, both share `model.path` | Actor is built from quantized weights | Separate `hf_weights_path` |
| MTP does not turn off cleanly (the provider appends the pattern back) | Extra untrained layers | Clear both `mtp_num_layers` and the MTP pattern |
| Expert bias is updated during training by default | Routing drifts from inference | `moe_router_bias_update_rate=0` |
| Framework's built-in QAT changes only the GPT layer spec | Hybrid models read the hybrid stack spec, so QAT has no effect | Attach the quantizer after building the model from the TE spec, using the same `requantize` as the exporter |
| QAT library incompatible with the TE grouped linear signature | Startup error or wrong quantization | Quantize weights only; unit test that forward = deployed bytes and the gradient reaches master |
| param + optimizer offload, HF safetensors export | Host OOM on save | Chunked optimizer offload, save only dist ckpt, see [checkpoint and host memory](checkpoint-resume.md#memory-budget-host-and-gpu) |
R3 checks: R3 in vLLM requires prefix caching and Model Runner V2; the capture is indexed by **absolute layer number** (non-MoE layers are zero columns), and the training side selects the column by global `layer_number`, which must also line up under PP; exclude the MTP router. Known pitfalls: omitting `routed_scaling_factor` causes a drift of about 0.28; a monolithic backend that bypasses top-k capture replays all zeros.
### Smoke criteria
- BI=1 quick: `rollout_probs_diff` max = mean = 0 at every step.
- BI=0 quick: mean has the same magnitude as the independent reference (Nemotron QAT about 0.010, BF16 about 0.025), max may approach 1, grad norm is finite.
- When a step's grad norm is exactly 0, check within-group rewards first: if the whole group has the same reward, advantage is 0 and it is not a numerical problem.
- After changing the BI=0 QAT implementation, compare probs_diff against the old implementation on the same quick workload, and print the number of replaced modules (identical on every rank).
## Evaluation contract and paired acceptance
One artifact answers one question. Bitwise logprob on fixed tokens proves only the scoring of the same token history; natural generation tasks (GSM8K, AIME) answer quality; fixed-length workloads (Rust Bench) answer throughput. Token sets and timing differ, so none substitutes for another. After the numerical path, deployed weights or backend change, every in-scope quality and speed check needs fresh evidence from the **current candidate**; out-of-scope items are marked not verified, and no acceptance tasks are added on our own.
### Evaluation contract
If any of the six items differs, it is not the same score: dataset version (AIME24 ≠ AIME25); prompt/chat template and thinking mode; sampling parameters (temperature/top-p/top-k, stop tokens; greedy n=1, sampled average and pass@8 differ; pass@1 at n=8 is not greedy accuracy); number of repeats; scorer (answer region detection, boxed extraction, tail truncation length, reward extractor); length cap (training and eval caps may differ; 4K vs 8K results are not a pure capability difference).
- Under greedy, n repeats give the same result and the score is a multiple of 1/number of problems: with 30 problems one problem is 3.33 percentage points, and smaller differences are invisible.
- Rising training reward, probs_diff of 0, and a held-out score gain are three separate conclusions; dynamic sampling changes the sample composition of the training reward. Read the truncation ratio together with accuracy.
- A gain from relaxing the token budget is not a training gain. Charts come from raw metrics; do not fabricate a step 0.
- W&B: the run config behind a curve, the step basis (rollout or optimizer step) and the checkpoint version must be traceable; for cumulative metrics, take per-phase differences.
- Changing the scorer or migrating a reward handler: run both scorers offline on the same responses and inspect only the samples they judge differently; do not switch scorers for the sake of "alignment" and mark another model's valid output wrong over a format difference.
### Paired acceptance
**Freeze the criteria first, then run.** Fix the model/checkpoint content, image digest, source/overlay, hardware/nodes, runner/parallelism, weight and KV precision, the actual attention/dense GEMM/MoE/all-to-all backends, and the resolved Graph and cache config; fix what is compared, how it is computed, the pass line and the noise floor; plant a known error to confirm the criterion can return FAIL. Tune BI=0 into a reasonably strong baseline at the agreed precision and have a human audit it; do not weaken it to narrow the gap. When the candidate changes quantization dispatch, attention, MoE, exported weights or image, re-evaluate which old evidence is invalidated ([evidence ledger](../practices/evidence-and-experiments.md#6-candidate-evidence-ledger)).
**List confounders first, then compare.** Tabulate every difference between the arms; each is either aligned or quantified separately; a delta from several simultaneous changes is not attributed to one switch. Example: the two Nemotron-H arms differed at once in PC on/off, `gpu_memory_utilization`, training backend (Bridge vs mLite) and the R3 return path.
**Quality pairing**: same problem IDs, template, seed, sampling parameters, max length and scorer, no retries; record per problem whether each side is correct; `both_correct`/`only_bi0`/`only_bi1`/`both_wrong` must sum to the problem count; report each side's score, the absolute percentage-point difference and the agreed threshold. Close scores do not mean identical per-problem results, and passing quality does not mean bitwise-equal logprob.
**Fixed-length throughput**: the agreed workload is vLLM Rust Bench 8192 input/1024 output, ignore EOS, c=1/8/32/64; record actual input/output tokens, completed requests, error count, warmup, and per concurrency point the wall time, throughput and TTFT/TPOT. List any real rollout workload separately; do not mix fixed-length and natural-EOS results in one ratio.
### Pairing and computation
Before testing, a human freezes: number of repeats, pairing order, aggregation formula, the budget/pass rule for each target c, and how an invalid leg is identified and re-tested. The suggested statistic is the median per-round paired loss `median_r(1 - T_BI1[c,r] / T_BI0[c,r])`; if `1 - median_r(T_BI1[c,r]) / median_r(T_BI0[c,r])` is chosen instead, label it "ratio of median throughputs". The two are not equivalent and cannot be swapped after seeing results; without a reviewed statistical basis, report only measured values and do not declare the budget passed. Historical results keep their original formula and are not silently re-judged.
1. For each leg, record on every worker the image/package hash, actual kernel/backend, resolved config, Graph state, cache hits, and request and token statistics; verify client and server exit normally, separately.
2. Same node, same resources, no interfering load, crossed order (AB/BA/AB), each side warmed up separately every round; do not pair each side's fastest run picked across node pools.
3. Check workload and exit receipts item by item first. A mismatched leg is `inconclusive`; do not compute BI loss from throughputs with different request volumes.
4. Report raw throughput and loss `1 - T_BI1/T_BI0` for every round, then compute the formal result per c with the reviewed formula; keep failed/invalid rounds and the reasons for re-testing. Averaging across c must not hide a required c over the limit; keep single-round and aggregate overruns apart.
Optimization screening may test only the previously failing concurrency points, but it only decides whether to continue. Take the median over multiple pairs and report single pairs separately: if a candidate round sits right at the budget line, it cannot pass on the old baseline's historical median; re-pair both sides in crossed order on the same node, same image, same time window. When the candidate changes the model output path, the old candidate's quality and post-update alignment evidence do not carry over. Crossed pairing does not guarantee a strong denominator: if the same config and workload show clear fast/slow tiers, keep per-round raw numbers and check tuning receipts, CPU/GPU state, Graph, cache and actual operator selection; an optimization is not effective just because BI=0 landed in the slow tier. Mark unexplained fluctuation's causal attribution `inconclusive`; any formal target point over budget by the predefined rule is still `fail`, with no cherry-picking rounds or switching the aggregation formula afterwards.
Profile only outside the formal profiler-free E2E measurement; GPU kernel durations may overlap and cannot be summed to stand in for E2E; a kernel name is not a module identity, see [Performance and profiling](performance.md).
### Failure and recovery
When the acceptance checker errors, first tell apart producer, numerics, scorer and parser. If the original worker/server receipts and score are sealed and the wrapper parsed them incorrectly after a schema change, keep the original artifacts and read them with a recovery checker that is **separately versioned with its hash recorded**; it must verify all workers, EngineCore, API server, package identity and workload, and keep the original failure status. It can correct only the control-plane conclusion; it cannot fill in missing values, unfinished requests or a failed service exit.
Keep failure records as they are and append new results; keep reruns and changed questions apart. The final report lists at least, separately: train-inference forward, module backward, real update, quality, performance, clean exit, final immutable artifacts; each cell is `pass`/`fail`/`inconclusive`/`not run`. Example scope: "five steps pass bitwise, GSM8K pairing passes, first-round c=32/64 performance over budget".
.agents/skills/train-inference-consistency/guides/rollout-engine.md
-108
View File
@@ -1,108 +0,0 @@
# Rollout engine and sampling
In RL the rollout engine (vLLM) pauses, sleeps and wakes next to the training process, its sampling randomness must be auditable, prefix cache and speculative decoding affect whether the same token sequence is bitwise identical, and RL-only features such as R3 constrain the engine config. Section 6 is a debugging procedure: save one rollout and replay the old-logprob forward. Measurement bases and A/B methods are in [Performance and profiling](performance.md); BI kernels and cudagraph are in [BI kernels](bi-kernels.md).
## 1. Sampling seed collision
**Mechanism.** Each vLLM DP engine resets the global RNG with the same `seed` at the end of warmup. Requests without a seed sample from the global stream, so the n-th request samples the same tokens on every engine. When GRPO spreads the siblings of one prompt across DP engines, many samples in a group repeat: before the fix, Nemotron had only 64-70 distinct responses out of 256 and up to 43% of sibling pairs identical; after the fix, 256/256. The fraction of zero-advantage groups barely changes with the fix; zero advantage mostly comes from task difficulty, not seed collision.
**Fix.** Reset the global RNG by `(seed, data_parallel_index)`; index 0 is bitwise identical to the old behavior (vllm-project/vllm#59788, merged). Use `data_parallel_index`, not `data_parallel_rank`: dense (non-MoE) models run independent engines under DP, and `reconfigure_for_independent_dp_rank()` resets `data_parallel_rank` to 0 (`vllm/v1/engine/core.py`), so 4 engines get the same seed stream and n=4 gets worse. Nemotron is MoE and goes through `DPEngineCoreProc` with a correct rank, so its 1024/1024 verification on verl holds, but only the index version fixes dense models.
**How to accept.** In the rollout dump, count distinct samples per group and the fraction of fully identical sibling pairs. With different seeds the baseline is about 0.1% identical sibling pairs, all short answers of <=8 tokens (the model outputs `\boxed{}` directly), and about 1-3% share a prefix of >=16 tokens. Anything above this is RNG correlation.
**Do not do this.** verl `rollout.full_determinism=True` gives the n responses in a group the same seed. Do not enable it.
**Greedy cannot verify the fix.** Under greedy, n repeats are identical (Nemotron: 32 of 32), and the score is a multiple of 1/number of problems (AIME with 30 problems: multiples of 1/30; 13/30 -> 19-22/30). A greedy A/B before and after the seed fix shows no difference because greedy does not sample. Rising entropy and the three config curves overlapping are visible only on sampling runs.
## 2. Prefix cache semantics
**Premise.** "Enabling PC does not change the numerics" requires the same request to be bitwise identical on hit, miss, and PC off. For pure attention models KV block reuse guarantees this. For stateful layers (Mamba2 conv/ssm state) it depends on which states the engine copies.
**The Mamba2 pitfall.** vLLM's Mamba align mode copies only the conv and ssm states. Under BI=1, Mamba2 uses exact replay (ported from vllm#55905, which stores three extra chunk buffers x/dt/B) and needs 5 states, so BI=1 Mamba2 originally rejected PC. The fix copies every state in each layer's declared state list; after that, hit/miss/PC off are bitwise identical. Early versions had a convolution window alignment bug in state export (written aligned to the tail, read from the head).
**Cost.** PC under BI=0 has a measurable cost: in Mamba align mode each decode step adds a batch of eager launches, and an 8K prefill is split into two segments. The approach in upstream #58851 is the best.
**Confounders between the two arms.** BI=1 Mamba2 rejects PC while the BI=0 R3 AuxOutput forces PC on, so PC on/off is a confounder in the BI=0 vs BI=1 throughput comparison. List confounders first, then compare; see [paired acceptance](recipe-and-acceptance.md#paired-acceptance).
**Relation to cudagraph.** A PC hit changes the token count of a step. Piecewise cudagraph chooses between the graph and eager fallback by total token count vs `max_cudagraph_capture_size`. Each path is batch-invariant on its own, but the two are not bitwise equal. Under BI, pin `FULL_AND_PIECEWISE`; see [BI kernels](bi-kernels.md).
**Checks for stateful layers and caches.** Mamba/SSM must also save chunk boundary state and partial input buffers; check aligned/unaligned split, cross-chunk and restore separately. Cover indexing and capacity of long-context KV/page/workspace near the boundaries.
When testing prefill, confirm whether the prefix cache was hit; during diagnosis, disable or isolate it if needed so the target path executes. Verify the production cache's hit, reuse, and preemption/recompute paths separately. Do not make a comparison pass by clearing the cache and then claim the original lifecycle is fixed.
For performance fixes on the cache path, use bitwise equality with "cache off" as the gold standard: hit positions outside chunk boundaries, every DP rank, recomputed hits after forced preemption, and chunks truncated by the scheduling budget; confirm in the trace that the target branch executed. Then add a hook that checks the device-side state index against the host prediction at every step, plus a per-problem quality evaluation. Two fixes that each pass alone may silently break when stacked; the combination must pass the gate on its own.
## 3. MTP / speculative decoding
- The Nemotron checkpoint has `num_nextn_predict_layers: 1`, and the MTP block is not quantized (bf16). It does not take part in training; the training side must explicitly skip its mapping ([training forward](training-forward.md)).
- Usable with BI=0. BI=1 + MTP requires the logprob from speculative verification to be bitwise identical to the full forward; without verifying this you cannot claim true on-policy.
- If used, use it in both arms; otherwise neither the throughput nor the logprob comparison holds.
## 4. R3 (routed experts replay) feature constraints
- **What is returned.** One `(num_moe_layers, top_k)` row of uint8 routed expert ids per token (Nemotron: 52 layers x 6 = 312 B), returned to the trainer with the request's CompletionOutput for route replay. 0 is a valid expert id; check distinctness, not non-zero.
- **PC must be on.** The R3 AuxOutput is keyed by KV block hash; with PC off there is no key. This is why the BI=0 arm in §2 is forced to enable PC; the keys-only form (PR 57968) does not lift it.
- **Why per-step return is expensive.** 312 B exceeds the 256 B msgpack zero-copy threshold, so each request's row per step becomes a separate frame, and the frontend accumulates millions of small objects per request before concatenating them at the end. The overhead is on the engine->frontend return path, not in the capture hook. The fix returns once at request finish (finish-only) and packs each step into one ndarray on worker->engine. `output_processor` already exposed `routed_experts` only at finish, so semantics are unchanged.
- **There is another floor below R3.** msgpack unpacking of three small ndarrays per token for sampling logprobs is a fixed frontend cost (about 25-30 µs/token), larger than everything else left after R3 is fixed. BI=1 does not see it when the engine is the bottleneck; BI=0 has a faster engine and hits it. Two arms tying on e2e is often due to this, not the kernel. Check the frontend asyncio busy fraction before comparing the arms ([Performance and profiling](performance.md)).
- **What does not change.** R3 adds no GPU->CPU sync, async scheduling is kept, capture writes into a statically preallocated buffer, and cudagraph is unaffected. flashinfer_cutedsl and Humming are both modular MoE (they receive the router's top-k), so moving R3 into the kernel saves nothing.
- **Limits of the equivalence check.** BI=1 cannot run the equivalence check (Mamba2 BI rejects PC while AuxOutput forces PC, until the §2 fix lands). BI=0 compares sha with strictly serial single-request greedy, but BI=0 routing is itself not batch-invariant, and an unpatched baseline also repeatedly shows partial inequality against the same reference. So first measure the drift with several baseline-vs-baseline runs; the patch counts as equivalent only if it fully equals one clean reference. Bitwise comparison of concurrent greedy dumps is not evidence.
## 5. Lifecycle and GPU memory ledger
**Three pause modes** (`EngineCore.pause_scheduler`): `abort` marks all running requests `FINISHED_ABORTED`, queues new requests, and optionally clears the cache; `wait` queues new requests but keeps stepping until running requests finish (not supported by the inproc engine); `keep` pauses everything, and the Future completes when the output queue is empty. In all three modes the Future returned by pause is the barrier; until it completes, workers still step. In `abort` mode the client receives aborted outputs and must resend them. Requests that arrive during pause are queued and released at resume, using the **new** weights. These are "orphan requests" (pause issued before the rollout tail finishes, or the agent loop sending during pause), and their tokens belong to no single policy version. Before each step's pause, confirm in-flight is 0; after resume, confirm the queue is empty. The HTTP endpoints `/sleep`, `/wake_up`, `/is_sleeping` and pause/resume are dev endpoints and need `VLLM_SERVER_DEV_MODE=1`, otherwise 404.
**Three sleep levels** (`EngineCore.sleep`, `Worker.sleep`):
| level | What it does | GPU memory | Host memory |
| --- | --- | --- | --- |
| 0 | Only pauses scheduling | Unchanged | Unchanged |
| 1 | Offloads weights to CPU, drops the KV cache, clears the prefix cache | Frees KV and weights | One weight backup per rank (DS4.1 78.61 GiB); the pinned cache is not returned after wake, see [Checkpoint and resume](checkpoint-resume.md) |
| 2 | Discards all GPU memory, saves only `named_buffers` to CPU | All freed | Small |
After L2, `wake(weights)` only reallocates GPU memory; the contents are uninitialized. "Reserved is very high but still no GPU memory" after wake is an allocator boundary problem: PyTorch's reserved-but-unallocated blocks cannot be used by another allocation path such as CuMem; see [Checkpoint and resume](checkpoint-resume.md#3-reserved-is-high-so-why-can-wake-still-have-no-gpu-memory).
**Set gpu_memory_utilization by training mode.** The training process does not fully release GPU memory during rollout. Write the value into the recipe per training mode; check the resolved value inside the container and the KV budget in the vLLM log, not the outer setting. Measured values and rules per mode are in [Checkpoint and resume](checkpoint-resume.md#5-gpu-memory-set-gpu_memory_utilization-per-mode-when-training-and-rollout-are-colocated).
**Metrics are cumulative.** `vllm:*_total` counts since process start; when one engine runs several phases, always use `after − before`. Why the first phase happens to be right and how to cross-check are in [Performance and profiling](performance.md#metrics-are-cumulative-values).
**DP traffic lands on rank 0.** When the Rust frontend sends one request at a time, all of them land on rank 0 (engine selection is `min_by_key(routing_score)`, which takes the first on equal load). "No events on rank 1" then means no traffic, not broken routing. Use `X-data-parallel-rank` to pin requests to each rank in turn. The two frontends interpret this header differently: Python in hybrid mode validates it as a local index in `[0, dp_local_size)`, Rust uses the global rank; do not send it to a single-rank service. Python hybrid/external LB with a dense model does not start (`DPAsyncMPClient` asserts `stats_update_address`, upstream #32252); test Python hybrid with an MoE model.
---
## 6. Save one rollout and replay the old-logprob forward
### When to use
For a difference that shows up once across many RL steps, save the failing input batch and matching policy early. Rerunning the rollout changes tokens, scheduling, lengths, cache, packing and later training data at once, hiding whether an intervention worked.
1. Save the framework batch and the policy identity after the real rollout ends and before the old-logprob forward; append the trainer scores after the comparison.
2. Verify that the save hook executed, that sequence / valid token counts match the log, and that the data reads back.
3. Use the framework's existing training-only/repeat entry point to read the same batch; restore its weights.
4. Return after old-logprob; forbid the optimizer step and regeneration; record the skipped stages and the weight fingerprints before and after.
5. Run A/A first, then change one layout factor. If A/A itself is unstable, first check randomness, mutable buffers, async streams and runtime state.
### Minimal matrix
| Variant | Change | Answers |
| --- | --- | --- |
| A | Original packing | Can the frozen input reproduce it |
| B | Exact repeat of A | Is it run-to-run nondeterministic |
| C | Sequence permutation/packing | Does it depend on layout |
| D | micro-batch size or target position | Does it depend on call shape/neighbors |
| E | Separate process or equivalent node | Does it depend on initialization/device state |
Run first the variants that distinguish the current hypotheses; the full matrix is not needed every time. After reordering, restore and compare by UID and response token position; the physical row no longer represents the same request. C failing does not prove the packing algorithm wrong; it may trigger a CP, page, tile or expert boundary error.
### Keep the trigger condition when shrinking the input
Keep the target sample and its necessary neighbors first, then reduce length and request count. After every reduction, rerun the criterion to confirm the failure remains. A CP problem may depend on the remainder from the end of the previous request to the start of the rank; a Graph problem may depend on short fillers and the capture bucket. Keeping only the failing request often makes the problem disappear.
Freezing the rollout and the weights does not freeze all runtime state. Training-mode dropout, FP8 amax history, RNG, route replay, caches and workspace must each be kept or reset according to the failure condition. A fresh eval forward is not a substitute for the original old-logprob path.
A resident model saves checkpoint loading, but state must not leak across variants. First verify A/A between a new process and a reused process; when native libraries, autotune or compile config are involved, restart the corresponding workers.
### Leaving the replay stage
Once the original reproducer is closed, return to normal generation, real backward/update and the next rollout step. Training-only success supports only the frozen-policy forward conclusion. If the error appears in backward or in long-running state, save that state and build a minimal reproducer for that stage.
This rollout replay is a debugging procedure. Nemotron's SSD exact replay is an algorithm in the model implementation that recomputes the partial chunk to match prefill. Do not call them by the same name.
.agents/skills/train-inference-consistency/guides/training-forward.md
-412
View File
@@ -1,412 +0,0 @@
# Training-side model support
How to build a training implementation cheaply and prove it computes the same math as inference. Where training weights come from (master -> training forward weights) is in [Precision, quantization and QAT](precision-quant.md).
## Three routes
The routes differ in where parallelism and kernels come from. In all three we write the model.
| | MCore (Megatron-Bridge) | FSDP2 (VERL's HF interface path) | mLite |
| --- | --- | --- | --- |
| Positioning | Mature TP/PP/EP/CP, main production RL path | Cheapest: write only the model, no parallelism | Built for consistency: forward calls vLLM kernels, backward reuses the training implementation |
| Forward source | MCore kernels (TE, grouped GEMM) | Our PyTorch model to vLLM semantics, in the Transformers `PreTrainedModel` interface for the VERL loader and PEFT | vLLM kernels (`visible_forward(visible, native, ...)`) |
| Achievable consistency | BI=0, via R3 + importance sampling | BI=0 | BI=1, logprob raw bytes identical |
| What to write for a new model | Bridge weight mapping + modules MCore lacks (Mamba/SSM, special norms) | Whole model: config class, checkpoint loading (with dequantization), per-layer forward, per-operator precision contract; autograd backward, filled in where missing | One `model/<name>/`: config, checkpoint mapping, assembly, protocol; modules from `primitive/modules` |
| Typical pitfalls | Mapping checks values, not semantics; MCore and vLLM math differ (MoE activation clamp) | bf16 autocast around the VERL forward rewrites module precision; `--system-site-packages` ABI; PyTorch reference (Mamba scan) too slow; loading a third-party implementation from the image by mistake | Private modules inside the model directory, bypassing the shared skeleton |
Choose after the [module inventory](#module-inventory-and-placement-rules), by: is BI=1 needed, how many backwards are missing, what parallelism is needed. Nemotron looked cheap by name; FA4, Mamba scan and W4A16 MoE all lacked a backward.
## Find by symptom
| Symptom / question | Where to go |
| --- | --- |
| New model loads and config parses, but unclear whether it can train | [Module inventory](#module-inventory-and-placement-rules) |
| Want to reuse vLLM forward code, but tangled in paged KV / cudagraph / scheduling | [Call kernels only](#reproducing-the-forward-call-kernels-only-do-not-port-the-scheduler) |
| Weight mapping shapes are right, output silently degrades | [Mapping checks values and semantics](#43-mapping-checks-cover-values-and-semantics-not-only-shape) |
| Inference kernel has no backward; should we write one | [Three-case decision](backward-validation.md#three-cases) |
| Wrote a VJP; how to prove it is right | [Module-level comparison steps](backward-validation.md#comparison-steps) |
| Unit tests pass but real training fails / step 0 occasionally not bitwise | [Two test-infrastructure false negatives](backward-validation.md#two-test-infrastructure-false-negatives); [BI kernels](bi-kernels.md) |
| Changes after modifying packing or CP/TP/EP; wrong only on long sequences | [Layout and length](#distributed-layout-length-boundaries-and-training-side-acceptance) |
| `ppo_kl` non-zero after rollout switches parallel layout (e.g. DP+EP to TP1 per GPU) | [Bitwise identity after changing the parallel layout](#bitwise-identity-after-changing-the-parallel-layout) |
| `modified by an inplace operation` after no longer recomputing old-logprob | [Entropy when old-logprob is not recomputed](#entropy-when-old-logprob-is-not-recomputed-bypass) |
| How to hook into VERL / Megatron | [Framework integration](#verl--megatron-integration) |
| Which kernel an operator uses in BI mode and its cost | [BI kernels](bi-kernels.md) |
## Key rules
1. Semantics follow the vLLM implementation and the checkpoint.
2. Before starting, list for every module whether its inference kernel has a backward.
3. Call kernels only; do not port the scheduler. mLite (BI=1): FP32 logprob raw bytes vs vLLM-captured trajectories, target 0 diff. FSDP2 / MCore (BI=0): no 0-diff target; judge each side against an independent reference, then look at precision.
4. Mapping checks cover values and semantics, not only shape; fused parameters are overwritten whole.
5. The precision contract lives in the model. Where vLLM computes in FP32 (position-encoding angles, router logits and bias, norm statistics, etc.), the training module sets the dtype and disables the outer autocast. Under VERL's bf16 autocast, casting only the inputs of matmul-type operators (including `F.linear` and `einsum` decomposed into bmm) does not stop the downgrade. Measure output dtype and values under the real autocast.
6. Backward follows the formula, not the forward kernel. Only rounding differs -> reuse. Quantization rounding -> fix the STE contract, reuse BF16 GEMM. No whole-segment backward -> write glue, GEMM still TE. "No backward kernel at the same precision" is not a reason to rewrite.
7. Reuse model-agnostic operators: mLite `primitive/modules` and `ops` first, then MCore/TE/mamba_ssm; only then is it new code.
8. Backward need not be bitwise; report bitwise and tolerance separately; the reference builds its graph independently with ATen autograd; tolerances are fixed before seeing results; synthetic inputs are labeled as not real activations.
9. Do not recompute activations; use the forward's saved intermediates (DS4: gate_up / LSE).
10. The model directory holds only config / checkpoint / model / protocol; reusable parts go into `primitive/`; no same-named `primitive/` under the model directory; impl name `lite`.
11. Long-sequence boundary points (KV page, packing prepend, the 7000/9216 tiers) need dedicated tests.
12. MTP block, EPLB redundant experts and draft are explicitly skipped in mapping and sync.
## Module inventory and placement rules
List per module: inference kernel, same-math training implementation, quantization rounding, discrete selection, backward. Then pick the route and each module's place in the code tree. This catches missing backwards before they cost weeks.
Training support means an implementation we maintain that backpropagates correctly. Config registration, same-family fallback, the vLLM inference implementation and third-party implementations in the image (newer Transformers modeling, NeMo AutoModel) do not count. Write it per [Reproducing the forward](#reproducing-the-forward-call-kernels-only-do-not-port-the-scheduler) into real source files and verify by loading through the actual loader in a new process. A REPL-only implementation or runtime monkeypatch is not evidence.
### 1. How to build the inventory
One row per module, every column filled; write unknown, not "should exist".
| Column | Question to answer |
| --- | --- |
| Inference kernel | The function vLLM actually calls (not the model-file class name), with quantization variants, attention backend, cudagraph branch |
| Same-math implementation in the training framework | mLite `primitive/modules`, `ops`; MCore / TE; mamba_ssm, causal_conv1d. Rounding-only differences count as same math; a different activation clamp does not. Third-party same-model implementations are not entered |
| Quantization rounding | FP8/NVFP4 rounding in the forward; weights and activations separately |
| Discrete selection | top-k expert IDs, sparse indices; IDs fixed and not differentiated, the continuous score needs a derivative |
| backward | Does the library have one; do its intermediates (LSE, rstd, token permutation) match what our forward saves |
| Conclusion | Reuse / glue / new code ([decision rules](backward-validation.md#three-cases)) |
Nemotron-H:
| Module | Inference kernel | Same-math training implementation | Rounding | Discrete | backward | Conclusion |
| --- | --- | --- | --- | --- | --- | --- |
| Mamba SSD scan | `mamba_chunk_scan_combined_varlen` | mamba_ssm | None | None | mamba_ssm `_mamba_chunk_scan_combined_bwd` | Reuse |
| conv1d | `causal_conv1d_fn` | causal_conv1d | None | None | Yes; dW needs fixed-order FP32 | Reuse + small glue |
| Attention (FP8 KV) | FA4 fixed split + FP8 Q/K/V | FA4 BF16 | Yes | None | FA4 BF16 backward only | Dequantize, recompute BF16 O/LSE, reuse |
| Mamba in/out_proj (FP8), shared expert (NVFP4) | FlashInfer / NVFP4 GEMM | TE `general_gemm` | Yes | None | TE | Fix STE, reuse |
| Routed experts (NVFP4 W4A16) | CuTe-DSL / Humming MoE | TE grouped GEMM | Yes | Yes | No whole-segment backward | Glue `routed_vjp` |
| RMSNorm / gated RMSNorm, routing score | BI kernel / grouped top-k | PyTorch formula | None | None for the score | autograd | Reuse |
| MTP block, draft | Not part of training | - | - | - | - | See §4 |
### 3. Placement rules
#### 3.1 mLite (Megatron `experimental/lite/megatron/lite`)
Read `experimental/lite/docs/porting.md` and `architecture.md` first. Three layers: `runtime` (lifecycle, training-step orchestration), `model` (registration + model-specific build/load/export protocol), `primitive` (reusable: parallel state, TP layers, checkpoint conversion, MoE utilities, optimizer wrappers). The runtime knows no model details: it gets the protocol from the registry, builds a typed config, and lets the protocol build the model.
| Rule | Content |
| --- | --- |
| The model directory holds only four things | `model/<name>/`: config, checkpoint mapping (WeightSpec), assembly, protocol. Impl name `lite` (`<name>/lite/`) |
| Modules go into primitive | norm / linear / router / experts / dispatcher / attention -> `primitive/modules`; generic VJP bridge -> `primitive/ops`; quantization encoding -> `primitive/quantization`; PP splitting, thd, CP -> `primitive/parallel`; HF read/write via `primitive/ckpt/hf_weights` (`SafeTensorReader`, `load_hf_weights`, `export_hf_weights`) |
| Model-specific math still goes into primitive and is registered | e.g. Mamba2 bridge, FP8-KV attention -> `primitive/modules/<x>.py`; no same-named `primitive/` under the model directory |
| No file for reused backward | Reusing TE / mamba_ssm backward means no `backward.py`; glue goes in the module file |
| Dependencies are one-directional | `model/` imports `primitive/`, never the reverse; top-level checkpoint/quantization does not import the impl subtree |
| Layering guards must be registered | Add the model to `tests/unit/runtime/test_layering_contracts.py` or the guard does not apply |
| No hybrid architecture packages | porting.md: "Keep these out of the first PR: Hybrid model implementation packages"; hybrid (Mamba + attention + MoE) modules are split into primitive |
| State third-party kernel dependencies | Document hard dependencies on vllm, flashinfer private APIs, mamba_ssm, causal_conv1d, and the `init_batch_invariance()` call in build_model |
DS4 is the positive example: `model.py` (650 lines) assembles te.RMSNorm, CSA, HCA/mHC, thd and the fp8 recipe with three model-specific deviations; `moe.py` (102 lines) = SigmoidTopKRouter + Experts + TokenDispatcher + SwiGLUMLP, hash-route the only model-specific code; `checkpoint.py` is one WeightSpec using shared load/export.
#### 3.2 MCore / Megatron-Bridge
- Bridge `models/<name>/`: provider + HF<->MCore mapping; modules MCore lacks (Mamba/SSM variants, special norms, HC) go into `megatron/core`, not Bridge.
- Mapping checks values, not only shape ([§4.3](#43-mapping-checks-cover-values-and-semantics-not-only-shape)).
- MCore and vLLM may use different formulas (MoE activation clamp); BI=0 cannot align those either, so list them first.
- The Megatron fork works only with Megatron-Bridge v0.6.2, pinned in the image.
#### 3.3 FSDP2 (VERL's HF interface path)
- Write a PyTorch model to vLLM semantics in the Transformers `PreTrainedModel` / config interface so the VERL loader, PEFT and HF save/load use it directly. "HF" is only the interface and weight format.
- Delivery: the model goes into the verl trunk, registered to the Auto classes (`AutoConfig` / `AutoModelForCausalLM`, etc.), and ships with the recipe as one verl PR ([Code and delivery §8](../practices/code-and-delivery.md#verl-a-thin-recipe-not-a-second-training-framework)). The verl FSDP extension guide's advice against copying modeling into `verl/models` targets models already in transformers. Register at the current checkout's FSDP load chain (`get_hf_auto_model_class` -> `from_pretrained`).
- Develop and deliver from the same verl source and registration entry. Do not modify the checkpoint's `config.json` / `auto_map` (the low-precision checkpoint is the single source of truth) or enable `trust_remote_code`. The load log must show our registered class was used.
- VERL's automodel engine (`automodel_actor`) uses NeMo AutoModel and is not this route; routes are distinguished by forward source, not by FSDP2 sharding.
- Under bf16 autocast, FP32 computations (position-encoding angles, router logits and bias, norm statistics, following vLLM) disable autocast inside the module and set the dtype.
- Backward prefers autograd; add operators next to the model file only when a backward is missing or too slow (the 160-line PyTorch Mamba scan reference).
- Common pitfalls: `--system-site-packages` ABI, the actual `use_remove_padding`, PEFT targets not going through nn.Linear.
- Under EP4 + FSDP2, all ranks' DCP expert shards share key and offset, so restore reads another rank's experts; the key must carry the EP rank ([checkpoint guide](checkpoint-resume.md)).
### 4. Parts that do not take part in training
Each gets its own inventory row:
| Part | Mapping | Facts |
| --- | --- | --- |
| MTP block | Not loaded for training; WeightSpec lists it as explicit skip, not missing-key tolerance; `_ignore_unexpected_suffixes` etc. only for listed suffixes | Nemotron `num_nextn_predict_layers: 1`, MTP unquantized bf16; usable with BI=0, then in both arms; BI=1 + MTP requires speculative-verification logprob bitwise equal to the full-segment forward |
| EPLB redundant experts | Training has logical experts only; physical-to-logical mapping is inference-side | Completeness counts logical experts, not physical tensors |
| draft / speculative model | Independent checkpoint and mapping | Shared weights (embedding, lm_head) state their owner |
Skipped parts appear in the load log and manifest as "explicitly skipped", not "no error reported".
### 5. Example: reuse audit of Nemotron-H against mLite
Nemotron-H is 3824 lines. From `primitive/` it reuses only protocol-layer orchestration utilities, `load_hf_config_dict` and `vocab_parallel_entropy`; every modeling module is rewritten under `model/nemotron_h/vllm/primitive/`.
| File (`N = model/nemotron_h`) | Lines | What it does | Verdict |
| --- | --- | --- | --- |
| `N/config.py` | 155 | HF config dataclass, validation | Keep (model dir) |
| `N/quantization.py` | 207 | NVFP4/FP8 requantize, E2M1 grid, FP8 encoding, hand-read index.json | Numerics -> `primitive/quantization/nvfp4.py` (parameterize E2M1 RNE vs mxfp4 ties-down); IO via `SafeTensorReader` |
| `N/checkpoint.py` | 486 | HF name mapping, Export spec, BF16 master invertibility check, θ0 threshold, own `load_hf_weights` | Ordinary tensors via `load_hf_weights` + WeightSpec; keep master check and θ0 threshold; importing the impl subtree is a layering inversion |
| `N/vllm/model.py` | 188 | Layer/Model assembly, PP stage | Keep; PP split via `build_pipeline_chunk_layout` |
| `N/vllm/protocol.py` | 409 | protocol, contract validation, BI check, FSDP2/dist_opt, post-step requant | Keep; hard-coded 52-layer pattern and (2688,128,1856,6) make it single-checkpoint |
| `N/vllm/primitive/dense.py` | 647 | visible forward + BF16-master VJP bridge; RMSNorm VJP; Humming/CuTe NVFP4 linear; projection factory; FP8/NVFP4 TrainingLinear | Bridge and RMSNorm VJP -> `primitive/ops/aligned_vjp.py`; `native_linear_vjp` duplicates `_te_general_gemm`; serving kernel adaptation stays |
| `N/vllm/primitive/logprob.py` | 36 | Chunked lm_head + vLLM `compute_token_logprobs` + FP32 log_softmax STE | Keep; the docstring's DS4 function does not exist |
| `N/vllm/primitive/attention/*` | 231 | `Fa4Fp8KVAttention` (FlashAttnFixedSplit 6768-token page, FP8 Q/KV), `_Fp8AttentionVJP` | Model-specific; reusable once GQAttention supports `rotary=None, qk_norm=False` |
| `N/vllm/primitive/mamba/*` | 316 | `_PackedConvVJP`, `_PackedScanVJP`, `SSMMeta`, `MambaMixer` | Keep both VJPs; generic Mamba2 -> `primitive/modules/mamba2.py`; derive `SSMMeta` from `PackedSeqParams`/`ThdPackMeta` |
| `N/vllm/primitive/moe/module.py` | 185 | `_FixedRouteVJP` + Router, `_CombineVJP`, ReLU² SharedExperts | Router VJP equals `SigmoidTopKRouter`, reusable, not hooked to router_replay; add an activation argument to `mlp.py` for SharedExperts |
| `N/vllm/primitive/moe/grouped.py` | 705 | FlashInfer W4A16 signature/sha pinning, `CuteDslRoutedExperts`, `Nvfp4ExpertWeights`, `_te_grouped_gemm` rewrite, ReLU² `routed_vjp` | Keep W4A16 forward + grouped VJP; TE grouped GEMM wrapper duplicated; `ep4_routed_experts` hard-codes 4 ranks |
| `N/vllm/primitive/moe/communication.py` | 251 | DeepEP normal-mode EP, `reduce_ep4_parts` (FlashInfer one-sided combine order), `EPRoutedExpertsVJP` | Use `_build_deepep_buffer`; keep serving-order reduction, parameterize EP size/topk |
Six placement violations: impl name `vllm` not `lite`; hybrid architecture package; same-named `primitive/` in the model layer; circular imports between `checkpoint.py`/`quantization.py` and the impl subtree; `nemotron_h` missing from the layering guard; third-party kernel dependencies and `init_batch_invariance()` undocumented.
Should have been reused: safetensors index reading (5 copies); HF load of ordinary tensors; DeepEP buffer; TE GEMM wrapper (2 copies); visible forward + reference VJP bridge; Router (`SigmoidTopKRouter` + `group_limited_topk`, forward switched to vLLM grouped_topk, hooked to router_replay); shared expert; E2M1 grid and FP8 encoding; PP splitting; `SSMMeta`; two docstrings citing nonexistent DS4 functions.
Model-specific, keep: the two Mamba2 VJPs + packed conv/scan; `Fa4Fp8KVAttention` + FP8 attention VJP; W4A16 MoE forward + ReLU² grouped VJP; serving-order reduction (parameterized); Humming/CuTe NVFP4 and FP8 serving linear; vLLM forward-value logprob; BF16-master invertibility check and θ0 threshold; ModelOpt 4over6 requantize rule.
## Reproducing the forward: call kernels only, do not port the scheduler
The first deliverable after the inventory is a training forward with vLLM's math. Logprob measurement basis and R3 route replay: [workflow](../workflow/README.md), [rollout engine](rollout-engine.md).
### 1. First separate the responsibilities of the model implementation and the training framework
A training framework (FSDP2, MCore, mLite runtime) shards a trainable model; it does not make an inference-only kernel differentiable or write the model. Write a minimal PyTorch model for the current loader interface covering the config, checkpoint loading, forward, gradients and save interface actually used; full Transformers integration, generation and all parallelism are not needed first. The semantic basis is vLLM and the checkpoint, not a third-party implementation in the image.
| Capability | What the training route still has to do | What in inference / Megatron is not copied |
| --- | --- | --- |
| Model math | Real semantics of attention/SSM, norm, position encoding, routing, residuals, special embeddings | vLLM paged KV, decode scheduling, cudagraph management; MCore provider/ModuleSpec wiring |
| Backpropagation | Input gradients of every training-path operator; trained-parameter gradients; recompute correctness ([Backward validation](backward-validation.md)) | TP/SP/CP-specific collectives and gradient reduction |
| Weights | Names, axes, fused order, expert IDs, sharing, derived state | HF<->Megatron shard layout, PP export protocol |
| LoRA | Target coverage, A/B direction and scaling, non-zero adapter taking effect | Parallel linear wrapper, TP adapter aggregation |
| Memory | Initialization peak, large frozen table residency, save/export peak | Megatron memory saver, PP scheduling |
dtype, dequantization, QAT and update format: [precision guide](precision-quant.md). Here each module states its input/output precision and differentiable boundary.
### 2. How to wire kernel-only calls
Each layer's forward runs once: call the vLLM kernel for the visible output, bitwise identical to inference, and save what the backward needs; see [backward wiring](backward-validation.md#how-to-wire-it-forward-once-backward-uses-what-the-forward-saved). Borrow only the kernel, no vLLM runtime object:
- No KV paging, decode scheduling or CUDA Graph management. Training runs a full-sequence teacher-forced forward; packing passes boundaries via thd/cu_seqlens.
- Reuse order: model-agnostic differentiable operators -> concise PyTorch reference to vLLM semantics -> ready-made operators with backward -> custom autograd/kernel only if necessary. This order only says where to look first; the training-step cost of each choice is checked at C6 against the C1 budget, not deferred to C10.
- Handle discrete top-k / index selection apart from the selected differentiable tensors. What is trained, frozen or given an auxiliary loss follows the model definition; do not detach to get it running or add router/indexer objectives.
- vLLM layer modules used directly (FusedMoE, ModelOpt linear) need `get_tp/ep_group`, so the training process creates vLLM process groups (Nemotron does; DS4's own modules do not).
- Take the same branch as inference: reproduce the quantization config, attention backend and BI-switch kernel selection, or the comparison is against a different path. Tests use the production `CompilationConfig(custom_ops=...)` ([the two false negatives](backward-validation.md#two-test-infrastructure-false-negatives)).
"The same kernel" still has pitfalls (Nemotron 7-layer localization):
1. Direct and compiled paths of the same gated RMS kernel differ; calling it through an opaque custom-op boundary removed the difference.
2. A differing reduction block was wrongly blamed alone; the same captured kernel showed both blocks match.
3. The last 4.768e-7 logprob difference came from the HF forward calling `.float()` after lm_head; keeping the real BF16 head output and sharing the scoring fixed it. This is not truncating an FP32 projection to BF16 to manufacture equality.
### 3. Read the inference implementation along the full support chain
Order: model file (layer types and order, residual and norm positions, position encoding, attention visibility, router/indexer, special embeddings) -> each layer's actual kernel and quantization variant -> attention backend and KV dtype -> weight loader (name mapping, fused order, derived scale) -> BI-switch branches. Confirm the parameter/buffer list from the same checkpoint/config.
Sizes sharing one implementation need no separate port. For ordinary dense models, look for a reusable generic skeleton first; we still write the model. Multimodal is validated per task; diffusion is out of scope.
### 4. The minimal implementation that really has to be filled in
#### 4.1 Reproduce the math
Confirm from the checkpoint/config and inference code: layer types and order, parameters/buffers, residual and norm positions, position encoding, attention visibility, router/indexer, special embeddings. Model-specific math (`sqrt(softplus)` scoring, SwiGLU with clamp, mHC Sinkhorn, ReLU² experts) is in no library; write it from the inference source, not the paper.
#### 4.2 A frozen base still has to propagate input gradients
```python
W.requires_grad_(False)
y = linear(x, W) + scale * B(A(x)) # no gradient for W, but still propagate to x
loss(y).backward()
```
Wrapping a frozen block in `no_grad()` cuts LoRA gradients of earlier layers; only a frozen table lookup needing no input/parameter derivative is exempt. With initial B=0 a zero gradient for A can be normal; test with a non-zero adapter or a real step. A finite loss does not prove backward.
<a id="33-shape"></a>
#### 4.3 Mapping checks cover values and semantics, not only shape
Per weight class: source name -> training name -> serving name, axis transforms, sharing. Focus on fused QKV, gate/up order, expert numbering, extra norm/bias, tie weights, non-persistent buffers, derived state (FP8 KV scale, NVFP4 alpha). Every missing/extra item needs a reason; quantization-derived items follow the precision guide and need not be byte-identical to the low-precision source. FSDP needs no HF<->Megatron bridge but may need checkpoint->training and export mappings.
#### 4.4 Packing boundaries for Mamba / stateful modules
Every stateful operator (conv1d, SSM scan, recurrence, PLE shift) must consume sequence boundaries; a correct attention mask proves nothing about them. Incidents:
- verl via Megatron-Bridge did not pass boundaries to Mamba; sequence states mixed and logprob was silently wrong, visible only as a larger mismatch.
- Qwen3.8 Flash-Next's PLE shift made the second sample reference the first sample's last token.
Check: unequal-length A/B run separately and merged; sentinel-initialized outputs to catch missed writes. In mLite, derive boundary metadata from `PackedSeqParams`/`ThdPackMeta`, not a separate `SSMMeta`. See also [Long-sequence boundary points](#long-sequence-boundary-points).
#### 4.5 Recompute and small-model trimming
With activation checkpointing on, verify recompute reads the same state and gives reasonable gradients. Trim to the smallest layer set covering every non-equivalent path and its dependencies, usually 4-6 layers; keep mixed attention, cross-layer sharing and layer-specific table lookups. Record kept layer indices and the config transform; a slice proves only covered paths, not full-model mismatch or score gain.
### 5. How to prove it computes the same math
Feed the training forward vLLM-captured trajectories (fixed tokens, weight version, position, mask, routing); never compare probability summaries or argmax. mLite (BI=1): FP32 logprob raw bytes, target 0 diff. FSDP2 / MCore (BI=0): judge each side against an independent reference, then check the remaining difference comes only from precision rules.
- Same policy weights. Differing dtypes are a contract mismatch; do not convert to fake bitwise.
- Reference artifact (Nemotron): frozen trainer image, full actor, 2048 tokens, FP32 0 diff; five real updates, 883,998 tokens of identical FP32 logprob.
- Old-logprob recomputation and the grad-enabled forward are two paths; evidence for each ([dual-path contract](../workflow/validation.md#old-logprob-and-training-forward-dual-path-contract)).
- Find the first divergence per [debugging](debugging-mismatch.md): per-module dump, single-request prefill, then the operator.
### 6. Small tasks handed to agents and stop conditions
Each outputs a short conclusion, source version, reproducible command and artifact.
| Task | Minimal acceptance | If it fails |
| --- | --- | --- |
| Implementation inventory | Training model class, operator backward, checkpoint loader; gaps listed | Write the minimal model; stop cold-starting the large model |
| Module comparison | Same weights/inputs; gradients vs a differentiable reference | Locate the earliest difference |
| Few-layer training | Non-zero loss, finite gradients, intended adapter changes | Separate broken path, operator bug, gradient-less target |
| Save and restore | Fresh-process restore: consistent forward and next step | Check state, adapter, optimizer, residency |
Distribute only after the first two pass; scale up after the few-layer two-step loop passes.
### 7. How to judge development cost
Training implementation exists: wire loading, LoRA, sharding. Inference only with differentiable references: fill in assembly, mapping, interfaces. Key operator lacks backward or clear semantics: separate blocker; reference and gradient validation first; no uniform person-day figure.
## Distributed layout, length boundaries and training-side acceptance
Record trainer and rollout topologies separately. TP/CP/EP, padding, packing and collectives change data mapping and reduction paths; long sequences hit untested boundaries; forward 0 diff does not vouch for gradients or the optimizer.
### Compute resources first, then request them
Write trainer and rollout process groups separately. At four GPUs per node, an inference DP8/EP8 group spans two nodes; two DP4/EP4 replicas are not one EP8 group and skip cross-node EP communication. The runtime decides grouping.
If dense DP excludes CP: `W = TP × PP × CP × DP`. Expert side: verify `W = ETP × EP × PP × EDP` and the CP mapping separately. **Do not multiply EP into the dense world again.** W=8, TP1/PP2/CP2 gives DP=2; a log calling DP×CP "DP" shows 4. Clarify naming and members; do not change GPU count to match a config string.
GPU memory follows the layout. With PP1 each rank holds all layers' activations, linear in sequence length. Per-route EP backward state stored for other ranks' slots grows by the average ranks per token; store only own routes. Rollout `gpu_memory_utilization` is set from what training still holds during rollout, per training mode, in the recipe ([Memory budget 02 §5](checkpoint-resume.md#5-gpu-memory-set-gpu_memory_utilization-per-mode-when-training-and-rollout-are-colocated)).
### Packing and CP
For the target UID reconstruct global request `[seq_start, seq_end)`, rank range `[rank_start, rank_end)`, local rows, cu_seqlens, positions, all compression/reordering mappings.
The intersection `[max(seq_start, rank_start), min(seq_end, rank_end))` contributes rows only when start < end. The overlap window must not give a disjoint request this rank's output. Host cumulative boundaries equal the rows the device compactor produces.
Cover: before/after/at the boundary, empty intersection, cross-rank, compression-ratio and tile remainders. Window constant 8: check 1…7, not only multiples. Start from a hand-computable counterexample, then enumerate nearby legal configs.
Passing `cu_seqlens` does not mean a fallback consumes it ([§4.4](#44-packing-boundaries-for-mamba--stateful-modules)).
### EP and collectives
Record expert mapping, tokens per rank, padding capacity, route order, shared/routed combine, quantization scales. Empty ranks may still join collectives; partially written reused buffers leave stale data. Use different sentinels per side; zero-filling both gives a false green.
Pinning NCCL algorithm/channels may remove reduction differences on one stack and cut throughput significantly; verify it per topology, not as a universal env list, before relaxing config or building a deterministic collective.
MCore dev's grad-ready assertion takes first-batch hook counts as golden and errors on a different count. Short-sequence micro-batches can make counts data-dependent (Nemotron r10 at step 2; r12 did not reproduce). With dp×cp=1 and expert_dp=1, `overlap_grad_reduce` off avoids it without changing values; root cause unknown. With EP4 + dist_opt and dp>1, multi-micro-batch reduction fires early (reported).
### Bitwise identity after changing the parallel layout
**Principle.** Floating-point addition is not associative. Layout sets reduction grouping, order and rounding: TP splits a matmul into per-rank partial sums; EP computes a token's top-k experts on different ranks and merges. Last-bit differences become different tokens after dozens of layers and sampling, so bitwise identity holds only for the verified layout; any TP/EP/DP change must be proven again.
**Rules.**
1. The reduction tree is part of the numerical contract: per level, accumulation dtype, order and rounding points. All legal rollout layouts and training walk the same tree.
2. A degenerate layout (TP=1, EP=1) must simulate the tree: per-rank partial sums, same rounding, same merge order.
3. After a layout change, compare per-token logprob raw bytes on fixed inputs before training. Non-zero `ppo_kl`/probs_diff is fixed, not avoided by reverting the layout.
4. Locate: dump per layer, find the first differing layer, confirm its inputs (MoE: router logits too) are bitwise equal; the bug is that layer's reduction ([Fault localization](debugging-mismatch.md#bisect-and-first-divergence)).
5. Measure the cost of simulating the order and report it apart from the layout gain.
**Examples.**
| Model | Layout change | Where the difference is | Fix and result |
| --- | --- | --- | --- |
| Nemotron-H (W4A16 MoE, BI=1) | rollout from DP4EP4 per node to TP1 per GPU | EP4: per-rank FP32 sums rounded to BF16, merged in FP32 as `((s0+s1)+(s2+s3))+(s4+s5)`. EP1: six outputs summed sequentially in FP32, no intermediate rounding. Max logprob diff 0.4; greedy tokens diverge at positions 10–129 | Without EP, BI mode computes all experts in one launch and reduces in EP4 order (vLLM `661cde6833`, not in the image). Both layouts match training bitwise; `ppo_kl` 0 over 58 steps; multi-turn replay 404 s vs 485 s, reduction cost about 2.5% |
| DeepSeek V4.1 | training TP=1, inference TP>1 | order of per-TP-rank partial sums | both sides use one fixed reduction tree, TP=1 included (`VLLM_TP_INVARIANT_LEAVES`) |
### Asynchronous communication and lifetime
Check streams, events, buffer owners and reuse timing. Passing with a global synchronize only proves a timing issue and is not the fix; find the missing dependency and accept on the same path.
### Long-sequence boundary points
Test at real lengths:
| Boundary | Fact | How to test |
| --- | --- | --- |
| KV page | training attention replicates vLLM's KV page size; page boundaries are numerical boundaries | sparse dY across the page boundary |
| Fixed split attention | fixed split count and keys per split; the resulting `MAX_SEQ_LEN` is an invariance bound, not a performance knob | above: error, not silent truncation; below: two real lengths × random/zero/sparse dY |
| Packing prepend | one prepended short sequence shifts every target row offset | two prepend controls vs none, bitwise |
| Micro-batch fits the longest sequence | verl needs each micro-batch to fit the longest sequence (`ppo_max_token_len_per_gpu`, Nemotron 16384 vs DS4 4096; Nemotron training has no CP) | OOM probe at formal length; do not extrapolate from 8K |
| Embedding backward | CUDA `embedding_dense_backward` accumulates at reduced precision with ≤3072 indices, e.g. a short sequence alone in a micro-batch | length-tiered tolerance |
| CP compressed boundary | see [Packing and CP](#packing-and-cp) | before/after/at the boundary |
<a id="attention-d"></a>
### The D term in attention backward
With P rounded to FP8, visible O = fp8(P)·V ≠ softmax(QKᵀ)·V. If `D = rowsum(dO⊙O)` uses visible O while dS uses recomputed P, `rowsum(dS_i) ≠ 0` and dQ gains a same-sign per-row bias (Nemotron: about 3% of dQ/dK, v2 FAIL at steps 150/250; fixed 0.006–0.009, PASS). Mainstream FP8 attention (TE fp8_dpa, SageBwd) uses visible O, so a human decides after being informed.
Four checks:
1. **FA4 interface signature**: `_flash_attn_bwd` derives D in preprocess from the `out` passed; D cannot be passed. The public interface errors under `requires_grad`; the actor calls `_flash_attn_fwd/_bwd`. Fix: recompute BF16 O and LSE from dequantized Q/K/V, pass to bwd (layer backward +4–15%, end to end <0.05%).
2. **varlen LSE layout**: ln vs log2, per head vs per token, softmax scale folded or not. On mismatch, recompute with the training kernel.
3. **GQA 32/2**: dK/dV summed within each Q head group; A/A needs fixed summation order.
4. **When the saved Q/K/V are released under recompute=full**: if only dequantized Q/K/V are saved, recompute must read the same state; early release changes the values.
### Microbatch split and gradient difference
Different microbatch splits give about 2e-3 relative gradient difference (loss normalization, accumulation order, reduced-precision embedding backward at ≤3072 indices). Not a bug, but:
- Criterion denominator = `max(B16 reference floor, noise floor of the implementation under test)`, never the FP32 ideal.
- Both arms use the same split; if it changes, rerun A/A.
- grad_norm is not comparable across different microbatch counts.
### Control matrix
With policy/tokens fixed, test original layout, reversed order, target first/last, changed micro-batch, CP-boundary sample, each reconstructed by UID/global token coordinates. Real TP/EP changes may legally change the reduction; the experiment contract decides whether cross-topology bitwise equality is required.
Safe packing is not a fix for a metadata bug; it may only lower trigger probability. After the fix, the original layout and nearby boundaries must pass.
<a id="v2-optauth"></a>
### Training-side correctness acceptance: gradient comparison criterion v2 and optauth
Two independent checks beyond forward 0 diff.
#### Gradient comparison criterion v2
- Freeze the criterion (tolerance, denominator, thresholds) before the run, record its sha, never change it after.
- Planted-error check first: gradients ×1.25 and ×1.05 must both FAIL.
- Accept on several snapshots (steps 0/50/150/250); the D-term problem appeared only at step 150.
- On failure, inject reference gradients block by block to find the first module over the limit.
- Extended to FSDP2: EP4 + FSDP2 all four snapshots PASS, worst-class median 1.08–1.39, all floors from B16.
- Tool bug: with `overlap_grad_reduce=False` gradients were reduced again before the EP4 dump (fixed). The export path needs A/A too: ‖g‖ bitwise equal, local shards bitwise equal to the dump, replica diff 0/0.
- Cross-process bitwise backward is not required. An occasional non-bitwise first step comes from the thread-local Inductor deterministic config (the autograd thread's first compile picks 512/1024 blocking by timing); fix in [BI kernels: determinism](bi-kernels.md#determinism-takes-effect-only-on-the-main-thread).
#### Optimizer acceptance (optauth)
Verify the optimizer consumes clipped `main_grad`, the update lands on master, and deployed bytes = requantize(master). Freeze the criterion first. Split failures into criterion modeling gaps (e.g. a check inapplicable without replicated parameters; the equal-magnitude first Adam step invalidating γ-type criteria while the per-coordinate bound holds), documented and accepted, and implementation errors, fixed. Cover all parameter group types: a first-step probe without an FP32 wd_mult=0 parameter misses weight decay applied to that group.
#### Facts outside the two criteria
- Frozen base with LoRA: base modules still propagate input gradients ([§4.2](#42-a-frozen-base-still-has-to-propagate-input-gradients)).
- The two-step RL loop (θ0→θ1→θ2) accepts the forward; three-arm SFT accepts training effect; report all three separately ([validation](../workflow/validation.md)).
## VERL / Megatron integration
Upstream config keys may differ from what was used here. Find the actual call paths of rollout dump, old-logprob and skip/repeat in the current checkout and record the commit; local-patch options are not upstream API. Process content (recipe and PR contract, review material, two-step acceptance, E2E targets) is in [validation.md](../workflow/validation.md#recipe-review-materials-and-two-step-acceptance-agreement); PR ownership in [Code and delivery §8](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image).
### Where the integration layer is
Megatron's `experimental/lite/examples/verl/verl_mlite/` provides the verl engine adapter (`model_engine=mlite`, `pkg://verl_mlite.config`), installable into the unified image without reimplementing the engine in verl. Check engine registration, model factory, batch conversion, loss normalization and optimizer path; list reusable interfaces and gaps. Consistency uses mLite; the training baseline uses the existing recipe's default MCore path; do not mix them.
### Capture location
Save the DataProto (or equivalent) after rollout returns and before the first old-logprob; append trainer scores after. Record policy version, UID, input/response IDs, mask, position IDs, reward, route replay data if needed, and the UID -> packed-row mapping if reordered. Keep a pre-update checkpoint bound to the batch by content fingerprint; after strict load compare structure and keys and reject missing/unexpected keys.
### Historical DS4 entry point
DS4 used patched VERL options (`skip.rollout.*`, `DS4_SKIP_UPDATE=1`); confirm in code and logs that any such switch takes effect.
### Minimal adapter contract
The synchronous entry point takes batch, checkpoint, variant and output directory, then:
1. Records actual policy, effective config, libraries and parallel layout.
2. Reads the same rollout, repacks per variant, keeps logical UIDs.
3. Runs the real old-logprob forward and exports per-token logprobs in UID order ([Saving and comparing](debugging-mismatch.md#compare-per-token-logprobs)).
4. Records no generation and no update, and mutable-buffer initialization.
5. Waits for workers/collectives, closes its resources, returns an exit code.
It keeps native `no_grad`/`forward_only` semantics and replaces neither the [grad-enabled forward comparison](../workflow/validation.md#old-logprob-and-training-forward-dual-path-contract) nor update acceptance. A/B/C may share one long-lived worker once weight and state isolation is verified. The runner does not take over Ray/Slurm; the launcher blocks on remote jobs and reclaims on timeout. For injected or overlaid adapters, record the actual class and source file and assert the callsite was hit; `PYTHONPATH` can shadow the overlay.
### Entropy when old-logprob is not recomputed (bypass)
Bypass uses rollout logprob as old-logprob: `actor/ppo_kl` then measures rollout vs training forward exactly (0 under BI=1), and entropy comes from the grad-enabled update forward.
- **Put metric-only quantities under `no_grad`.** With `entropy_coeff=0` the loss still holds `0 * entropy_loss`, so entropy's backward runs and saves a logits copy.
- **Chunked entropy in the verl Megatron engine cannot backpropagate.** With `entropy_from_logits_with_chunking=True`, each chunk view's backward overwrites the shared saved logits in place; the shared version counter then raises `one of the variables needed for gradient computation has been modified by an inplace operation`. Hit whenever chunking meets entropy. Fix: [verl-project/verl#8195](https://github.com/verl-project/verl/pull/8195); until merged, with `entropy_coeff=0` compute entropy under `no_grad` with the same chunking.
- **Re-verify gradients after changing forward computation.** Nemotron chunked logprob: BF16 bitwise equal to unchunked; FP32 max diff about 1e-8 at max gradient about 4.
### Training extension
Pairing a shared inference forward with a native VJP: check recompute, CP metadata lifetime, PP payload, parameter registration, through the standard training entry point. After a real update, export a new rollout and trainer scores; old bundles replay only with the old policy, and new-vs-old output differences are expected, not a consistency failure. Modules using vLLM layers need vLLM process groups. `param_offload=True` changes the quantized master's `data_ptr` and fails the pre-export check; use optimizer offload only.
.agents/skills/train-inference-consistency/practices/code-and-delivery.md
-90
View File
@@ -1,90 +0,0 @@
# Code, experiment snapshots and delivery
Code ownership across three repos for cross-stack train-inference consistency delivery and "three PRs, one image" (§8), and PRs and cleanup (§9). Lifecycle, cold-start and model acceptance details are not repeated.
## 8. Three-repo delivery: from experimental implementation to three PRs and one image
Cross-stack RL delivery for a model is one vLLM PR, one Megatron PR, one verl PR, and one unified image built from the versions the three pin. Default targets are `vllm-project/vllm:main`, `NVIDIA/Megatron-LM:dev` and `verl-project/verl:main`; a fork may host the head but does not replace the main repo's target branch. This section covers change ownership and delivery checks; it does not replace the [ten-step flow](../workflow/README.md), the [numerical gates](../workflow/validation.md) or human recipe review.
### Before changing code: give a change map first
Fix the inference-side change locations in step 3 and the training side and integration interfaces in step 5; step 1 need not design every interface. Submit a short table of "feature → existing module/interface → proposed change location → reason for reuse or addition", covering at least model construction, forward/backward, quantization, parameter registration, checkpoint and the build entry.
Read the implementation and callers in the selected checkout before choosing file names; the paths below are design references, not a requirement to copy DeepSeek's directories. Do not write a standalone engine first and discover at the end that the existing integration layer or the real trainer cannot use it.
### vLLM: implement in the formal execution path, do not deliver a patch loader
- Development environment and code sync follow [same-commit worktree, per-file sync and build verification](../guides/environment-provenance.md#boundary-between-overlay-and-native-wheel); development sync is not how the final image is installed and does not change the PR's target branch.
- Model-specific selection goes in the model construction/execution entry; changes to shared kernels, backends and metadata go in the existing owning module, stating which models and callers are affected.
- BI=0 keeps the reviewed native path; BI=1 explicitly selects the aligned implementation. Check the actual backend and fail early when a required BI capability is missing; never silently fall back while still claiming consistency.
- Kernel numerics, binding and build registration ship together; CUDA/C++ changes enter the normal build, and JIT changes such as CuTe DSL record the compiler and specialization. Do not submit only the Python call while depending on an uncommitted `.so`, source file or build script inside the experiment container.
- Verify the default path is unchanged and cover the modified shared interfaces; when only a certain model/runner/topology is supported, restrict it explicitly rather than implying other models support BI.
Reference points in [vLLM #54955](https://github.com/vllm-project/vllm/pull/54955): the formal kernel integration in `csrc/batch_invariant/`, the DeepGEMM BI capability check, and the stable tie-break/prefill-decode semantics of sparse selection. These address operator numerics, backend authenticity and execution shape change respectively; they are more than a BI switch.
### Megatron mLite: separate the model, the aligned operators and the runtime
The train-inference consistency implementation uses mLite; the training baseline uses the default MCore path of the existing verl recipe, without the aligned implementation. Record the training baseline and the vLLM BI=0 inference performance baseline separately; they are not the same comparison. When target-model support has gaps, state them and get them reviewed; do not change the training scheme on your own.
| Change | Preferred owner | Should not go in |
| --- | --- | --- |
| Model layer structure, parameters and checkpoint mapping | model/protocol/checkpoint modules of `experimental/lite/model/<model>/` | a temporary model implementation inside a verl recipe |
| forward aligned with vLLM, autograd/STE, model-specific quantization | that model's aligned implementation and primitive modules | miscellaneous top-level mLite scripts, global monkeypatch |
| Reusable training/communication/optimizer capabilities | existing shared primitive/runtime, changed only for a real interface gap | one model's recipe constants and diagnostic branches |
| verl engine adaptation, batch/loss bridging | existing `experimental/lite/examples/verl/verl_mlite/` interface | another equivalent engine inside verl |
- Select native vs aligned implementations through the existing protocol/factory/explicit config; do not scatter environment checks across low-level forwards or replace instance methods/global functions at runtime. Small-scope branches can stay; do not copy the whole model for a tidy directory; reuse shared math and backward, and implement independently only primitives whose numerical semantics really differ.
- Do not assemble production operators with `inspect.getsource`, AST rewriting, or string replacement plus `exec`, and do not register models via `sitecustomize` or temporary `PYTHONPATH` hijacking. Import stable interfaces from the real dependency, and commit necessary changes to the owning module.
- Keep the parameter registration, state_dict/checkpoint, distributed reduction and recomputation contracts; do not replace the training master parameters with an inference cache the optimizer does not manage just to align forward. Backward is accepted separately per the [gradient contract](../guides/backward-validation.md#write-down-the-gradient-contract-first); it need not be bitwise identical with inference, and a kernel without autograd does not automatically need STE.
[Megatron #7050](https://github.com/NVIDIA/Megatron-LM/pull/7050) separates DeepSeek-V4's `lite/` from `vllm/primitive/` and reuses the existing `verl_mlite`, illustrating these boundaries. `verl_mlite` exists in official Megatron dev, not only in a fork; still check the version this task pins and its actual capabilities.
### verl: a thin recipe, not a second training framework
What is delivered depends on the route:
- **mLite route**: for non-observability changes, deliver ideally only one recipe script and one image build script, reusing the real trainer of the current version and the installed `verl_mlite`.
- **FSDP2 Day-0 route**: verl lacks the model, so the deliverable is the self-written model implementation in verl mainline, its Auto class registration and the recipe, in one verl PR; the model implementation itself is the deliverable and needs no separate "generic interface gap" pre-review. See [Training-side model support §3.3](../guides/training-forward.md#33-fsdp2-verls-hf-interface-path).
On both routes, when there is a real generic interface gap, first give call-chain evidence, the interface to change and the scope of impact, **get human review before implementing**, and make the minimal fix in the owning repo; do not hide code in scripts or move it to the wrong repo to make up the file count.
The two-file form of [verl #7688](https://github.com/verl-project/verl/pull/7688) and the `verl.trainer.main_ppo`/`model_engine=mlite` chain are the reference; its development mounts and overlay are not how the final image is delivered. For data, config differences and review rules, use the [DAPO reference and review materials](../workflow/validation.md#dapo-reference-and-review-materials); do not invent an "approximate DAPO" script that bypasses the trainer.
### Step 10: from clean PRs to verifiable delivery
1. **Three complete diffs.** For each repo, record the target owner/branch/base SHA agreed with the user and the PR head; check the merge-base, commit history and actual diff. Required model support must be in the PR or already merged into the target branch; another unmerged feature branch cannot be the base. Do not carry over a preview PR's history drift or unrelated deletions; pin cross-repo dependencies by commit. After changing the base or rebasing, re-review the diff; evidence from the old head does not transfer.
2. **One rebuildable image.** Build from three clean checkouts, pinning the base image digest, dependencies and necessary native/JIT build inputs, and install all runtime packages normally, including external adapter packages. Training and inference workers use the same immutable platform digest and the same Python/Torch and runtime dependencies; different nodes/processes are fine, but not two independent environments inside the image to dodge compatibility. It must run without host source mounts, startup-time patches or uncommitted overlays.
3. **Re-accept from a new container.** Check the actual worker's interpreter, imports, libraries and versions; verify the proxy, full-model zero-update and the real RL loop of the reviewed recipe per the existing gates; then record performance/quality. A successful build, offline disk replay and historical 0 diff do not replace verifying the final artifact; on a difference, shrink back to a fixed failing replay first instead of relaunching full training to locate it.
4. **One mapping.** Use the [report template](../workflow/report-template.md) to list the three PRs, image digest, recipe review version and acceptance receipts, marking each item pass/fail/inconclusive/not run. Update it after code cleanup, build or dependency changes; an unverified clean implementation does not inherit the overlay prototype's pass.
When slimming the implementation PR, non-essential diagnostic tools, docs and experiment artifacts stay in the private repo/model workspace; necessary regression tests can go with the PR. Do not delete tests that exist on the target branch, and a slim PR still runs tests. Follow each target repo's contribution requirements, record the regressions and model evaluations actually run, and submit for human review; if the gates do not pass, keep the not-accepted status.
### Reference versions, not proof of passing
All three PRs were open when reviewed, and the first two were marked DO NOT MERGE:
| Reference | Reviewed head SHA |
| --- | --- |
| vLLM #54955 | `4b5c627c23e6690b8987100d1cbf78e34ac14bf9` |
| Megatron #7050 | `327a238239448de6771beea185dc6902a8de3eaa` |
| verl #7688 | `640c5f0a317307e99636924804c40624242fa19b` |
A PR body may describe a different historical pin; when citing the design, distinguish the reviewed head, the experiment version the body claims and the version measured now, and do not treat the reference as fully correct or merged. Its Triton implementation also does not override this workflow's constraints on added/replaced kernels.
## 9. PRs and cleanup
- Three-repo delivery follows the target branches in §8.
- The PR description records: minimal purpose, base and patch stack, related existing PRs/why it is not a duplicate, upstream/private PRs, RL-only content, test commands and results, model impact, rollback method, AI assistance and human review requirements.
- Split a large PR into small independently reviewable PRs and prove equivalence before and after the split with the tree hash; open PRs only from your own fork.
- Resolve conflicts on the head branch of the PR the user specified; do not open a similarly named new PR to pass the original off as fixed.
- Cross-repo order is usually kernel/runtime → framework integration → recipe/pins → final image acceptance, with each later stage locking the actual revisions of the earlier ones.
After a patch is merged and a new image is built, archive the worktree, delete the temporary overlay, and keep the PR, commit, manifest and verification report.
## Minimal changes, contribution and delivery
Change the config along the entry script the user specified, so multiple temporary launchers do not drift apart. End with one launch entry that actually runs, with defaults, necessary overrides and image/code pins traceable; exports that exist only in chat are not delivery.
A fix lives with the state/feature owner: fix vLLM backend/selector problems in vLLM, not with a global monkeypatch in the training framework; a numerically matched forward must keep the real backward. Call existing reusable functions instead of building synonymous logprob metrics and compatibility switches.
An authorized independent audit does not transfer responsibility to the auditor. Before declaring completion, check: are all of the user's goals covered? Do the results belong to the final code/image? Does the wording hide any untested item? If the evidence is insufficient, write down the gap; do not wrap it in a percentage or a success marker.
.agents/skills/train-inference-consistency/practices/evidence-and-experiments.md
-179
View File
@@ -1,179 +0,0 @@
# Evidence and experiment rules
Rules for what counts as evidence, how experiments are designed and recorded, and when old evidence stops holding. Mechanisms live in the guides.
## 1. Evidence labels
PASS / FAIL / INVALID (did not enter the target path) / VACUOUS (condition had zero hits) / SUPERSEDED / HISTORICAL. These label evidence, not lifecycle states. The tables in [validation.md](../workflow/validation.md) and the [report template](../workflow/report-template.md) have only `pass / fail / inconclusive / not run`. PASS→`pass`, FAIL→`fail`; INVALID and VACUOUS become `inconclusive` with the reason in the notes; SUPERSEDED and HISTORICAL stay out of the current table and live only in the [candidate evidence ledger](#6-candidate-evidence-ledger) and historical reports. One SUPERSEDED piece of evidence makes the gate it covered stale. A successful run, train-inference consistency and a held-out score gain are different conclusions and cannot substitute for each other.
## 2. Evidence that must not be confused
- Enabling `VLLM_BATCH_INVARIANT` is not cross-framework bitwise alignment; the same text is not the same logprob.
- Zero difference in probability space, zero KL or identical printed numbers do not replace comparing valid logprobs in the original dtype.
- Replaying fixed inputs is a localization tool; statistics from old trajectories are not an unbiased KL under the current policy.
- The same function name, generated code or environment variables do not prove the same actual kernel, launch config and shared libraries.
- Output replacement, forced routing, disabling Graph and bypassing bad packing are valid causal experiments if you state what they change.
- Similar quality scores are not identical per-problem results; same-problem quality, bitwise logprob and 8K/1K performance are accounted separately. When the acceptance script fails, check the producer receipt first; do not resample or declare a pass automatically.
- Verify quantization rules on the deployed bytes, not dequantized weights ([precision pitfall P2](../guides/precision-quant.md#p2-verifying-a-rule-with-dequantized-weights-gives-a-false-positive)). QAT/low-precision training only narrows the weight gap; 0 diff comes only from training calling the same kernel as inference; see [How much QAT can shrink the mismatch](../guides/precision-quant.md#4-how-much-mismatch-qat-can-remove).
- Do not manufacture a pass by changing tokens, weights, model geometry, precision or tolerance. If the comparison baseline really must change, start a new experiment and re-verify quality.
- Do not carry DS4's EP/MoE config over to dense Qwen; do not manufacture a performance gain by weakening the BI=0 config, locally disguising BI as off, or removing existing correctness patches.
- Source tests passing is not the final image passing. Development overlay results belong only to their experiment; the delivered image digest, installed source and build artifacts must match the acceptance object, and the evidence chain is rebuilt after repackaging.
## 3. Inferences no longer carried forward
Historical numbers and criteria stay with the run that produced them; a new task does not inherit a past run's model, dataset, topology, performance statistics method or pass status. Lessons from retracted or narrowed claims:
- A number in a summary is not evidence: find the raw numerator and denominator for the same policy/workload/config, or retract the assertion.
- Execution completion is not numerical acceptance: finished steps, a started job, an existing adapter file or finite nonzero gradients do not show the target class ran or the comparison passed. Record the instantiated class and its import source, and accept non-empty original-dtype token comparison, module VJP against an independent reference, actual update, shutdown and long-term stability separately.
- Evidence does not transfer across candidates: a source pin, the same Docker tag, an earlier proxy, an experiment stack or a direct operator call does not prove the current candidate, image or compiled module dispatch. Verify branch, dirty patch, wheel, `.so`, dependencies and digest separately, and re-establish entry-gate evidence for the current candidate.
- A fixed split count or a pass on one path does not cover other batches, prefill/decode, the first round after resume or cross-node communication; each needs its own evidence.
- When an acceptance wrapper fails, find the failing layer before resampling or rerunning; a control-plane parsing error does not discard sealed numerical evidence.
Each new case records only "contract → minimal counterexample → discriminating experiment → root cause/minimal fix → regression on the original reproducer → scope/source", with the small raw receipt for disputed numbers. Merge status, recommended dependency versions and best-performing config are volatile; look them up in the target repo.
## 4. Experiment contract
### Define the objects of equality first
State the checkpoint/adapter/policy version, tokenizer and chat template, raw token IDs, position IDs, attention/response mask, padding/packing, and the logits processing on both training and inference sides.
For response token `y[t]`, compare `log p(y[t] | prompt, y[:t])`. Make explicit the prediction position of the first response token, whether EOS counts, and the exclusion rules for truncation and padding. Do not align logits rows of different positions.
Distinguish raw logits, raw logprobs, and logprobs after temperature/repetition penalty/top-k/top-p, and record which one the interface returns. The same seed does not guarantee the same scheduling or RNG consumption: freeze the tokens instead of resampling.
### Record these goals separately
| Goal | Valid evidence | Cannot infer |
| --- | --- | --- |
| Single-side repeatability | A/A with the same input and state | Consistency across batches or processes |
| batch/layout invariance | Same semantic sample identical under different positions, neighbors and packing | Training and inference are the same |
| prefill/decode parity | Two execution paths with the same forced tokens | backward or the training update is correct |
| Cross-stack forward parity | Real training forward against a frozen rollout | Full-vocabulary logits match, if only sampled tokens are compared |
| Effective training loop | Non-degenerate gradients, parameter change, post-update parity | Convergence or long-term stability |
When the user requires bitwise, compare original-dtype bytes and record mismatching/valid token counts; do not reach a pass through a float cast, probability-space rounding or a tolerance loosened afterwards. If the goal is an error budget, fix the metric and threshold in advance and measure each side's own noise; generic numerical alignment need not be forced to bitwise.
A finite logprob is necessary; an empty mask, all-zero placeholders, missed capture, NaN or the wrong weight version cannot pass. A zero may be legitimate; tell it from a placeholder by hook count, score source and a negative control.
### Recipe and coverage record
When building a train-inference recipe, also fix these inputs, recovering missing items from the existing config first:
| Item | What to record |
| --- | --- |
| Workload | Training context length, rollout input/output length distribution, concurrency, packing, training token batch |
| Hardware | GPU architecture/count, GPU memory, interconnect and target TP/CP/EP/PP/DP, trainer and rollout separately |
| Numerics | master/deployed weight, activation, KV, accumulation and gradient dtype; quantization granularity, scale rule and update method |
| Scope | Raw sampled logprob or full-vocabulary logits, applicable hardware/software, Graph/cache path |
| Budget | Throughput, peak GPU memory, quality and gradient criteria, experiment count and stop condition; with an unknown threshold, report the measurement and do not declare a pass |
Record the measured baseline of the current config, then build a strong performance baseline per the main flow, without exhaustive tuning first. Rust Bench 8192/1024, ignore EOS, c=1–64 is a fixed comparison and does not replace the target context and natural-length workload; a user-specified workload governs, with the difference reported.
A proxy records "full-model function → proxy position → test case → uncovered items" and keeps the hidden/head/expert/quantization geometry and layout that affect kernel selection. Reuse existing modules and tests, but after changing model, dtype or parallelism re-check the contract; changing the config does not carry over the old acceptance.
### Environment fingerprint
Keep the code commit, dirty patch hash, untracked runtime files, image digest, GPU architecture, Torch/Triton/CUDA, actual import path, and native `.so` path and hash. Record TP/PP/CP/EP/DP for trainer and rollout, dtype/quantization/scale, kernel backend, and Graph/compile/cache and autotune config.
Read the resolved config and actual modules/libraries inside the worker, not launcher settings. For weights, record content version and update status of derived weights/scales, not just the path. Keep immutable known-good and known-bad artifacts before diagnosing.
### Scope of conclusions
For each result, record model scale, valid sample count, independent token count, lengths and boundaries, process count, GPU/parallel config and number of updates. Four samples repeated into batch32 are still four independent samples; mixed scheduling a small model never triggers cannot cover the full model.
Record machine budget, maximum experiment count and stop condition. GPU or system faults give `inconclusive`; a failed numerical comparison gives `fail`; `pass` for a scope requires every declared acceptance item.
### Experiment discipline
Write these nine rules into the record before comparing two arms; if any is unmet, the conclusion is "pending re-check".
| Rule | How to do it | Mechanism and example |
| --- | --- | --- |
| One image per recipe | Both arms use the same image and recipe, differing only in the recipe switch; inject changes with an overlay and record the hash | [One recipe, one image](../guides/recipe-and-acceptance.md#one-recipe-one-image-deliver-true-on-policy-and-the-regular-baseline-with-the-bi-switch) |
| One seed per config | One seed by default; add more only when the user asks for variance, reported separately | Extra seeds spend budget on variance; most problems are systematic |
| Small-workload iteration | Iterate with small batch, short outputs and few steps; run the recipe default workload only for final acceptance | The default workload allows only one or two iterations a day |
| Freeze the criterion first and plant an error | Fix comparison window and tolerance in writing before the run; confirm with a known error that it can judge FAIL; never adjust the threshold after a failure | A threshold chosen afterwards can always pass; [paired acceptance](../guides/recipe-and-acceptance.md#paired-acceptance) |
| No conclusion from a single e2e run | Measure noise first; on the same hardware and image, report the median of multiple pairs and single pairs separately | [Performance and profiling](../guides/performance.md) |
| Same fabric | Before comparing communication bottlenecks, confirm both arms use the same transport and record the relevant environment variables | [Same-fabric fairness](../guides/debugging-mismatch.md#same-fabric-fairness) |
| Same-node comparison | Run paired arms sequentially on the same node or node group, with node names recorded | Memory fragmentation, resident processes and interconnect topology differ between nodes |
| JIT pollution: check `tl_after_bench=0` | After the run, confirm no new compilation after the bench; results with recompilation are void | [Performance and profiling](../guides/performance.md) |
| Take metrics by key name | Read numbers from raw metrics (jsonl/W&B) by full key name and cite the key in the report; never by column position in log lines | An `actor/grad_norm` column once read as entropy reversed a conclusion |
List confounders before comparing; see [paired acceptance](../guides/recipe-and-acceptance.md#paired-acceptance).
## 5. Execution state
For every comparison, record the actual execution state, and keep results from different execution states as separate evidence.
### Fields that must be recorded
- Resolved Graph mode, capture sizes, whether the step is prefill/decode/mixed, and kernel call evidence; a successful capture does not mean every step uses the same graph, nor numerics identical to eager.
- Whether the prefix cache hit, and whether it was disabled/isolated during diagnosis; passing with the cache off does not mean the native lifecycle is fixed.
- Weight lifecycle phase: cold load, sleep/wake, Graph replay, reload in a new process, one row each.
- Whether chunk boundary state and partial input buffers of stateful layers such as Mamba/SSM are saved.
### Keep separate evidence
Eager, compiled with Graph off, and production Graph each keep their own artifacts; hooks run at capture, not replay, so a capture dump does not represent replay. List separately: diagnostic runtime override vs formal source candidate, Graph vs no-Graph, forward-only vs real `train()`+autograd. A new process passing does not prove the old process's cache/Graph is fixed; two fixes that each pass go through the gate again when stacked. Once CUDA state is corrupted, end that experiment and collect no more numerical evidence from it.
### Where the mechanisms and methods are
| Question | Where to go |
| --- | --- |
| Graph capture vs replay, the Graph/eager 2×2 comparison with fixed composition, which bucket the filler falls into | [BI kernel appendix](../guides/bi-kernels.md) (cudagraph row) |
| prefix cache semantics, Mamba chunk boundary state, bitwise comparison with "cache off as the gold standard" | [rollout-engine](../guides/rollout-engine.md) |
| Localizing CUDA asynchronous errors | [debugging-mismatch](../guides/debugging-mismatch.md#isolate-the-side-first) |
## 6. Candidate evidence ledger
Goal: keep an old pass from becoming a new pass.
When an experiment changes model scale, quantization dispatch, attention, MoE, Graph or policy weights, create one ledger row per **candidate execution closure**, then decide what to rerun.
### What one row of evidence is bound to
| Field | What to record and check |
| --- | --- |
| Object | Original layer numbers and real adjacency, parameter/quantization geometry, tokenizer, checkpoint content, policy version; a proxy's manually spliced boundaries listed separately |
| Execution on both sides | Trainer/server image digest, source and overlay hashes, actually imported classes and files, resolved config, native/JIT artifacts |
| Path | `train()`/autograd/whether backward runs, actual dispatch of attention/norm/linear/MoE, compile/Graph/cache, TP/EP/PP layout |
| Workload | Content digest of prompt/response tokens, valid mask, length distribution, packing, concurrency, reward/advantage and actual accepted sample count |
| Criterion | Original dtype, finite valid token count, bitwise mismatch, sample count, first difference, and the independent states of gradient/update/normal exit |
| Artifacts | Inputs, score, worker/rank receipt, command, time, parent version, hash and `pass`/`fail`/`inconclusive`/`not run` |
`policy_manifest_sha256`, the exported `policy_content_id` and the rollout receipt hash answer three separate questions, **manifest version, weight content, run artifact**, and are not comparable with each other. If the exported content ID is the hash of the export receipt, follow the receipt to check shard hashes against the files the server actually loaded; one equal version string is not enough.
### Entry gates from cheap to expensive
1. **Module.** Use a real triggering input and a non-unit scale/real layout; A/A first, then the candidate's forward against an independent reference; for gradients, check dX/dW/other differentiable quantities separately with the same dY. A pass covers only the shapes, arithmetic and VJP tested.
2. **Minimal real-adjacency proxy.** Cover the actual order of layer types, not just the operator set. With a fixed 8192+1024 tokens, check inference BI/PD first, then per the [dual-path contract](../workflow/validation.md#old-logprob-and-training-forward-dual-path-contract) verify native old-logprob and the grad-enabled actor forward separately; for a fixed-token comparison only, do not add GSM8K/RL to the proxy. Record source closure, actual class, Graph, non-empty valid tokens and original-dtype mismatch. When any item changes, re-examine whether this step still transfers.
3. **Full-model zero-update.** Generate a real-workload rollout with the full model on the **current** serving candidate, saving tokens and raw scores; the actor replays the same policy and history before the optimizer update. Check the earliest failing sample, then expand to the agreed sample size; if unequal, stop here and narrow the boundary instead of working around it with five-step RL or a quality evaluation.
4. **Per-version update.** For each θk, compare on-policy first, then check reward/advantage, backward, master parameters and the optimizer step, then sample θk+1; "the last training step exited" is not "the last version's comparison passed".
5. **Independent acceptance.** Module gradients, whole-model gradient coverage, real RL, quality, performance, clean exit and the final image each get their own result. Numbers from a previous candidate or another execution mode stay as historical comparison and never fill the current candidate's pass cell.
A fixed-weight serving overlay without RL lifecycle operations such as sleep/wake is only a serving screening candidate, even if BI/PD and 8K/1K paired performance pass; once the lifecycle is restored, rerun the affected proxy, zero-update and post-update comparison. Do not splice historical throughput and another candidate's RL pass into one "completed" version.
If a phase fails midway, classify the failure as numerical, control, environment or resource plane; keep the sealed rollout and producer hash. A control-plane acceptance bug is handled by a separately versioned recovery acceptor reading the original artifacts; never edit the old receipt, or resample and call it the same trajectory.
### When evidence must be invalidated and rerun
| Change | Minimum regression; old conclusions that cannot be carried over |
| --- | --- |
| Proxy's original layer numbers/adjacency or hidden, head, expert geometry change | Coverage table, BI/PD, training forward; the old proxy's 0 diff no longer counts for promotion |
| Quantization algorithm, scale, actual dispatch or JIT/compiler change | Real module input/quantized bytes, proxy, full-model zero-update; a direct-call unit test does not replace the compiled module |
| Attention/MoE operator or its training adapter changes | Real class and source, same-input operator, relevant VJP, proxy and full-model zero-update; a pass on a substitute backend is diagnostic only |
| Graph/cache/packing/parallel layout changes | The corresponding BI/PD, replay and state lifecycle; an eager pass does not cover FULL Graph |
| Optimizer or master→deployment export changes | Gradient and one-step update, export byte hash, post-update comparison; a θ0 forward pass is not enough |
| Final image/overlay/native build changes | Verify the closure in the worker and rerun affected numerics and performance; development-mount evidence does not transfer |
If you can prove a change does not affect a gate, write the dependency closure and reason in the ledger to skip a mechanical rerun; equivalence cannot be decided from a file name, PR description, same tag or "looks like a cleanup".
### Separate cells for performance and quality
Tune BI=0 well under the agreed weight/KV precision and quality requirements, record the real dense GEMM, MoE expert, all-to-all and attention backends, and freeze the denominator after human audit. Pair BI=1 and BI=0 on the same full workload, nodes and candidate closure; Rust Bench 8192/1024, ignore EOS, c=1/8/32/64 is one fixed throughput workload, not natural DAPO lengths or GSM8K quality. Check effective tokens/error count first, then compute the loss per the [pre-reviewed statistical protocol](../guides/recipe-and-acceptance.md#pairing-and-computation); do not reach the budget by slowing BI=0, changing weight precision or changing prefix cache/Graph parameters.
Quality needs the same problems, prompt, chat template, seed, generation limit, scorer and no-retry policy; report per-problem pairing and the absolute score difference. A quality pass does not transfer to a candidate whose numerical path changed. If the user changes the threshold, the current one governs; keep the old threshold and change date for historical reports.
## 7. Failure records
Keep failure records as they are; append new results, with no overwriting, deletion or editing. For W&B, use a separate project per comparison and read back through the API to check no step is missing. For freezing the criterion, planting a known error and "no conclusion from a single e2e run", see [§4 Experiment discipline](#experiment-discipline).
.agents/skills/train-inference-consistency/workflow/README.md
-122
View File
@@ -1,122 +0,0 @@
# Train-inference consistency workflow: C0–C10
Turn a train-inference difference into a repeatable experiment with fixed input and fixed weights, locate the first divergence, fix or optimize, then return to the real training loop. Keep the user's model, precision, parallelism and consistency target. Switches, thresholds, passing scope and implementation paths from historical cases, reference PRs and implementations already in the training framework (such as an existing true-on-policy path) are experience, not defaults for the current model, dataset or hardware; measure their numerics and training-step time before reusing them.
Forward consistency, backward conformance to the gradient contract, training quality and performance are accepted separately; none replaces another. User config comes first. Recover missing information from the source and resolved config, and ask only for what cannot be recovered and affects the approach. Choose the low-precision recipe by hardware and existing paths, not by a default precision, model or framework. Changing the model or precision requires rechecking coverage and acceptance evidence.
Work in one directory per task, with state changes in an event log and artifacts carrying provenance sidecars. Details are in [validation.md](validation.md) (entry gate table, dual-path contract, tiered criteria, SFT three-way control, two-step RL sequence) and [numerical contracts](#numerical-contracts-and-shared-forward).
Archive materials as you go: each run ID links config/scripts, execution identity, raw results and conclusion; the step index states the current candidate, failed/excluded items, unverified items and the human-confirmed version. C1 keeps the full tuning comparison table; C10 references the reviewed baseline ID. Layout: [archive experiment materials by step](../guides/debugging-mismatch.md#archive-experiment-material-by-step).
## Gate table
Iterate as "single operator / fixed replay -> proxy -> staged full model"; do not rerun every gate each time, and do not skip the human stop points at C1 and C9. Inference-only tasks apply only the relevant gates and state that the training loop is not accepted.
| Gate | Input | Action | Pass condition | Artifacts | See |
| --- | --- | --- | --- | --- | --- |
| **C0 Module inventory and route** | checkpoint, vLLM implementation, target (BI=1 / BI=0), hardware | Per module: inference kernel, whether training implements the same math, quantization rounding, discrete choices; choose FSDP2 / Megatron-Bridge / mLite | Inventory covers all modules; route fits "BI=1 needed? how many backwards missing? what parallelism?" | Module inventory, route and rationale | [Training-side model support](../guides/training-forward.md) |
| **C1 Recipe and strong baseline** | Model, target context length, hardware | Decide weight/activation/KV precision, quantization rules, runner, parallelism, execution config; decide the training master start (BF16 original or dequantized weights; do θ0 bytes equal the official checkpoint); agree scoring scope, consistency standard, performance budget (inference throughput, and limits for training forward and forward+backward against the native training path on the main topologies; a human decides which parallelism to enable and which topologies are main); tune BI=0 with **vLLM Rust Bench** at 8192 in / 1024 out / ignore EOS / c=1–64; for MoE, measure MoE / all-to-all backend combinations | Baseline = the formal recipe's BI=0 config (including R3, PC, gpu_mem), not the serving-only optimum; multi-round measurements and backends recorded; **freeze the candidate and submit for human audit** | Tuning table, resolved config, image digest, audit materials | [Precision, quantization and QAT](../guides/precision-quant.md), [Recipe and evaluation](../guides/recipe-and-acceptance.md), [Performance](../guides/performance.md) |
| **C2 Proxy model** | Full model structure | About 4 layers covering all operator types, needed adjacent-layer interactions and state paths; real hidden/head/expert/quantization geometry and needed parallelism constraints; coverage table "full-model feature -> proxy layer/module -> test case" | Uncovered items listed; the proxy is the daily development environment | Coverage table, proxy config | [Training-side model support](../guides/training-forward.md) |
| **C3 Inference BI/PD** | Proxy, fixed weights and tokens | Verify batch and prefill/decode invariance separately, covering chunks, mixed scheduling, cache state; consider kernel suitability for training / large token counts. On failure, coarse-scan layer boundaries, bisect, capture identical operator inputs, fix kernel/metadata/layout/state | Bitwise identical in all three cases under fixed input; **proves only consistency inside inference** | Comparison matrix, minimal failing replay | [BI kernel appendix table](../guides/bi-kernels.md), [Fault localization](../guides/debugging-mismatch.md); before changing inference read [implementation ownership](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image); sync source per [runtime provenance](../guides/environment-provenance.md) (same-commit worktree + per-file hash) |
| **C4 Structural performance** | Proxy, real operator shapes | Profile prefill/decode (and training forward and forward+backward once wired): launches vs tokens, CPU gaps, sync, repeated quantization/copies/allocation, abnormal workspace; target CUDA Graph captures/replays correctly | No per-token Python loop or many-small-kernels defects (autoregressive decode steps do not count); consistency retest passes after the fix | Profile, defect list | [Performance](../guides/performance.md) |
| **C5 Training forward↔inference** | Proxy, fixed weights, input/output tokens, state | Reuse inference primitives and quantization, layout, mask, position, routing, reduction, logprob semantics; keep the training parameter container and native entry point; same-packing A/A and different-packing A/C first; then compare native old-logprob recomputation and the grad-enabled actor forward separately per the [dual-path contract](validation.md#old-logprob-and-training-forward-dual-path-contract) | Both paths: valid logprob mismatch=0 in original dtype; record non-empty valid token count and execution identity. Inference BI/PD, single operator or diagnostic substitute paths do not replace this gate | Dual-path receipt | [Training-side model support](../guides/training-forward.md) |
| **C6 backward / STE** | Forward implementation from C5 | Reuse existing backward, autograd, communication, recompute; where the gradient passes quantization rounding, decide the STE surrogate (identity/clipping, master or dequantized value, scale rule). A kernel without autograd does not force an approximate gradient; no AST rewriting or monkeypatching in the production path | Every module has a "reuse / glue / newly written" verdict and rationale, and its backward uses only what the forward saved; the training step is within budget on the main topologies fixed at C1 | Module-level backward inventory, training-step time table for the main topologies | [Backward validation](../guides/backward-validation.md), [STE contract](../guides/precision-quant.md) |
| **C7 Independent gradient validation / SFT control** | C6 implementation, same input, weights, saved state, same-dtype upstream gradient | Compare against a same-contract reference that does not call the backward under test: dX, dW, other trainable gradients, one optimizer step; fix tolerances first; report bitwise and tolerance separately. When required, compare production, an independent surrogate with the same visible forward, and native training per the [three-way SFT control](../guides/backward-validation.md#layer-3-proxy-sft-three-arm-comparison) | Independent comparison passes; without SFT, mark "quality not covered", a loss decrease is not quality. Forward 0 diff or finite gradients do not replace this gate | Gradient audit, three-way SFT results | Same as above |
| **C8 Proxy RL loop** | Proxy, real framework entry point | Follow the [exact sequence](validation.md#two-step-rl-sequence): θ0 compare -> update to θ1 -> θ1 compare -> update to θ2 -> θ2 compare; check finite, meaningful gradients, actual master change, normal exit | No cross-version comparison; a forward skipped by zero advantage does not count; proxy limited to logprob alignment: mark "loop not verified"; proxy asked to verify forward only: full-model zero-update comparison first, then the training gates; never start multi-step RL directly | Two-step RL receipt | [Two-step RL sequence](validation.md#two-step-rl-sequence) |
| **C9 Recipe review and full-model entry gate** | Launch script, image digest, code version, full resolved config, differences from the reference recipe (verl #7688 etc.), proxy smoke | **Human review; start the full model only after the recipe version is confirmed** (data, length, batch, parallelism, precision/quantization, optimizer); fill in the [entry gate table](validation.md#full-model-entry-gate-pass-proxy-train-inference-forward-first); then add layers, length, concurrency, parallelism one at a time and verify mixed scheduling, cache/Graph, multi-step loop; with checkpoint restore in scope, check alignment after the first restore | No escalation while C5 is fail / inconclusive / without artifact; two-step acceptance is the formal recipe with only fewer steps, same sequence/batch/data/algorithm; on full-model failure, shrink to a minimal replay first | Review materials, entry gate record, two-step results | [Checkpoint, resume and memory](../guides/checkpoint-resume.md), [recipe review materials and two-step agreement](validation.md#recipe-review-materials-and-two-step-acceptance-agreement) |
| **C10 Optimization and final E2E** | Full-model consistency passed, delivery PRs, reviewed recipe | E2E profile -> BI=1 hotspots (shape, dtype, layout, frequency, time share) -> representative micro bench -> correctness and performance on all target shapes -> E2E. **Triton is forbidden**; use CuTe DSL or better, CUDA/CUTLASS, or an existing implementation. Each round regresses inference, training forward and gradients by impact | On the delivery image through the real trainer entry: θ0 compare, update, θ1 compare, update, θ2 compare; 0 diff = equal original-dtype logprob for the same weight version, token history and valid mask; 8K/1K Rust Bench BI=0/BI=1 per the [pairing protocol](../guides/recipe-and-acceptance.md#pairing-and-computation); no added long runs or gates | Delivery manifest, pairing results, loop receipt | [Performance](../guides/performance.md), [Paired acceptance](../guides/recipe-and-acceptance.md), [Delivery](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image) |
The consistency implementation uses mLite by default: it separates model, operators and runtime, so the inference forward can be reused, backward/STE wired in independently, and the existing `verl_mlite` interface and MCore distributed optimizer wrapper kept. The training baseline follows the default MCore path of the existing verl recipe without the alignment implementation, and is accounted for separately from C1's vLLM BI=0 inference baseline. **mLite does not inherently guarantee 0 diff or better performance**; verify the target model's parallelism, recompute and checkpoint support with the actual recipe. Both arms share one recipe and one image, and `VLLM_BATCH_INVARIANT` selects the arm (see [bi-switch-recipe](../guides/recipe-and-acceptance.md)); the delivery contract is in [validation.md](validation.md#recipe-and-pr-delivery-contract).
## Run the cheapest failure criterion first
1. host metadata / shape arithmetic: minimal counterexample of a pure function and boundary parameterization.
2. Single operator: same input, real dtype/layout, non-finite detection, repeat / cross-process / Graph.
3. frozen batch replay: the original failing A/C and A/A, full valid-token comparison.
4. engine: batch composition, prefill/decode, chunk boundaries, cache/Graph, actual parallelism scope.
5. Real training: non-degenerate reward/advantage, finite non-zero gradients, optimizer step, parameter change, frozen-token alignment after the update.
6. Production acceptance: full model, agreed multiple steps, restore lifecycle, quality evaluation and paired performance.
If a second run adds no localization information, change the probe. Not reproduced, empty artifacts and an unexecuted path are not passes. When the failure is in backward, replaying the forward is not enough: save checkpoint, batch, RNG and state and move to a minimal backward experiment. Details are in [validation.md](validation.md).
## RL delivery gates
Training and inference use **the same immutable image digest and the same Python/Torch and runtime dependencies**. They may run on separate nodes and processes, but two environments inside one image cannot be used to dodge compatibility. Complete at least two real updates through the target RL framework's actual entry point and verify each round: original-dtype logprob 0 diff before the update, an actual parameter update, and 0 diff for the new policy after it. Saved rollouts, offline replay, separate images and cold start are diagnostics only. This is a delivery agreement, not a technical claim about all RL systems.
Deliver one vLLM PR, one Megatron PR, one verl PR and one unified image, targeting `vllm-project/vllm:main`, `NVIDIA/Megatron-LM:dev` and `verl-project/verl:main`. Non-observability changes in verl should be limited to the recipe and the image script where possible. Follow [from clean PR to delivery](../practices/code-and-delivery.md#step-10-from-clean-prs-to-verifiable-delivery): check the target branch and full diff, pin cross-repository commits, build the image from clean source, and fill in the [delivery manifest](report-template.md#delivery-manifest).
## Baseline and consistency run as two parallel tracks
- **Baseline track**: C1's BI=0 tuning, backend combinations and multi-round performance evidence. Do not sacrifice default performance for consistency in advance; after reasonable tuning, freeze the candidate and wait for human audit.
- **Consistency track**: do not wait for the baseline; advance through C2–C8 on the proxy. C9 still requires the entry gate; parallel progress is not a reason to go to the full model early.
- Share the model/precision/hardware/image contract; isolate code, environment and artifacts; keep consistency patches out of BI=0. While the baseline is not final, mark the scope consistency covers and re-verify after config changes.
- Do not add nodes beyond the budget. On a single node, GPU experiments time-share; during a formal benchmark, run no GPU, heavy CPU or I/O work from the other track. The final performance gap is accepted against the human-confirmed strong baseline with full consistency evidence.
## Two human stop points
**After C1: baseline audit.** Once measurements support "reasonably tuned", freeze the candidate, report and wait; do not keep chasing the performance limit. Materials: model/weight version, hardware and node count, image digest / build time, code and dependency versions, reproducible launch and Rust Bench commands; weight/activation/KV precision, runner, parallelism and scheduling, Graph/cache state, actual operator backends, and a statement that alignment patches do not change BI=0; configs tried and why, valid tokens per concurrency, warmup, repeat count, throughput/latency and variation, failed items and uncovered scope; the config proposed for locking, the [statistical measurement basis](../guides/recipe-and-acceptance.md#pairing-and-computation), repeat count, per-concurrency criteria and budget, and open questions for humans. An unconfirmed candidate baseline must not be the final denominator. Re-audit after changing model/precision/hardware/image; while waiting, hand over or release nodes as needed.
**Before C9: recipe review.** See the gate table. Changes to workload or numerical/update semantics are reviewed again; pure path or log fixes are not.
## Find by problem
| Current problem | Read first |
| --- | --- |
| Setting the consistency target or evidence basis | [Experiment contract](../practices/evidence-and-experiments.md#4-experiment-contract) |
| Reproducible input not yet saved | [artifacts](../guides/debugging-mismatch.md), [rollout-engine](../guides/rollout-engine.md) |
| rollout and old logprob differ: side unknown / stable failing sample / identical operator inputs still differ | [Fault localization](../guides/debugging-mismatch.md), [numerical contracts](#numerical-contracts-and-shared-forward) |
| Results change with packing, CP/TP/EP | [Layout and length](../guides/training-forward.md) |
| Fails only under Graph, compile or cache | [Execution state](../practices/evidence-and-experiments.md#5-execution-state) |
| Wiring in backward/STE, comparing gradients, SFT | [Reuse backward and backward validation (three layers)](../guides/backward-validation.md) |
| Going to the full model, shrinking a full-model failure, declaring completion | [validation.md](validation.md) |
| Candidate/backend/image changed: is old evidence still valid | [Candidate evidence ledger](../practices/evidence-and-experiments.md#6-candidate-evidence-ledger), [inferences no longer carried over](../practices/evidence-and-experiments.md#3-inferences-no-longer-carried-forward) |
| Aligned; recovering performance, profiling, ablation, BI kernels | [Performance](../guides/performance.md), [BI kernel appendix table](../guides/bi-kernels.md) |
| Separating quality, fixed-length Rust Bench and normal exit | [Evaluation and paired acceptance](../guides/recipe-and-acceptance.md) |
| Image, source, wheel, submodule, overlay versions confused | [Runtime provenance](../guides/environment-provenance.md) |
| Three-repo ownership, PRs, unified image | [Three-repository delivery](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image) |
| GB200 devices, NVSHMEM, DeepEP, cross-node initialization failure | [Communication fault ladder](../guides/debugging-mismatch.md) |
| checkpoint save/load OOM, `unexpected pos`, host memory rising, first resume misaligned | [Checkpoint, resume and memory](../guides/checkpoint-resume.md) |
| BI=0 baseline arm, two modes in the same recipe | [bi-switch-recipe](../guides/recipe-and-acceptance.md) |
| Requantization rules, QAT selection, activation fake quantization | [Precision, quantization and QAT](../guides/precision-quant.md) |
| Many repeated samples within a rollout group | [rollout-engine](../guides/rollout-engine.md#1-sampling-seed-collision) |
| Different seeds give the same samples (greedy n repeats are identical; SGLang deterministic mode without a seed always uses `sampling_seed=42`) | [rollout-engine §1](../guides/rollout-engine.md#1-sampling-seed-collision) |
The ten kinds of evidence that must not be confused are in [evidence and experiment rules §2](../practices/evidence-and-experiments.md#2-evidence-that-must-not-be-confused). Writing comparisons and replays: [save and compare](../guides/debugging-mismatch.md#compare-per-token-logprobs). Record experiments and report with the [report template](report-template.md). This workflow grants no extra GPU, job stop, image publish or cross-repository write permission.
## Numerical contracts and shared forward
### Check the arithmetic that actually executes
The same formula does not guarantee the same floating-point result. Check accumulation dtype, cast position, FMA, reduction tree, chunking/split-K, approximate activations, quantization scale/rounding, epilogue and the input dtype of the final logprob. The same source can also produce different execution configs through autotune, compile specialization or a different GPU.
A shared kernel still needs its caller, arguments, layout and compile boundary checked; a direct call and a torch.compile fused path may differ. If needed, keep the same forward behind an existing opaque custom-op boundary; do not introduce unproven arithmetic rewrites to imitate the output.
Quantization execution checks (scale taken from the actual runner, quantization dispatch, quantizer outside the `nn.Module` tree, training vs deployment quantization grid) are in the [precision, quantization and QAT guide](../guides/precision-quant.md#training-side-numerics), the grid in [P8](../guides/precision-quant.md#p8-training-grid-differs-from-the-deployment-grid); format and scale records are in [conversion chain §1](../guides/precision-quant.md#1-four-stages).
First check that the semantics exist on both sides. When a feature is inside the kernel on one side and external or missing on the other, the difference does not go away with precision settings. Run a two-layer minimal comparison per model feature (attention sink, RoPE variant, final norm, etc.), enabling one at a time. In the external MiMo case SWA alone gave 0 and SWA+sink gave `3.906e-03`, because only one side's kernel included the sink.
### Discrete decisions come before last-bit error
After router/indexer logits, also compare the selected set, the order and the tie-break. Top-K with the same set in a different order still changes downstream gather/reduction. A deterministic tie-break defines the index order for equal scores and keeps the invalid/padding index contract.
For MoE, also fix or verify dispatch, expert capacity, grouped GEMM geometry, quantization and combine. Do not hide a router difference by forcing expert indices; forced routing only isolates the downstream.
The top-k/top-p truncation set is also a discrete decision: a last-bit logit difference moves tokens near the cutoff in or out of the set, and the normalization denominator of the processed logprob jumps with it. Compare the candidate set before processed logprob. MiMo-V2.6 replays the candidate set together with R3 in production: rollout records a fixed-shape, full-vocabulary bitmap (avoiding host sync and truncation), and training renormalizes within the set. If the contract adopts this replay, state in the report that it masks the set difference, and count the set flip rate separately.
### Test shape and process boundaries
For the same victim, construct batch1 / multiple neighbors, different physical positions, lengths on both sides of a boundary, cross-process autotune, prefill/decode and Graph/eager. Different launches when only M changes are not a numerical difference; measure the output. Finding no difference covers only the shapes tested.
SSM chunk recurrence and per-token recurrence can be two legal but different rounding paths. Unifying them requires matching the chunk grid, boundary state precision and gate order, not just a deterministic switch.
For history-dependent operators, a single token's scan/conv input and output can make an upstream history difference look like an error in that operator. First find the first differing token at each layer over the **full prompt batch**, then capture norm, projection, conv, scan, gate and output projection in order at that layer. A single-position probe only narrows candidates; it does not prove the first cause. A plain Python hook in a compiled graph may hold stale warmup data; verify capture freshness with a predictable sentinel such as embedding/residual, and if needed use the custom-op boundary that actually executes and a separate buffer.
### Boundary between shared forward and backward
With a shared inference forward + native differentiable VJP, verify separately: forward parity, a finite gradient for every trainable parameter, and gradient comparison under a cotangent representable in the same dtype. Operator-level VJP verification does not prove the whole-model gradient is bitwise equal to the original training implementation.
Do not detach the inference output and use it as a trainable substitute. A finite gradient does not mean effective learning; also check that advantage/reward is not degenerate, the gradient is non-zero, and parameters and optimizer state are updated.
### Changing the arithmetic baseline
If you deliberately change both sides' arithmetic baseline to align them, keep the old baseline, build a new frozen rollout, state that "old outputs are not guaranteed to stay bitwise unchanged", and run model quality and performance regressions. Success numbers from the old baseline do not carry over to the new implementation.
.agents/skills/train-inference-consistency/workflow/report-template.md
-119
View File
@@ -1,119 +0,0 @@
# Train–inference consistency report
## Outcome and scope
State pass/fail/inconclusive for the requested scope. Record model/layers, independent samples, valid tokens, dtype, topology, Graph/cache paths, update count and untested conditions.
Link the human-approved acceptance checklist: required vs out-of-scope checks, recipe revision, two-update stop condition, quality/resume coverage and target-repository submission requirements. Out-of-scope checks are unverified, not passed; finishing the short run does not complete missing required checks.
## Immutable experiment
Record source commits and dirty patch hashes, loaded libraries, image digest, GPU, checkpoint/policy identity, input bundle digest, score semantics, resolved config, commands, resource budget and artifacts.
Separate source, installed wheel, loaded native libraries and JIT specializations. State whether the measured runtime and the final delivered image are identical; if not, name the missing verification instead of inheriting the result.
For recipe work, record training context, rollout length distribution, concurrency, packing/token batch, hardware/interconnect and master/deployment/activation/KV/accumulation/gradient dtypes, plus quantization granularity, scale/update rules and the quality/throughput/memory budgets defined before testing. Unknown budgets stay unaccepted even when measurements exist.
## Delivery manifest
For the three-PR / one-image delivery, fill this from the final clean build, not development overlays. Use full SHAs and immutable image identity; mark missing evidence as not run.
| Deliverable | Review / source identity | Build or acceptance evidence |
| --- | --- | --- |
| vLLM PR | URL; target owner/branch/base SHA; head SHA | Build command and wheel/native/JIT inputs |
| Megatron PR | URL; target owner/branch/base SHA; head SHA | Installed model / primitive / adapter identity |
| verl PR | URL; target owner/branch/base SHA; head SHA | Recipe and image build script paths |
| Unified image | Base and final platform digests; dependency pins | Build log; actual trainer and rollout worker imports/libraries; no development source mounts |
| Approved recipe | Script SHA, resolved config and human approval | Changes from reference recipe; dataset and topology |
| Final validation | Commands and artifact IDs tied to the above image | Proxy / full zero-update / θ0–θ2 / quality / performance status, each reported separately |
Each PR is complete against the agreed base, not stacked on an unmerged feature branch; cross-repository dependency pins are expected. Every required runtime file outside the PRs must be a pinned existing dependency; an uncommitted production patch is a delivery gap. A changed base, build input or recipe requires reassessing affected evidence.
## Proxy coverage and gradient contract (when training is in scope)
Map each full-model feature to its proxy layer/module and test artifact; name omitted depth, geometry, topology and execution paths. Per module, record visible forward semantics, trainable variables, master vs dequantized surrogate values, STE/clipping/scale rules, fixed discrete choices and differentiable scores, VJP accumulation/collectives, and saved/recomputed state with its memory cost.
Link the independent reference and explain why it does not call the tested backward. Record shape/dtype/layout and random, zero, sparse and captured upstream-gradient coverage; per input/parameter: missing/nonfinite gradients, max absolute error, relative L2, norms and cosine where defined. Include predefined thresholds with rationale, original forward-byte comparison and one-step master-parameter deltas.
## Symptom and localization
Record per-request/token differences, A/A and A/C results, side-isolation evidence, last equal / first unequal semantic boundary, identical-input operator comparison, competing explanations and counterexamples.
## Cause and change
Explain the violated arithmetic/metadata/state contract and the smallest triggering input. Separate the actual repair from diagnostic interventions. Link code changes and the preserved regression. Write unknown if evidence does not establish a cause.
## Verification
Record the proxy upgrade decision before escalating to full-model BI consistency or training. Do not merge inference-only and cross-stack parity into one result; a failed, missing, empty or stale forward comparison blocks escalation. If the user scoped the proxy to fixed-token forward parity, mark SFT/RL not run and list the later training gates.
| Proxy upgrade gate | Status (pass/fail/inconclusive/not run) | Bound command / artifact and valid-token count | Decision |
| --- | --- | --- | --- |
| Proxy feature/operator coverage | | | |
| Proxy inference BI and prefill/decode parity | | | |
| Proxy inference ↔ native old-logprob recomputation | | | |
| Proxy inference ↔ actual grad-enabled actor forward | | | |
| Backward / SFT / RL, if in proxy scope | | | |
| Upgrade to full-model BI work | | | |
For each cross-stack row, follow the [dual-path contract](validation.md#old-logprob-and-training-forward-dual-path-contract): model geometry, frozen policy/checkpoint identity, source/image digest, dtype, topology, model/grad modes, output requires_grad, actual backends, raw logprob mismatch count and first mismatch. Keep native no-grad scoring semantics; its pass does not prove grad-enabled forward parity. Capture scores before updates and loss transforms.
| Gate | Command / artifact | Result and count | Scope limitation |
| --- | --- | --- | --- |
| Original failing replay | | | |
| Operator / metadata regression | | | |
| Engine and production layouts | | | |
| Independent module VJP and one optimizer update | | | |
| Controlled proxy SFT (A/B/C paths) | | | |
| Real backward/update | | | |
| θ0 / θ1 / θ2 parity | | | |
| First post-checkpoint-resume export / parity | | | |
| Model quality | | | |
| Clean shutdown / sustained run | | | |
| Final immutable image, without development overlays | | | |
For two-update RL, keep a receipt per policy θ0, θ1, θ2: input/mask identity, valid token count, original-byte mismatches, actual trainer-forward execution and update norm/count. Record both forward-path results at matching frozen policy versions. Worker exit is checked at final shutdown, not at each policy transition.
## SFT controls and quality (when training is in scope)
Separate SFT-A (production), SFT-B (same visible forward at the same snapshot plus an independent surrogate) and SFT-C (verified native training); these labels are independent of replay packing variants. Link same-checkpoint/same-batch A/B gradients and optimizer deltas; once trajectories diverge, later forward values are not assumed identical. A/C differences include forward/quantization changes.
Record initial master and optimizer state, data/order, labels/masks, packing, effective-token normalization, batch/accumulation, learning-rate schedule, clipping, precision/loss scaling, RNG/dropout/recompute and intentional differences. Keep raw train, held-out and fixed-probe loss, gradient norm, update RMS, clipping/nonfinite/skipped-step counts, throughput and baseline/seed variation when available. Use one optimizer-step or effective-token axis; report predefined windows/thresholds, missing controls and whether evidence is only a short smoke run.
## Performance
Record matched work, warm/cold timing, node pool and paired order, repetitions, latency/throughput, kernel/rollout/full-step results and numerical receipts. Separate measured improvements, rejected candidates and estimates.
For 8K/1K acceptance, link the [approved aggregation protocol](../guides/recipe-and-acceptance.md#pairing-and-computation): formula, repetitions, per-concurrency thresholds and invalid-leg replacement rules. Include all raw pairs, exclusions with reasons and the approved aggregate per c; median paired loss and loss from median throughputs are different statistics. Other phase timings are diagnostics, not gates.
Where relevant, include training time, peak memory, launch count vs token batch, CPU gaps and repeated quantization/allocation. Separate early feasibility profiles from final full-model optimization; normal autoregressive decoding is not evidence of an avoidable per-token prefill/training loop.
State whether this is a BI-only causal comparison or each side's best valid stack, and list all differing flags/backends. Include prompt count vs expanded requests, actual input/output lengths, completed output tokens, wall-time denominator and whether the run is fixed-length synthetic or natural-EOS. Do not sum overlapping GPU ranges or percentages from separate ablations.
## Evidence corrections and retention
List superseded claims with the correcting artifact and reason. Record which original failure bundles, checkpoints and small receipts are retained, their storage budget/location, and any missing eval steps. Never fill a curve with a different policy lineage.
## Remaining work
List concrete gaps and the smallest next experiment. State whether production hooks/config were restored. Keep raw evidence; do not upgrade historical or partial results into current completion.
## Experiment record (YAML)
```yaml
# Train-inference consistency replay experiment record: each variant changes only one variable, everything else is fixed.
purpose: "frozen-policy old-logprob replay"
target: "bitwise selected-token raw logprob"
rollout: "/absolute/path/to/frozen-rollout" # original framework batch, not exported scores
checkpoint: "content id + pre-update recovery path"
environment: "image digest, source+patch hashes, libraries, GPU"
topology:
trainer: "TP/PP/CP/EP/DP"
rollout: "TP/PP/CP/EP/DP"
state: "RNG, mutable buffers, Graph/cache and reset policy"
budget: {max_runs: 3, timeout_seconds: 600}
variants:
- {name: A, change: "none (A/A baseline)", cmd: "", output: "", exit: null, result: ""}
- {name: B, change: "none (A/A repeat)", cmd: "", output: "", exit: null, result: ""}
- {name: C, change: "packing reversed", cmd: "", output: "", exit: null, result: ""}
```
.agents/skills/train-inference-consistency/workflow/validation.md
-169
View File
@@ -1,169 +0,0 @@
# Tiered regression and completion criteria
## Full-model entry gate: pass proxy train-inference forward first
When backends, compile paths or delivery images have changed, check against the [candidate evidence ledger](../practices/evidence-and-experiments.md#6-candidate-evidence-ledger) whether old results still apply. A passing old candidate is history only and does not fill the current candidate's entry gate cells.
Before full-model correctness or training experiments, provide the evidence below. Missing or not executed counts as not passed; repeated full-model launches are no substitute. Proxy inference-internal BI/PD and proxy training forward↔inference logprob comparison are two independent results: the former passing does not mean the latter passes.
| Gate | Minimal evidence |
| --- | --- |
| Proxy coverage | Inventory of operator types, necessary adjacent-layer interactions, quantization geometry, state and parallel paths; uncovered items stated |
| Inference BI/PD | Batch, prefill/decode and mixed-scheduling comparison on the proxy with fixed weights and tokens; proves only consistency inside inference |
| Training forward↔inference | Same proxy, same policy weights, fixed tokens: verify native old-logprob recomputation and the grad-enabled actor forward used in training separately, each against the target production inference path; keep evidence per the dual-path contract below |
| Backward and quality | [Independent module gradient comparison](../guides/backward-validation.md#layer-1-module-level-vjp-comparison) with reference implementation, tolerances and raw results; controlled SFT per task scope, or an explicit statement that quality is not covered |
| Real RL update | At least two effective updates, θ0/θ1/θ2 same-policy comparison, actual parameter change, normal exit; sequence below |
| Performance feasibility | launch/sync/memory overhead on real shapes and the proxy profile, ruling out bottlenecks such as per-token launches in batchable prefill/training; the target Graph path works |
This is an escalation gate, not the completion criterion: two-step RL does not prove convergence, and the proxy does not prove full depth, capacity, long context or cross-node correctness. On the full model, extend coverage one item at a time. Inference-only tasks apply only the relevant gates, state that training is not accepted, and add no training work.
Before escalating, record per gate in the [report template](report-template.md#verification): `pass`/`fail`/`inconclusive`/`not run`, command, artifact path and the bound model geometry, weights, source/image, precision and parallelism. Accept only current-candidate evidence. For training forward↔inference, the bitwise target requires original-dtype mismatch of valid tokens to be 0; tasks with an agreed error budget use the predefined standard and cannot use it to report a bitwise failure as a pass. An unexecuted training forward, an empty mask, or passing only single-operator or inference BI/PD tests must not release to the full model. Keep the first failing sample on the proxy for layer-by-layer localization, and after the fix rerun the original failing case and the affected matrix.
### old-logprob and training forward dual-path contract
- **Native old-logprob recomputation:** follow the framework's actual path and its `no_grad`/`forward_only` semantics; do not force autograd on for acceptance. Its result proves only that this scoring path matches rollout.
- **Actual training forward:** with the same frozen θ, token history and valid mask, run the grad-enabled actor branch training actually uses, capture logprob before the loss transform, and compare with rollout independently. Keep the target packing, recompute, operators and parallelism; do not substitute a diagnostic forward. Passing does not replace the later independent backward validation.
- Record per path: model mode, grad mode, output `requires_grad`, operator class / backend / source hash, original dtype, non-empty valid token count and mismatch. Diagnostics (substitute attention/MoE, Graph off) prove only that path; `forward_only` alone is not grounds for rejection, and native old-logprob passing does not prove the grad-enabled branch.
- The comparison happens before the corresponding policy update, for both proxy and full-model θ0 entry. If the formal recipe has several minibatch updates per round, do not require 0 diff between the updated actor and the old-θ rollout, and do not modify the algorithm to manufacture equality; record the real parameter version and compare only within one version.
When an adapter / overlay is injected dynamically, the receipt must also record the actual object's class, source file and key config and assert that the target callsite was hit. A script directory placed before the overlay in `PYTHONPATH` can silently load an old adapter; a 0 diff from such a run cannot vouch for the new candidate.
If the user limits the minimal proxy to fixed-token logprob comparison (or forward only), do not require GSM8K, SFT or two-step RL; mark those gates `not run`, not passed. Scale up as agreed only after the proxy train-inference forward passes. Before any multi-step full-model RL, first complete a zero-update same-policy logprob comparison with **full-model θ0 real workload, real training mode, target production inference mode**, then the necessary gradient validation and effective-update checks. A fixed 8K/1K proxy pass does not replace the full-model variable-length workload. When full-model θ0 fails, stop optimizer updates, keep the failing trajectory and shrink to the first divergence. This scope adjustment does not waive the proxy forward gate.
### Limited exceptions that allow going to the full model early
- Establishing the BI=0 performance/quality baseline of the full target model is not approval for full-model BI debugging or training.
- When a fault depends on real depth, capacity, cross-layer state or cross-node topology and cannot be reproduced on a semantics-preserving proxy, targeted full-model diagnosis is allowed. Before the experiment, write down why the proxy is insufficient, the hypotheses to distinguish, the minimal scale and the stop condition.
- Scaling up does not scale up resource permissions: no extra GPUs, no stopping other jobs, no bypassing the model/runner scope the user set.
### Faster iteration and fallback on failure
Judge each change's impact first: minimal counterexample, single operator or fixed replay, then the relevant proxy gates, then the staged full model; do not repeat all ten steps per commit. Changes to shared arithmetic or gradients regress the affected forward/backward; final delivery still runs the full agreed acceptance.
On a full-model failure, first save the input before the first divergence, the weight version, metadata, cache/routing and necessary RNG state, and extract a module or few-layer reproducer. Iterate on it only after confirming the fault survives shrinking. If the fault disappears, a passing shrunk test is not a fix, and the full-model reproduction is kept.
Fixed rollout / module inputs may be cached and replayed in an isolated resident process, bound to fingerprints of weights, input, config and necessary state; a new implementation still reruns the comparison. Each scale-up states the existing evidence and the open question the new scale answers.
If a candidate fix changes the sampling trajectory or EOS stop point, save the new token history and replay training old logprob on it with the same weight version. The old trajectory remains a fixed-input diagnostic; comparing new inference against the old response length is not an end-to-end pass. List evidence separately for a diagnostic runtime override, a formal source candidate, Graph/no-Graph, forward-only and real `train()`+autograd. An isolation package may be used for localization; final acceptance returns to the delivered code and the default target execution mode.
## Start from the cheapest failure criterion
1. host metadata/shape arithmetic: minimal counterexample of a pure function and boundary parameterization; if needed, exhaustively check equivalence before and after the fix.
2. Single operator: same input, real dtype/layout, non-finite detection, repeat/cross-process and Graph.
3. frozen batch replay: the original failing A/C and A/A, full valid-token comparison; UID mapping passes.
4. engine: batch composition, prefill/decode, chunk boundaries, cache/Graph, actual parallelism scope.
5. Real training (when in scope): non-degenerate reward/advantage, finite and meaningful non-zero gradients, optimizer step, actual parameter change, frozen-token alignment after the update.
6. Production acceptance: full model, agreed multiple steps, update/restore lifecycle, quality evaluation and paired performance.
Not every task runs every tier. Report the tiers completed; when the user asks only for diagnosis, do not start full training. Extend existing test suites first; do not replace behavior tests with assertions on source strings.
## Preventing false greens
- Zero valid tokens, comparison not executed, missing rank/sample: invalid experiment.
- All-zero score, mask all false, the same sample repeated: check the producer and coverage; a passing comparison is not enough.
- probability/KL/logPPL is zero: still check the original logprob bytes.
- Zero gradient or identical rewards: maybe no effective update; check advantage, update count and parameter difference.
- Small model / short output passes: confirm the full model's mixed-step and long-context trigger conditions.
- A shrunk proxy's error magnitude is not the full-model baseline: in the external MiMo case the 12-layer proxy was `1.866e-12` and full-depth real weights were `1.143e-02`. The proxy only verifies the execution path.
- In-training mismatch metric is 0: confirm the micro-batch really ran the trainer forward. The framework may skip the forward for a zero-advantage batch and return 0, a trivial match; count only reports with non-zero advantage.
- A success marker written before the process was killed: separate a completed numerical experiment from normal exit / long-term stability.
- probs_diff is 0 but many sibling samples within a group are identical: a sampling diversity defect, independent of the consistency metric. Count distinct rollout samples per group; the mechanism and fix are in [rollout-engine §1](../guides/rollout-engine.md#1-sampling-seed-collision).
- The launcher shows `COMPLETED 0:0`: read the recipe's own exit code. When tests report passed, check the skipped count; GPU cases may be skipped wholesale by a marker or environment gate.
## Unified measurement basis and matrix for forward consistency
Fix policy, token history, position, mask and tensor semantics, and record per item the original dtype/shape, valid token count (greater than 0), non-finite count, bitwise mismatch count, maximum absolute difference and first failing position. A dtype difference is a contract inconsistency first; do not convert and then claim a bitwise pass. After aligning semantics, compare contiguous tensor bytes that preserve the bit pattern or an equal-width integer view; floating-point `==` does not distinguish positive and negative zero. NaN/Inf in a valid logprob is a failure even with identical bytes. Sampled-token comparison does not extrapolate to the full vocabulary.
Keep cases and artifacts per item for the target execution scope, and list unsupported / not run items:
| Dimension | Comparison under fixed semantics |
| --- | --- |
| A/A | Repeated execution with the same weights, input and state; in a separate process if needed |
| BI | Single request, different batch sizes, companion requests, permutation/position and packing |
| PD | Full prefill, chunked prefill and step-by-step decode on the same history, covering mixed scheduling |
| Length and branches | Real lengths, chunk/causal boundaries, and cases that trigger sparse selection, MoE and special modules; not only random tokens |
| Execution state | Target Graph capture/replay, cache reuse/reset and the actual parallel layout |
The bitwise target requires mismatch=0; error-budget tasks report against the predefined standard and are never recorded as a bitwise pass. Gradient tolerances are defined independently, see [module VJP](../guides/backward-validation.md#layer-1-module-level-vjp-comparison) and [SFT control](#proxy-sft-three-way-control).
## Two-step RL sequence
This sequence is the lower bound for entry; "container started / training script exited 0" is not a pass.
When the task includes the full train-inference loop, complete at least this sequence; every actor recomputation uses the same policy and token history as its rollout:
```text
θ0 → rollout0 → actor(θ0) recompute and compare
→ backward0 / optimizer update → θ1
θ1 → rollout1 → actor(θ1) recompute and compare
→ backward1 / optimizer update → θ2
θ2 → rollout2 → actor(θ2) recompute and compare
```
The θ2 comparison needs no third optimizer update. To compare outputs across policies, use the same frozen token history; do not rely on sampling happening to match.
For each stage save the policy content identifier, token/mask, non-empty valid comparison count and raw byte comparison result; for each update save finite loss/gradient, coverage of parameters that should have gradients, master update norm and optimizer step.
For a bitwise task the θ0, θ1 and θ2 comparisons must all have mismatch=0, with two effective updates and a normal exit actually completed. Skipped updates, both sides reading an old cache, an empty mask or comparing only probability summaries are not allowed.
A small master update may not cross a quantization bucket, so deployment bytes need not change every step. Judge an update by master update norm and optimizer step, not by step count or byte change.
## Proxy SFT three-way control
The three paths, controlled variables, recorded items and localization order are in [Backward validation layer 3](../guides/backward-validation.md#layer-3-proxy-sft-three-arm-comparison). If the module gradient comparison fails, fix the implementation before starting SFT.
## Performance and model quality
When forward arithmetic, precision, scorer or inference path changes, rerun the same quality evaluation within the human-reviewed scope and the target repository's requirements. Record model/data, problem count, prompt format, generation settings, statistical variation and unfinished parts; mark out-of-round items unverified, and add no SFT or long runs. A loss decrease or two-step reward change does not prove no quality regression.
Micro benchmarks, fixed-token parity, generation quality, normal exit and multi-step stability are separate. A run with four non-zero updates and 139,260 bitwise-equal tokens that ended in a DataLoader error supports four-step alignment, not clean shutdown or 50-step stability.
The full model covers the reviewed scope (real depth / weights, target context, concurrency, packing, parallelism/hardware, cache/Graph) and reruns affected correctness gates; proxy curves do not extrapolate to full-model quality. Final performance is accepted on 8K/1K per the [paired statistics protocol](../guides/recipe-and-acceptance.md#pairing-and-computation); training and full-step times may be listed as diagnostics and add no performance gate. Scope and stop condition of the two-step short run are in the [general two-step acceptance agreement](#general-two-step-acceptance-agreement).
## Failure classification and delivery
Use `pass`, `fail`, `inconclusive` with a reason. Keep source/patch/image/bundle hashes, commands, per-item counts, earliest difference, excluded hypotheses, and before/after results on the original reproducer. Cleaned-up code needs re-verification or explicit behavior-equivalence evidence; old tests do not vouch for new code.
Passed items stand despite unfinished ones, but the final conclusion is limited in scope. Use the [report template](report-template.md). Formal upstream contributions also follow the target repository's duplicate check, human line-by-line review, tests and AI assistance disclosure; historical cases do not mean the review was done on the author's behalf.
If the delivery includes an image, the gate object is the final immutable artifact, not a development image plus a source mount. Verify the closure per [runtime provenance](../guides/environment-provenance.md); accept save per [01 §3.2](../guides/checkpoint-resume.md#32-reproduce-and-discriminate), load per [§4](../guides/checkpoint-resume.md#4-load-enough-capacity-does-not-mean-enough-for-the-transient-peak), and the first resume per [§6](../guides/checkpoint-resume.md#6-how-to-prove-resume-is-correct). Passing environment checks proves only the preconditions for running, not GPU numerics or performance.
## Recipe review materials and two-step acceptance agreement
The integration point is in [training-side model support](../guides/training-forward.md); only process agreements are here.
### Recipe and PR delivery contract
Locate the existing integration layer in the selected version first. Megatron's `experimental/lite/examples/verl/verl_mlite/` provides the verl engine adapter; an external integration package installs normally into the unified image, and removing the runtime overlay does not require reimplementing the engine in verl. Check the actual engine registration, model factory, batch conversion, loss normalization and optimizer path, and list reusable interfaces and specific gaps.
- **verl PR** contains by default the model's DAPO recipe and the unified image Dockerfile. Public interface changes beyond these two: submit the gap and plan first, implement after human review; necessary regression tests may stay.
- **Megatron PR** contains the model, operators and necessary support in the existing `verl_mlite` integration layer; do not move generic code that does not belong in Megatron there just to keep the verl PR at two files.
- **vLLM PR** contains the necessary inference-side model, kernel, state and weight loading changes.
- Per the [general delivery target](../practices/code-and-delivery.md#8-three-repo-delivery-from-experimental-implementation-to-three-prs-and-one-image), submit complete PRs to the three main repositories, including the required model adaptation and not stacked on another unmerged feature PR; pin cross-repository dependency versions and state the build order. Consistency uses mLite and the training baseline uses the existing verl recipe's default MCore path (see [the workflow](README.md#gate-table)); do not merge them into one training scheme.
### DAPO reference and review materials
[verl #7688](https://github.com/verl-project/verl/pull/7688), at reviewed version `640c5f0a317307e99636924804c40624242fa19b`, uses Megatron's integration package through `model_engine=mlite` and `pkg://verl_mlite.config`, so the verl diff needs only the recipe and the Dockerfile. Follow its real `verl.trainer.main_ppo` execution chain; do not assume its config keys, DeepSeek precision/topology, mounts or overlay fit the current model.
The reference case's model, data and evaluation set are not defaults. Choose training data / reward by target model and human agreement, and keep the training set distinct from the quality evaluation set. Finish config resolution and the necessary proxy smoke, then submit for human review:
- Executable launch script, unified image digest, commit of each repository and full resolved config.
- Parameter differences from the reference recipe with reasons: data/reward, prompt/response length, rollout count, batch/microbatch, parallelism, quantization / KV precision, optimizer.
- Proxy check results, known uncovered items, and the acceptance scope of the full-model zero update and the two real updates.
Formal full-model E2E starts after a human confirms a specific recipe version and the proxy entry passes. Workload or numerical/update changes are reviewed again; path and log changes that leave input and execution semantics unchanged are only recorded. This gate does not block already authorized builds or proxy development.
### General two-step acceptance agreement
- **Freeze versions first.** Record the heads of the three delivery PRs, the unified image digest and the full recipe config. Changes that affect results require a version update and the related verification; old artifacts cannot be reused for acceptance.
- **Review the recipe first** (materials above). Confirming process or scope does not approve an unsubmitted recipe. The review checklist also marks which quality, restore and long-run checks are required or out of scope for this round, and the evidence the target repository requires; out of scope does not mean passed, and the short run ending does not hide missing required items.
- **Run two effective updates by default.** Use the formal recipe and shorten only the execution length to two steps; do not downgrade data, sequence length, rollout count, batch, optimizer, algorithm, parallelism or backend. If the stop step count drives the learning-rate schedule, keep the formal recipe's learning-rate/warmup semantics for the first two steps instead of compressing the full schedule into two; the implementation is reviewed with the recipe.
- **Keep the existing correctness gates.** After the proxy real training forward passes, run the full-model [two-step RL sequence](#two-step-rl-sequence) with actual parameter updates; a fixed-weight replay does not count instead.
- **Performance is accepted only on 8K/1K.** On the frozen version, use Rust Bench at 8192 input / 1024 output, ignore EOS, c=1/8/32/64, against the human-confirmed BI=0 baseline and performance budget. Training forward, backward and full RL step times may be recorded but add no performance gate.
- After the two updates, the θ2 comparison and the agreed 8K/1K performance pass, stop the short run and report. Overall acceptance is still judged item by item against the review checklist; if a required item is missing, state that it is not complete. Do not extend to 100 steps, add an MCore long-run comparison or new quality gates, and do not extrapolate the two-step result to long-run training quality.
### E2E comparison target
Run the [two-step RL sequence](#two-step-rl-sequence) through the real trainer entry point. 0 diff means equal original-dtype logprob for the same weight version, token history and valid mask; θ0 and θ1 outputs may differ. Verify the actual master update each round (not step count or byte change). Disk replay is a localization tool and does not replace acceptance through the real trainer entry point.
Finally re-verify with the image built from the delivery PRs and the reviewed recipe. Record standalone engine smoke, old image numerical results and historical performance separately; they do not replace full evidence from this execution entry point.
.claude/skills/train-inference-consistency
-1
View File
@@ -1 +0,0 @@
../../.agents/skills/train-inference-consistency