Four Bodies, Four Precision Policies

Why IDA Stopped Treating FP8 as a Universal Answer

TL;DR: A correct FP8 implementation is not automatically the correct production policy for every model body. In IDA's current native trainer, family-specific re-ablation found that Edge and Swift run faster with BF16 linears, while AI and MoE run faster with FP8 linears. The result changed production defaults and exposed a real launcher bug that had silently forced Edge onto AI's FP8 policy.

Low-precision training first presented itself as a correctness problem.

Then it became a policy problem.

In "When Activation Checkpointing Calls a Stateful Quantizer Twice" (https://huggingface.co/blog/KissTheHabit/fp8-fp4-delayed-scaling-recompute-bug), I described a failure involving delayed-scaling FP8 and packed-FP4 attention under activation-checkpoint recomputation.

The quantizer was stateful. The original forward pass could use one delayed scale and advance the live history. Backward recomputation could then re-enter the same quantizer, observe already-advanced state, and advance it again for the same logical training step.

The repair required:

• independently owned persistent state;
• pre-mutation scale and descale snapshots;
• immutable metadata paired with each quantized tensor;
• and a hard rule that recomputation may reproduce arithmetic but may not reproduce state transitions.

That established one boundary:

«A faster numerical format is not useful unless its state remains correct under the real training lifecycle.»

Once that boundary was enforced, however, a second assumption survived longer than it should have.

The training system still treated precision too much like a global engine switch.

FP8 was enabled or disabled. BF16 was the baseline or fallback. A precision profile appeared to describe the trainer as a whole.

That framing did not survive contact with four different computational bodies.

IDA currently trains four model families:

• Edge
• AI
• Swift
• MoE

They share institutional lineage, curriculum structure, operational tooling, and a common native training engine.

They do not share the same dimensions, attention-kernel compatibility, memory behavior, deployment purpose, or ratio between useful matrix multiplication and quantization overhead.

When all four families were measured independently, the result was not one universal winner.

It was four body-specific precision policies.

---

The result in one figure

Observed throughput advantage of the selected linear format within each family's matched comparison window

Edge BF16 +6–8% ███████ AI FP8 +9–11% ██████████ Swift BF16 +8–11% █████████ MoE FP8 +1–6% ████

Bar lengths are illustrative mid-range summaries. The reported percentage ranges are the evidence-bearing values.

Family| Required attention backend| Selected linear precision| Observed throughput advantage Edge| "hopper_wgmma_packed_fp4"| BF16| approximately 6–8% AI| "hopper_wgmma_packed_fp4"| FP8| approximately 9–11% Swift| "scalar_flash"| BF16| approximately 8–11% MoE| "scalar_flash"| FP8| approximately 1–6%, increasing at deeper accumulation tiers

The conclusion is not that BF16 is always better for small models or that FP8 is always better for large models.

The defensible conclusion is narrower:

«In this trainer, on these bodies, under the measured schedules, Edge and Swift benefit from BF16 linears, while AI and MoE benefit from FP8 linears.»

That is the level at which the production policy now lives.

---

How the comparison was made

The family-specific re-ablation was recorded on July 22, 2026, against the active Hopper training path.

The public-safe methodology was:

1. Hold the body fixed. Each comparison stayed within one family architecture. Edge was compared with Edge, AI with AI, Swift with Swift, and MoE with MoE.

2. Hold the required attention path fixed. The experiment did not pretend every family could use the same attention kernel. Edge and AI remained on the Hopper WGMMA packed-FP4 path. Swift and MoE remained on the shape-agnostic scalar-flash path.

3. Change the linear precision profile. The relevant comparison was FP8 linears versus BF16 linears inside the family's otherwise compatible execution path.

4. Compare matched work states. Throughput was compared at corresponding optimizer-step and gradient-accumulation tiers. The decision was not made from one cumulative tokens-per-second average spanning unlike portions of the curriculum ramp.

5. Require a valid training trajectory. A faster path was not accepted merely because it produced more tokens per second. The run still had to remain numerically usable rather than converting throughput into skipped or invalid optimization work.

6. Promote only after resolving the production path. The selected profiles were written into the family-aware launcher and then checked on the training server to confirm the resolved runtime policy matched the source decision.

The public record supports the result ranges and the matched-tier method. It does not provide enough public-safe raw windows to claim statistical confidence intervals, universal scaling laws, or transfer to different hardware generations.

That limitation is intentional.

This is a production engineering result, not a claim that four measurements define all low-precision training.

---

Why one precision setting failed

FP8 is not free.

Before lower-precision matrix multiplication can provide an advantage, the system must perform work around it:

• determine or retrieve scale metadata;
• quantize higher-precision values;
• convert or pack data;
• preserve paired scale and descale state;
• and move the resulting representation through a compatible kernel path.

Some of this overhead can be fused.

Some can be amortized.

Some remains real work.

For a body with sufficiently large linear operations, FP8's execution advantage can dominate that preparation cost.

For a smaller body, the quantization machinery can consume the margin that FP8 was supposed to create.

That is what the four-family comparison exposed.

AI

The AI body has enough linear-layer work for FP8 to produce a clear gain.

The current re-ablation measured FP8 ahead by approximately 9–11%.

This reversed an older AI result that had favored BF16. The engine had changed since that earlier measurement, including later fusion work around the BF16-to-FP8 snapshot path. The old result remained historically valid for its engine state, but it was no longer entitled to govern the current one without being measured again.

MoE

The MoE body also benefits from FP8.

Its measured advantage is smaller—approximately 1–6%—but the margin increases at deeper accumulation tiers.

That pattern matters. One cumulative average could make MoE's gain look negligible while hiding that the format becomes more favorable as the schedule gives it more useful work over which to amortize conversion cost.

Edge

Edge behaves differently.

Its useful GEMM share is smaller, so the time saved by FP8 execution does not offset the conversion and state-management overhead.

BF16 wins by approximately 6–8%.

Swift

Swift shows the same broad small-body pattern, although it uses a different attention backend from Edge.

BF16 wins by approximately 8–11%.

The important distinction is:

«“The engine supports FP8” and “this body should use FP8” are different claims.»

The engine supplies capabilities.

The body policy decides whether using one of those capabilities is rational.

---

Attention compatibility is a separate axis

Precision choice is not the same decision as attention-kernel compatibility.

Edge and AI can use the Hopper WGMMA packed-FP4 attention path because their attention head dimensions fit the kernel's supported shapes.

Swift and MoE cannot.

The current WGMMA implementation requires:

head_dim ∈ {64, 256}

Swift's real head dimension is:

head_dim = 32

MoE's real head dimension is:

head_dim = 512

Those values are not configuration mistakes that should be normalized away.

They follow from what the bodies are for.

Swift is shaped for constrained and portable deployment. MoE carries a larger hidden body with sparse expert routing. Both therefore use the shape-agnostic "scalar_flash" attention path.

Within that shared attention constraint, their linear precision policies still diverge:

Swift: scalar_flash attention BF16 linears

MoE: scalar_flash attention FP8 linears

The same separation appears on the WGMMA path:

Edge: WGMMA packed-FP4 attention BF16 linears

AI: WGMMA packed-FP4 attention FP8 linears

This is why a single global “low-precision mode” is inadequate.

Attention compatibility, linear precision, optimizer policy, activation-memory strategy, and deployment purpose are related.

They are not interchangeable.

Two bodies may require the same attention backend and still need opposite linear formats.

---

The production bug hidden behind the correct comment

The re-ablation did more than change a performance recommendation.

It exposed a real production bug.

The launcher contained a comment explaining that precision profiles should be split by family class.

The executable branch directly below it did not preserve that split.

Edge was grouped into AI's FP8 branch.

The BF16 path intended for Edge was unreachable dead code.

Nothing necessarily crashed.

The model could still train. Status files could still update. Throughput could still appear plausible. The system simply ran Edge with a numerical policy that was slower than the measured alternative.

This class of failure is dangerous because the explanation and the implementation appear to agree until the actual branch is traced.

The comment described the desired system.

The runtime implemented something else.

After the defect was identified, the production defaults were corrected:

Edge → hopper_bf16_packed_fp4 AI → hopper_fp8_packed_fp4 Swift → legacy_bf16 MoE → legacy_fp8

The family-resolved profiles were then checked directly on the training server.

That last step matters.

A source-code fix is not yet production evidence.

A production claim should remain traceable through at least four states:

requested policy → accepted configuration → observed runtime configuration → persisted checkpoint/evidence configuration

The same drift can affect:

• microbatch size;
• gradient accumulation;
• optimizer type;
• attention backend;
• precision profile;
• dataset identity;
• tokenizer identity;
• and parent-checkpoint lineage.

A comment is not an execution contract.

---

Precision became part of the body contract

The deeper change is architectural.

IDA's four families are not merely different-sized exports of one universal model.

They are different computational bodies carrying related institutional roles into different operating environments.

• Edge prioritizes small-footprint execution.
• Swift is shaped for constrained and portable deployment.
• AI carries a larger dense body.
• MoE distributes capacity through sparse expert routing.

A body contract therefore has to specify more than:

• hidden size;
• layer count;
• head count;
• vocabulary;
• and checkpoint path.

It must also resolve how that body inhabits the training engine:

• which attention kernels accept its shape;
• whether FP8 conversion cost is justified;
• how delayed numerical state is owned;
• how activation memory is handled;
• how optimizer state is represented;
• which checkpoint and resume rules apply;
• which telemetry constitutes valid evidence;
• and which failures invalidate the run.

Precision is not an acceleration switch applied after architecture selection.

It is part of the architecture's relationship with the hardware.

---

Performance evidence has a software-version boundary

The AI result demonstrates another rule.

An earlier experiment had favored BF16 for AI. That result became the production default.

The trainer did not remain static after that experiment.

Later work changed the numerical path. Once conversion and snapshot behavior changed, the balance between FP8 overhead and FP8 execution changed with it.

The new re-ablation reversed the decision:

AI: old measured winner → BF16 current measured winner → FP8 by approximately 9–11%

This does not make the earlier experiment dishonest.

It makes it version-bound.

A performance result belongs to a specific combination of:

• model shape;
• kernel implementation;
• precision implementation;
• batch and accumulation schedule;
• hardware generation;
• orchestration state;
• and trainer revision.

When one of those changes materially, the conclusion may need to be earned again.

The correct response is not to protect the old result.

It is to preserve it as historical evidence and rerun the decision under the current system.

This is especially important in a custom CUDA trainer, where one fused conversion kernel can change which cost dominates the step.

---

Why matched-step measurement matters

Training schedules are not uniform.

Gradient accumulation, sequence handling, memory pressure, checkpoint cadence, and curriculum stages can create different throughput regimes during one burn.

A single cumulative average combines those regimes into one number.

That can hide the actual comparison.

Suppose FP8 and BF16 are observed during different portions of a ramp:

candidate A: measured mostly during a shallow accumulation tier

candidate B: measured mostly during a deeper accumulation tier

The resulting cumulative throughput values are not directly comparable, even if they came from nominally similar runs.

The family re-ablation instead matched corresponding work states.

That is particularly important for MoE, where the FP8 advantage grew with accumulation depth.

The methodological rule is:

«Compare numerical formats at equivalent work states, not merely inside the same wall-clock period.»

The unit of comparison is not “one run.”

It is the matched region of the run in which both candidates performed equivalent scheduled work.

---

What this changes for distributed training

The next infrastructure phase extends the native trainer beyond one GPU process and eventually beyond one machine.

The current project direction is synchronous distributed execution with:

• one native process per GPU;
• fixed rank assignment;
• globally partitioned data;
• NCCL collectives;
• rank-0 checkpoint authority;
• whole-job fail-stop behavior when a required rank disappears;
• distributed checkpoint and exact-resume work;
• and later Blackwell and GPUDirect RDMA validation.

That work is not presented here as completed capability.

The current precision result still changes how that system must be designed.

A distributed launcher cannot apply one numerical profile to every family merely because every rank participates in the same collective system.

The body contract must be resolved before rank creation.

Every rank training the same body must agree on:

attention backend precision profile optimizer policy accumulation schedule dataset partition contract checkpoint identity

That resolved contract becomes part of the distributed job identity.

Otherwise, more GPUs would multiply ambiguity rather than throughput.

Distributed execution is not simply the creation of more processes.

It is the reproduction of one governed training state across more hardware.

---

What this result does not establish

These measurements do not establish that:

• FP8 is always best for large models;
• BF16 is always best for edge models;
• MoE architectures inherently prefer FP8;
• the current result transfers unchanged to Blackwell;
• Hopper results automatically predict multi-node behavior;
• or one precision policy should remain fixed as the trainer evolves.

They establish a current production policy for four current IDA bodies on the current native engine.

That policy is useful because it is:

• narrow;
• executable;
• measured;
• falsifiable;
• and allowed to change.

A future kernel revision may reverse one of the results.

A different sequence length may reverse one of the results.

A Blackwell port may reverse one of the results.

When that happens, the policy should change with the evidence.

---

The larger lesson

Custom training engines often expose precision as a global command-line argument:

--precision fp8

That is convenient.

It is also increasingly inadequate.

Once one engine trains several architectural families, a global precision switch collapses decisions that belong at different layers:

• hardware capability;
• kernel compatibility;
• model shape;
• quantization overhead;
• memory pressure;
• delayed-state correctness;
• training schedule;
• and production evidence.

IDA now treats those as a resolved body policy rather than an operator guess.

The result is less visually simple than one universal FP8 mode.

It is operationally more honest.

Edge uses BF16 because BF16 is faster for Edge.

AI uses FP8 because FP8 is faster for AI.

Swift uses BF16 because BF16 is faster for Swift.

MoE uses FP8 because FP8 is faster for MoE.

No family is required to imitate another family's relationship with the hardware.

That is the actual conclusion:

«A shared training engine does not require a shared numerical identity.»

The engine provides capabilities.

The body contract decides which capabilities should govern the run.

---

Disclosure

This article reports public-safe findings from a private native CUDA/C++ training project.

It intentionally omits:

• proprietary repository paths;
• private configuration files;
• unreleased model weights;
• complete raw ablation logs;
• infrastructure credentials;
• and implementation details that would expose private operational systems.

The reported percentage ranges come from the July 22, 2026 family-specific precision re-ablation and the production policy change that followed it.

No claim is made that the result transfers unchanged to another trainer, another model family, another accelerator generation, or another training schedule.

---

Related IDA research

• "When Activation Checkpointing Calls a Stateful Quantizer Twice" (https://huggingface.co/blog/KissTheHabit/fp8-fp4-delayed-scaling-recompute-bug) Delayed-state correctness, recomputation invariants, and independently owned packed low-bit state.

• "The Kernel Was Not the Whole Bottleneck" (https://huggingface.co/blog/KissTheHabit/the-kernel-was-not-the-whole-bottleneck) Persistent worker state, orchestration repair, Hopper profiling, and why end-to-end performance is a property of the complete execution system.

• "The Economics of Precision" (https://huggingface.co/blog/KissTheHabit/the-economics-of-precision) Why lower-bit arithmetic matters only when it reduces the complete cost of accepted, useful capability.

• "Beyond the Model Checkpoint: Building AI as a Governed Institution" (https://huggingface.co/blog/KissTheHabit/beyond-the-model-checkpoint) The larger architectural thesis behind differentiated bodies, persistent institutional roles, and governed succession.