Hardware-Aware FP4 FlashAttention-4

• The report covers full-FP4 forward attention.Source: REPORT FOCUS · p. 1
• The report covers quantized causal backward.Source: REPORT FOCUS · p. 1
• The report measures GB200/B300 speed–accuracy trade-offs.Source: REPORT FOCUS · p. 1

1 Introduction

• Softmax reduction, exponentiation, probability construction, and publication remain between attention’s matrix products and can become the bottleneck as those products accelerate.Source: §1 · p. 2
• Direct-P maps normalized scores directly to MXFP4 probability codes and normalizes the same rounded values consumed by the PV product.Source: §1 · p. 2
• Causal backward reconstructs probabilities from forward-saved quantized Q/K state, while learned projections remain a separate precision boundary.Source: §1 · p. 2
• 2.13× is the maximum BF16 forward throughput achieved by Direct-P on favorable Blackwell shapes.Source: §1 · p. 2
• The causal path’s single-GPU 8B update reaches up to 1.14× speedup, while distributed training retains FP8 P/V because every tested MXFP4 P/V trajectory diverges.Source: §1 · p. 2
• At D128, tensor-memory ownership and operand readiness limit QK, softmax, and PV overlap, motivating more allocatable overlap capacity.Source: §1 · p. 3

2 Previous Work

• Direct-P retains HAO’s outer schedule and changes the stage that converts a completed score fragment into the scaled probability operand for PV.The contribution is the score-to-PV interval, not a replacement of the inherited pipeline.Source: §2 · p. 3

2.1 Why FlashAttention uses an online softmax

• Online softmax avoids storing full score and probability matrices in HBM, but each score tile must become probabilities before multiplication by V.Source: §2.1 · p. 4

2.2 Implementation baseline

• HAO identifies online P quantization and scale movement as costs, while SageAttention3 supplies numerical ideas but targets a different hardware interface.The execution schedule here is inherited from HAO and FA4 [4, 11, 12, 13].Source: §2.2 · p. 4

2.3 Two-query execution model

• At D128, HAO’s 16-warp CTA advances two query stages while a load warp reuses a K/V stream and softmax warpgroups process stages in ping-pong fashion.The protocol interleaves stage work when dependencies permit and publishes the first K64 probability operand before the tail [p61].Source: §2.3 · pp. 4–5

2.4 Why the D128 layout cannot look farther ahead

• Two score banks and two output banks fill the 512-column TMEM allocation, so the next QK cannot claim a score destination until PV releases overlaid storage.This is a storage-ownership handoff; a barrier can report availability but cannot create another destination.Source: §2.4 · p. 5, Figure 1 · p. 5

2.5 Why probability tiles are published in halves

• The scaled-FP4 PV instruction consumes K64, so two adjacent N32 probability quarters must be ready before useful work issues.Publishing after Q1 lets first-half PV overlap Q2/Q3 and work from the other query stage, without reducing the score bank allocation.Source: §2.5 · p. 6

3 What Prevents a Full-FP4 Speedup?

• Replacing HAO’s FP8 probability–value product with FP4 leaves both a timing problem and a numerical-range problem.Source: §3 · p. 6

3.1 Problem 1: P lies on the critical path

• FP4 speeds QK and PV, but not score reduction, exponentiation, synchronization, or scale publication; only work removed before the first legal P tile can shorten the critical path.Work removed after that publication may not change latency [11].Source: §3.1 · p. 6

3.2 Problem 2: P needs both range and cheap scales

• When an NVFP4 E4M3 scale rounds to zero, its whole block disappears; a row-level stabilizer fixes range but adds online scale and inverse-correction state.Source: §3.2 · p. 7
• MXFP4 uses a power-of-two block amplitude at N32 granularity; with the largest E2M1 code of 6, the reconstruction convention corrects an MXFP4 P/V product by 1/36.The choice is a latency–range trade-off, not a claim of greater intrinsic accuracy than E4M3.Source: §3.2 · p. 7
• Table 3 diagnoses probability-format range at D128 using zero-scale fraction and lost probability mass, rather than kernel timing.Its setup compares matched N32 blocks after renormalization, with stabilized NVFP4 as the high-fidelity control and unscaled NVFP4 exposing underflow.Source: Table 3 · p. 7, §3.2 · p. 7
• Direct-P chooses MXFP4 over stabilized NVFP4’s row-level correction and FP8’s loss of the FP4 PV instruction, redesigning production and normalization for coarser probabilities.Source: §3.2 · p. 7

4.1 Design overview

• Direct-P changes only the interval from a ready FP32 score fragment to a legal FP4 probability operand, preserving HAO’s outer schedule.The retained policies use NVFP4 QK, MXFP4 P/V, K64 PV issue, and the two-query HAO layout.Source: §4.1 · p. 8, Table 5 · p. 11

4.2 Fix 1: map scores directly to E2M1 codes

• Direct-P classifies normalized scores directly into E2M1 codes, avoiding accurate exponentiation whose extra precision is discarded before PV.Packed FMA and Blackwell’s native converter were retained over threshold trees, lookup tables, and custom packing, which generated more machine code.Source: §4.2 · pp. 8–9
• The score map’s A and B constants calibrate E2M1 decision boundaries; layer-wise maps were rejected because isolated substitutions did not predict the composed all-FP4 trajectory.The fast fit uses A = 1.50, B = 1.20; Wan activation evaluations use A = 1.60, B = 0.95.Source: §4.2 · p. 9

4.3 Fix 2: normalize the represented probability

• The denominator is computed from the same rounded FP4 codes and block scales consumed by PV, keeping it consistent with the approximate numerator.Four packed payload words are reduced with byte operations and DP4A.Source: §4.3 · p. 10

4.4 Fix 3: guard only extreme model logits

• A sampled Q/K-row guard handles extreme Wan logits, while a reordered calculation fixes the layer-39 subnormal failure without a second scan, barrier, or stable-softmax fallback.The 1.3-billion- and 14-billion-parameter Wan models use (M, L) = (110, 16) and (112, 14), respectively.Source: §4.4 · pp. 10–11

4.5 Two operating points

• The retained policies share NVFP4 QK, MXFP4 P/V, K64 PV issue, and the two-query HAO layout; optional Wan guard preparation is outside the timed kernel.Anchor permutations and folded Q/K scale preparation should be fused into an upstream layout or quantization step for deployment.Source: Table 5 · p. 11, §4.5 · p. 11

5 Experimental Setup

• Kernel speed excludes projection and optimizer work, and short fixed-token updates do not establish training quality.Table 6 states what each comparison supports.Source: §5 · p. 11, Table 6 · p. 11

5.1 Metrics and comparison rules

• Relative-L2 is reported alongside cosine and RMSE because cosine can hide magnitude error and RMSE depends on output scale.Principal rows pair latency and numerical error from one deterministic case.Source: §5.1 · p. 12

5.2 Operator benchmark suites

• The operator suite compares noncausal D128 cases using shared prepared Q/K/V, fixed seeds, and matched timing controls across TK and HAO kernels.A saturation suite adds H64 cases; D64 uses a distinct tile and ownership regime and is excluded from the main comparison.Source: §5.2 · p. 12

5.3 Timed scope and operand contract

• Timed kernel scope includes attention work through the output epilogue, but excludes dynamic Q/K/V quantization and optional K/V permutations.The fast path requires adjacent NVFP4 Q/K scale blocks folded for K64 access; ordinary block-16 scales violate its operand contract.Source: §5.3 · p. 13

5.4 Fixed-input model evaluations

• Fixed-input evaluations replace eligible attention calls in ViT and BERT, while the ViT-MAE and Wan tests pair outputs against BF16 references.These are held-input inference comparisons, not training-stability claims; Wan leaves cross-attention and the rest of the model in BF16.Source: §5.4 · p. 13

5.5 Reproducibility

• Tables are generated from JSON manifests, with commands, policies, source identities, hardware metadata, and evidence checksums collected in the appendices.Historical format paths and rejected timing experiments are kept separate from the primary comparison.Source: §5.5 · p. 14

6.1 Operator performance and accuracy

• 2.023× geometric speedup over HAO BF16 across nine GB200 D128 rows; fast peaks at 2998 TFLOP/s.Accurate reaches 1.669× and 2416 TFLOP/s; at S32768/H24, fast takes 4.400448 ms.Source: Table 7 · p. 14, §6.1 · p. 14
• HAO NV/FP8 has mean cosine 0.9899, versus 0.943789 for fast and 0.951669 for accurate; TK mean relative-L2 is 0.336602 and 0.327225, respectively.HAO does not publish relative-L2 for these rows.Source: §6.1 · p. 14, Table 7 · p. 14

6.2 Speed–accuracy behavior across shapes

• At S8192/H64, fast takes 0.758336 ms versus 1.611488 ms for BF16 (2.125×); in this suite, local TK NV/FP8 reaches 1.118× BF16 versus 1.783× for NV/MX fast.Figure 3 notes that the fastest full-FP4 points accept more operator error than HAO NV/NV or FP8 PV.Source: §6.2 · p. 15, Figure 3 · p. 15

6.3 ViT and BERT fixed-input evaluations

• On ViT S4096, fast matches BF16 top-1 accuracy at 88.5% with 95.5% prediction agreement; across 2,272 classification examples, it changes 32 predictions.Thirty-one changed predictions fall in the lowest quartile of BF16 top-two logit margins; the suite does not establish universal inference or training safety.Source: §6.3 · pp. 15, 17, Table 8 · p. 17

6.4 Wan video diffusion

• Fast is 2.09× faster than CuTe BF16 on the 14-billion-parameter Wan model, but its 20-step output shows more drift than HAO NV/NV.Both 14-billion-parameter twenty-step runs complete all 1,600 attention calls without a scan or stable-softmax fallback; the results use paired BF16 outputs.Source: Table 9 · p. 17, §6.4 · p. 18

6.5 Paired image reconstruction

• After all twelve ViT-MAE encoder layers, fast output remains 0.99973 cosine and 2.03% relative-L2 from BF16, despite the largest reconstruction displacement among listed routes.The paired PSNR comparison is FP4 minus BF16; Figure 5 amplifies the residual eightfold for visualization.Source: §6.5 · p. 18, Table 10 · p. 18, Figure 5 · p. 19

7.1 Pass the forward quantization into backward

• Backward reconstructs probabilities from saved quantized Q/K payloads, scales, and LSE, while projection epilogues publish FP8 layouts for gradient products.The represented values are exact with respect to the quantized forward payload, not the original unquantized BF16 inputs.Source: §7.1 · p. 20

7.2 How the backward pass works

• Probability reconstruction and dP must finish before dQ and dK; the implementation reuses rounded P, publishes both dS layouts, and uses E5M2 dO.E5M2 reduced the observed dO zero fraction from roughly 97% with E4M3 to roughly 14%, with less than 1% publisher overhead.Source: §7.2 · pp. 21–22, Figure 6 · p. 21

7.3 Isolated backward

• The isolated backward reconstruction core saves 29% of BF16 backward time, but the training-safe dO and row-statistic publisher consumes almost the same saving.Paired projection-inclusive operator checks report output and input-gradient cosine of 0.998 and 0.956; these are not convergence results.Source: §7.3 · p. 22

7.4 Combined forward and backward

• The projection-inclusive boundary times the attention sublayer with projection, RoPE, layout publication, attention, output projection, and gradients, excluding RMSNorm, residual, loss, and optimizer.BF16 and quantized paths use identical projection weights and the same B1/S4096/D128 shape.Source: §7.4 · p. 23, Table 13 · p. 23, Figure 8 · p. 24

7.5 Eight-billion-parameter end-to-end result

• Up to 1.14× complete-update speedup is reached at B4; FP8 and MXFP4 are tied end to end, despite MXFP4’s modest forward-only timing edge.At B4, FP8 throughput rises from 19,173 to 21,795 tokens/s per GPU; the complete-update comparison is performance-only, not evidence of training quality.Source: Table 14 · p. 25, §7.5 · p. 25

7.6 Training stability selects FP8 P/V

• Every tested MXFP4 P/V training trajectory diverges, while both FP8-P/V arms remain non-divergent through their common 55.5-billion-token horizon.In a matched B4 launch, MXFP4 loss rises from 7.01 at update 325 to 16.25 at update 350.Source: §7.6 · p. 25

7.7 Matched distributed training

• 1.112× median throughput for the FP8 route versus BF16 across 874 logged coordinates; the single-route trajectories complete the 100-billion-token schedule but differ in final validation loss.At the last common validation update, losses are 2.3048 for BF16 and 2.3948 for FP8, a 0.0900 gap; one trajectory per route does not estimate run-to-run variability.Source: Figure 10 · p. 27, Figure 11 · p. 28, §7.7 · p. 28

8.1 Takeaway 1: TMEM capacity and ownership bound the overlap window

• Under the current D128 layout, two score banks and two output accumulators occupy all 512 TMEM columns, preventing the next QK from using a legal destination until PV releases the overlay.The other query stage still provides partial overlap; reducing shared memory did not improve latency while the TMEM allocation still allowed one CTA per SM.Source: §8.1 · p. 29, Figure 12 · p. 29

8.2 Takeaway 2: low tensor activity is a readiness symptom

• Profiling attributes 1,715 of 2,669 not-issued samples to dependency stalls; tensor-core underutilization reflects operand readiness, not a lack of nominal FP4 throughput.Even fixed-P saves only 5.23%, with score loads, K64 publication, scale delivery, correction, accumulation, and epilogue remaining.Source: §8.2 · p. 30, Table 17 · p. 30, Table 16 · p. 30

8.3 Takeaway 3: faster arithmetic does not scale the whole kernel

• B300’s faster arithmetic does not scale the whole kernel: at S32768/H24 it reaches 2945 TFLOP/s versus 2998 on GB200, and other shapes depend on launch geometry.Per-SM-clock normalization favoring B300 is an inference, not an observed GPU-throughput win; the tested B300 does not increase TMEM capacity.Source: §8.3 · pp. 30–31, Table 18 · p. 31

8.4 Implications for hardware and kernels

• The proposed next steps are hypotheses, not measured wins; another allocatable score bank must preserve two-query K/V reuse and legal issue semantics.A naive ping-pong prototype doubled query jobs and was 52.8% slower; moving scales alone did not create a full score destination or demonstrate speedup.Source: §8.4 · p. 31

8.5 Scope and limitations

• Noncausal model results are fixed-input inference, and hardware conclusions are shape-specific; other head dimensions require their own tile and ownership strategy.One distributed trajectory per route does not estimate run-to-run variability, and the historical format controls do not provide a paired BF16 quality comparison.Source: §8.5 · pp. 31–32

8.6 Conclusion

• Direct-P shortens the score-to-value serial path, while causal training reuses quantized forward state and reaches up to 1.14× speedup for a single-GPU 8B update; distributed controls select FP8 P/V.Every tested MXFP4 P/V trajectory diverged; faster matrix arithmetic alone does not resolve the overlap limit at D128.Source: §8.6 · p. 32

Generative-AI Disclosure

• The author reports using generative-AI tools for code, experiments, analysis, figures, and manuscript work, then reviewing the material and verifying measurements.The author takes responsibility for the paper.Source: Generative-AI Disclosure · p. 32

A Historical NV/NV Investigation

• The historical NV/NV investigation is separated from the principal result because its kernels use different P scales, approximation policies, and sometimes denominators.Those timings must not be mixed with final NV/MX headroom.Source: A Historical NV/NV Investigation · p. 33

A.1 Initial stabilized path

• At S4096/H24, the stabilized path measured about 0.1884 ms after computing floating probabilities, reducing block maxima, encoding scales, and packing E2M1.Its scale selection waited for the floating probabilities, making the sequence conservative and serial.Source: §A.1 · p. 34
• A global NV encode factor G cancels when applied consistently to numerator and denominator, but rounding, underflow, and saturation prevent it from serving as a free accuracy knob.Sweeping G alone gave little benefit; full stabilization restored the expensive path.Source: §A.1 · p. 34

A.2 Code-directed polynomial development

• The code-directed approach fit base-two E2M1 threshold locations and targeted QE2M1(f(x)), yielding a faster and more accurate approximation than full exponentials and early linear fits.The cubic development also used packed Horner operations; an intermediate hybrid mixed native EX2 samples with cubic or affine pairs, coupling numerator and denominator sampling.Source: §A.2 · p. 34
• The later affine endpoint reduced non-native work to one packed FMA per pair, a mechanism retained in NV/MX with policy-specific coefficients and an exact represented denominator.The provided passage does not state the displayed equation's relation to these later code-directed steps.Source: §A.2 · p. 34

A.3 Sampled denominator

• Sampling eight of 32 probabilities saved a denominator pass but could miss concentrated maxima or misestimate tails; the retained NV/MX path sums all represented E2M1 payload words.Two- and three-word reductions reached respectively 0% and 75% top-1 agreement with BF16 in a 20-image ViT gate.Source: §A.3 · p. 35

A.4 Policy ladder and distribution dependence

• Fast NV/NV could look accurate on Gaussian scores yet produce non-finite or distorted model outputs when row-reference assumptions changed, while stabilization consumed the speedup.Fixed-scale sweeps did not reproduce the accuracy of the stabilized E4M3 control.Source: §A.4 · p. 35

A.5 Why NV/NV was not retained

• NV/NV was superseded because its E4M3 block scales could underflow, independent block-16 scale work fit the N32 producer poorly, and robust stabilization was too slow.The HAO-derived 512-column layout and MXFP4 scale handoff offered a cleaner way to retain small P blocks without another score bank.Source: §A.5 · p. 35

A.6 B300 D64 sweep and range failure

• 1.124–1.284× faster than matched B300 BF16 across S4096-and-longer rows, retained NV/MX stayed finite across all 24 D64 shapes while raw shiftless NV/NV was finite on only 23.The NV/NV control could be more locally accurate, but the retained route was chosen for finiteness across the matrix.Source: §A.6 · p. 36, Table 19 · p. 36
• At H32/S32768, one seed produced 64 non-finite raw NV/NV outputs while BF16 and NV/MX stayed finite; another seed passed.Saturating the E4M3 scale encoder repaired the failing distribution at 6.670336 ms and 1.186× BF16, with 0.962921 cosine and 0.275742 relative-L2; bounded mode cost roughly 2–3%.Source: §A.6 · p. 36, Table 19 · p. 36

A.7 Measured downstream failure mechanism

• On model activations, the shiftless scale quantity exceeded E4M3’s finite maximum of 448 in every tested workload, contaminating P and the denominator path without a satfinite guard.With row-max stabilization, small blocks could still underflow but carried negligible probability mass relative to the row anchor; HAO remained finite on every full fixed-input evaluation.Source: §A.7 · pp. 36–37, Table 20 · p. 37

B.1 SM103 launch and EX2 controls

• Table 21 is a tuning control rather than promoted policies: its corrected rows use 148 workers and vary EX2 density with the numerical contract held fixed.The initial 152-worker binary created a partial second wave on the tested 148-SM SKU.Source: §B.1 · p. 37, Table 21 · p. 37

B.2 Kernel directions

• The appendix preserves major rejected kernel directions as constraints on interpreting the retained method, while compile-time probes remain default-off.The table captions identify these as rejected directions, not promoted performance results.Source: §B.2 · p. 37, Table 22 · p. 38, Table 23 · p. 39

B.3 Quarter-local speed-of-light experiment

• 12.288 µs was saved by removing transforms from all quarters of the historical NV/NV path, but accelerating one N32 producer alone did not advance a K64 tensor command.The 12.288 µs total is not the current NV/MX ceiling; the paper reports a 2.336 µs raw-score-pack gap for that path.Source: §B.3 · p. 40

B.4 Scale-footprint experiments

• Smaller source scale tensors did not create free tensor-memory columns because tcgen05 still consumes an instruction-facing scale layout.The retained folded K64 Q/K representation reduces loads and arithmetic, but does not claim a new 128-column score slot.Source: §B.4 · p. 40

B.5 Joint affine-route search

• Fresh model processes rejected every proposed affine-route calibration change, despite reused-worker gains on fixed prompts.The joint optimizer did not improve throughput, and the observed numerical margins were below process-to-process variation; future calibration changes require fresh-process, unseen-prompt validation.Source: §B.5 · p. 40

C Accuracy-Matched FP8 Control

• 0.9897 cosine at 8.854 ms and 1490 TFLOP/s for exact TK NV/FP8 matched HAO’s published 0.9899 cosine, but took longer than HAO’s 4.929 ms and 2677 TFLOP/s.The intermediate optimized TK NV/FP8 route was faster but reached only 0.9573 cosine; matching HAO’s accuracy with exact TK FP8 gave back the retained NV/MX advantage.Source: C Accuracy-Matched FP8 Control · p. 41, Table 24 · p. 41

D Complete Fixed-Schedule Format Matrix

• The complete fixed-schedule format matrix compares Q/K/P/V formats inside forward-only attention, reporting latency and output error together against the same canonical BF16 reference.It does not include learned QKV/O projections or imply a pure-FP4 end-to-end training route.Source: D Complete Fixed-Schedule Format Matrix · p. 41, Table 25 · p. 42

E Reproduction and Compile-Time Contract

• The release contains kernels, configurations, evidence, and build scripts, while the measurement planner checks supplied artifacts and emits a build-and-run graph without downloading data, running CUDA, or submitting jobs.Historical measurements remain receipt-backed where original external asset revisions were not recorded.Source: E Reproduction and Compile-Time Contract · p. 43, §E.1 · p. 43

E.1 Canonical command graph

• The generated command graph requires new or empty output and build roots; blocked nodes are reported rather than replaced with guessed commands.The noncausal suite preallocates outputs, cycles provider order, and records every timing window.Source: §E.1 · p. 43

E.2 Fixed-input model evaluations; E.4 B300 aggregate reconstruction

• Historical Wan tables remain receipt-backed because some calibration and provider inputs are absent, and recomputing B300 aggregates from raw captures awaits a checksummed capture bundle.Fresh Wan runs can produce replacement evidence; the committed B300 summary can render tables, but raw-capture recomputation is not yet available.Source: §E.2 · p. 44, §E.4 · p. 44

G.1 Projection precision and attention precision are different

• In an isolated 8B block, E4M3 QKV projection had about 0.0375 relative-L2 error versus 0.1459 for NVFP4, while the native attention kernel added only 0.0216 relative-L2 on those represented operands.Attention-output error was already 0.3614 after exact BF16 attention on the decoded low-precision operands; projection precision is treated separately from attention precision.Source: §G.1 · p. 46

G.2 Two independent backward scale failures

• The backward scale fixes were independent: omitting +8 from the lifted LSE statistic reconstructed probabilities at 1/256 scale, while the gradient-output epilogue required its own 1/256 correction.The passage identifies the second correction as a separate matrix-layout convention.Source: §G.2 · p. 46
• Publishing dO in E5M2 reduced its observed zero fraction from about 97.1% with E4M3 to about 14.4%, at a 0.7–0.9% matched-gate slowdown.The wider exponent range motivated using distinct FP8 formats rather than treating all 8-bit values as interchangeable.Source: §G.2 · pp. 46–47

G.3 How the native backward schedule evolved

• Performance counters and clearing requirements cited for predecessor backward schedules are not evidence of speed for the retained v509 route, which always clears outputs.The predecessor profile reported 21% warp occupancy, 30% tensor activity, and 9% external-memory activity despite 95% SM activity.Source: §G.3 · p. 47

G.4 Historical isolated backward timings

• The longest matched BF16 sequence sweep used a generated CuTe kernel, and the D128 comparison isolated scheduling improvement rather than comparing low precision with BF16.The cited historical timings concern predecessor or generated schedules, not the retained implementation.Source: §G.4 · pp. 47–48, Table 27 · p. 48, Table 28 · p. 48

G.5 Why MXFP4 V in backward was not retained

• The retained training route keeps E4M3 V views for backward and selects FP8 P/V in forward because the exact MX consumer was too slow and the complete approximate-consumer recipe failed the observed distributed numerical gate.The shared MX route was only 0.194% faster at the median and 0.102% faster in sustained time, while the full MX update was 1.978 ms slower than the FP8 update.Source: §G.5 · p. 48, §G.6 · p. 49

G.6 Historical local speed brackets

• Historical local speed brackets used older projection formats and cut-cross-entropy, so they are not timings of the current dense-cross-entropy recipe.The retained comparison keeps projection formats and loss boundary explicit.Source: §G.6 · p. 49, Table 29 · p. 49

G.7 Historical training context

• At 55.5 billion tokens, one-billion-token EMA losses were 2.534 (E4M3 projections with FP8 P/V), 2.542 (NVFP4 projections with FP8 P/V), and 2.756 (historical BF16), but the runs are not comparable.The BF16 run used a different topology, global batch, sample order, and history-sampling density.Source: §G.7 · p. 49, Figure 13 · p. 50

G.8 What the distributed failures established

• The first long MX run did not isolate a cause: at a matched checkpoint its loss was 7.0805 versus 6.5887 for FP8, with gradient norms of 157,696 versus 6.41.Its gradient norm later exceeded 1.5 million despite clipping.Source: §G.8 · p. 50
• At update 500 (131.1 million tokens), both MXFP4 arms reached loss 7.97 versus 5.41 for their FP8 controls, separating by forward P/V format.This is strong factorial evidence, but not proof that one kernel instruction causes divergence.Source: §G.8 · p. 50
• In a later B4 campaign, both MXFP4-P/V runs diverged while BF16 and NVFP4-projection/FP8-P/V remained stable; the tests do not identify which MXFP4 stage is responsible.The E4M3-projection MXFP4-P/V run was cancelled at update 2550 after repeated finite gradient explosions.Source: §G.8 · p. 51, Figure 14 · p. 52
• At the 55.5-billion-token horizon, E4M3+MX and NVFP4+MX losses were 8.94 and 8.18, versus 2.53 and 2.54 for FP8 controls; MX gradient norms repeatedly reached millions.The earlier retained-observation markers were truncation markers, not estimates of when divergence began.Source: §G.8 · p. 51, Figure 15 · p. 53

G.9 Evidence boundary

• FP8 stability is observed only through the stated horizon; one trajectory per route cannot establish run-to-run variability or statistical equivalence, and historical MX throughput is not a matched speedup.The historical rows were route-selection diagnostics, not checkpoint-aligned validation or aggregate throughput for the retained recipe.Source: §G.9 · p. 54, H Causal-Training Result Provenance · p. 56