fmt=7: fp8-e4m3-b128 passthrough weight format — CPU + Metal + repack tool (#524)

[monotophic]

Authored by Fable 5 in Claude Code (We/our), analysis in partnership with @monotophic (me/my).

Revised per review (2026-07-23): collision policy inverted; see the comment of this date for the full revision + evidence.

fmt=7: native FP8-e4m3 passthrough — CPU path, Metal kernel, repack tool

Branch: fmt7/fp8-passthrough. Base: dev 9baae9b. Closes the FP8-passthrough
Feature Request; implements the format the maintainer assigned fmt=7 on
#524. Companion PR (fmt7/stamp-registry, stacked on this one) adds the
container metadata stamp and the docs/FORMATS.md registry — split out per
the maintainer's own scoping request, so this PR is CPU path + repack tool +
collision/refusal logic + the Metal kernel only.

What / why

Z.ai ships GLM-5.2's release-native checkpoint in FP8-e4m3
(zai-org/GLM-5.2-FP8) — to our knowledge there is no public bf16/fp16
upstream of it. Every existing path into this engine from that checkpoint
first dequantizes to f32 and requantizes to a lossy format (int4/int3),
paying a real second quantization error on top of whatever the FP8 already
cost. This PR adds a different path for tensors where the extra fidelity is
worth the extra 8 bits/weight over int4: repack, not requantize. The
repack tool copies the fp8 weight bytes AS-IS (.view(torch.uint8), a pure
bit-reinterpret, verified byte-identical against torch.float8_e4m3fn's own
storage) and only renames/reshapes the weight_scale_inv sidecar to the
engine's .qs convention. Same 8 bits/weight streaming cost as the FP8
source, zero additional quantization loss — a resident/quality-core tier
sitting above int4/int3 in fidelity, below f32 in size.

Selection: which tensor roles get fmt=7

tools/repack_fp8_passthrough.py reuses convert_fp8_to_int4.classify() —
the engine's existing tensor-role taxonomy — rather than inventing a new one:

role (classify() kind)	selected for fmt=7?	why
shared expert (sh)	yes	resident, always-active — fidelity matters every token
o_proj (o)	yes	resident attention output projection
other attention projections (attn, incl. q_a_proj)	yes	resident
dense-MLP first layers (dmlp)	yes	resident
generic resident fallback (q)	yes	catches anything else resident-shaped
kv_b_proj (kvb)	no	resident, but no fmt=7 case in the CPU/CUDA MLA-absorb path — see below
routed experts (x)	no	stays on the existing int4-g64 streaming path (58k+ tensors — the streaming-cost tradeoff is different for a per-token top-K subset)
routers/norms/biases (f32)	no	never FP8 in the source, always f32
embed/lm_head (io)	no	BF16 in source, never FP8

Blending selected fmt=7 residents with int4-g64 routed experts in the same
loadout directory is separate, deferred "loadout index-rewrite blending"
work — out of scope here; the tool only ever emits a fmt=7 container for
the roles in the table.

Design

Decode LUT (quant.h's E4M3_LUT, 256 entries): a static compile-time
table, cross-checked byte-for-byte against torch.float8_e4m3fn for all 256
codes including both signs of zero, subnormals, and the OCP E4M3-FN NaN
encoding (exp==0xF && mant==0x7 only — exp==0xF is NOT reserved for
infinity in this format). Static rather than lazily initialized because
matmul_fp8 runs inside #pragma omp parallel for and a lazy-init global
would race under concurrent first use. The Metal mm_gemv shader's fmt==7
branch does the identical bit-decode in-kernel (no LUT texture — BW-bound
kernel, decode ALU is free) and is verified to byte-match the CPU LUT via an
exhaustive 256/256-code GPU test (see Metal validation below).

Block scales: one f32 per 128×128 block of the [O,I] weight matrix
(ceil(O/128)*ceil(I/128) entries, row-major block order), not one per
output row. Dequant is w[o,i] = e4m3_decode(byte) * scale[o/128, i/128]
(multiply, not divide) — matches tools/convert_fp8_to_int4.py's own
dequant(), the authoritative reference for how Z.ai's checkpoints read.

Scale encoding is a declared property of the format, not a hardcoded
constant. This was an open design question we'd have gotten wrong by
defaulting to "f32, full stop" — the maintainer's own finding, surfaced on
#524, is what settled it: DeepSeek-V4 ships this identical weight geometry
(FP8 E4M3, 128×128 blocks) with UE8M0 scales instead of f32 — one byte per
block, a power-of-two exponent, dtype F8_E8M0. Concretely, in this PR:

• fmt=7's scale encoding is named explicitly in the QT struct comment and
  qt_resolve_fmt's own documentation; f32 is the implemented value
  (what Z.ai's GLM-5.2-FP8 checkpoints ship, and what this PR's CPU kernel,
  Metal kernel, and repack tool all read/write).
• The resolver recognizes the UE8M0 byte signature and refuses by name
  rather than misreading it: the two encodings differ cleanly in
  scale-sidecar arithmetic (ceil(O/128)*ceil(I/128)*4 bytes for f32 vs
  *1 for UE8M0), a real, distinct discriminator (never coincidentally
  equal for any O,I >= 1, since one is exactly 4× the other). An
  unstamped container with a 1-byte-per-block sidecar gets "fp8-e4m3-b128
  with ue8m0 scales recognized but not implemented; only f32 block scales
  are supported in this build" instead of garbage.
• UE8M0 decode itself is invited follow-up work, same posture this project
  already applies to the CUDA tier: whoever reaches DeepSeek-V4 first (that
  may be us, later) lands it into a seam that already exists rather than a
  near-duplicate format. The MXFP4 routed-expert question (E2M1 + UE8M0 g32)
  raised alongside this finding is a separate format and a separate
  conversation — but the same scale-encoding-as-property mechanism serves
  it too.

The collision/refusal analysis

qt_resolve_fmt infers a tensor's format from byte arithmetic alone — the
container carries no format ordinal. fmt=7's weight bytes are
byte-identical to fmt=1 (int8): O*I raw bytes either way. That creates
three genuine byte-collision shapes this PR's resolver must catch and
refuse rather than silently misread:

1. fmt=1 vs fmt=7 (THE DESIGN LANDMINE). For small O (≤128) and/or
   small I, ceil(O/128)*ceil(I/128) can equal O exactly (e.g. O=1, I≤128; O=2, I∈[129,256]; O=256, I∈(16256,16384]) — a scale array
   that satisfies BOTH fmt=1's per-row convention (O*4 bytes) and fmt=7's
   per-block convention (ceil(O/128)*ceil(I/128)*4 bytes) at once.
   Guessing either way silently corrupts every weight. Refused.
2. fmt=6 vs fmt=7 at I=98 (SECOND DESIGN LANDMINE). Upstream's own
   fmt=6 (E8/IQ3, engine: fmt=6 E8/IQ3 lattice container — CPU decode kernel and dispatch (#452 step 3) #465) tags itself with ns==4 && nb==O*e8_rowbytes(I),
   checked unconditionally before any fmt=7-aware logic runs.
   e8_rowbytes(I) = ceil(I/256)*98 is the constant 98 for every I in
   (0,256], so that check collapses to nb==O*98 — which is ALSO fp8's
   raw weight-byte count (O*I) at exactly I=98. At that same I=98, a
   single-block fp8 tensor (O≤128, f32 scales) carries exactly one
   f32 block scale (ns==4, the same tag fmt=6 uses); a four-block
   fp8 tensor (O∈(384,512], UE8M0 scales) also lands at ns==4 (4
   blocks × 1 byte/block). A straight, un-reconciled rebase onto post-engine: fmt=6 E8/IQ3 lattice container — CPU decode kernel and dispatch (#452 step 3) #465
   dev would silently misread either as E8/IQ3-lattice-decoded garbage —
   demonstrated fail-before during development (a standalone probe at
   O=64, I=98 returned fmt=6 with no error before the fix). Refused
   after the fix, for both the f32-scaled and UE8M0-scaled variants.
3. O==1 triple collision. At O=1, I=98, fmt=1's own per-row tag
   (O*4) is ALSO 4 — a genuine three-way ambiguity (fmt=1 / fmt=6 /
   fmt=7) at one shape. Refused.

This PR's refusal posture is unconditional: it carries no container
metadata stamp (that's the companion PR), so every one of these three
collisions refuses regardless of the tensor's true identity — "nothing real
is lost: no real checkpoint tensor has these shapes" (the maintainer's own
framing, #524). The companion PR's stamp is the only thing that can turn a
refusal into a resolution, and even then only for the f32-scaled candidates
— a stamp can confirm a WEIGHT format, it can't grant a decoder this build
doesn't have, so the UE8M0 corners of collisions 2 and 3 stay refused even
under a (correct) stamp.

Test coverage: tests/test_fp8_load.c's Part A (15 cases: golden paths, 6
degenerate collisions, 4 boundary-adjacent non-collisions) + Part A2 (the
fmt=6 collision at I=98, both boundary shapes, the three-way O=1
case, plus 2 non-colliding fmt=6 regression fixtures proving dev's own
E8/IQ3 inference stays unbroken) + Part A3 (2 UE8M0-signature refusal
cases, including the small-O collision variant).

kv_b_proj exclusion

kv_b_proj (kind kvb) is a valid resident tensor classify() would
otherwise select, and it's deliberately excluded anyway: colibri.c's
MLA-absorption CPU path (qt_addrow/qt_matvec_rows, the always-available
fallback whenever the Metal fused decode kernel isn't used) has no fmt=7
case and would silently misread fp8 bytes as int2-packed data; the CUDA
absorb kernels (coli_cuda_attention_absorb/_kvdev in backend_cuda.cu)
are similarly int4-specific. Repacking kv_b_proj to fmt=7 needs that
CPU+CUDA absorb-path work first — out of scope for this PR, so the tool
refuses to produce a container the engine cannot safely read
(tests/test_fp8_repack.py::test_kv_b_proj_excluded is the regression
guard).

Metal validation record

This is the evidence we'd ask you to merge on, absent Apple hardware on
your side — every fmt=7 code path exercised on real M-series Metal, not
just described:

• LUT-through-GPU exactness: 256/256 byte codes, run through the
  ACTUAL compiled mm_gemv shader (not a CPU-only spot check) — proves the
  in-kernel bit-decode agrees with quant.h's reference LUT for every
  possible byte value, including both zeros, all subnormals, and the NaN
  encoding.
• Adversarial shape suite: block edges (O,I both non-multiples of
  128), degenerate 1×1 (single sub-block), a non-square block grid
  (nblkO=3, nblkI=48 — a stride-audit case that would surface a
  transposed or swapped block index), worst observed relative error
  3.0e-8 to 7.0e-8 across all cases (tolerance 1e-4).
• S>1 gate test: the spec-shaped [2048,6144] case run at S=4
  (batched decode), not just S=1.
• Two independent negative-gate tests: coli_metal_gemm explicitly
  refuses fmt=7 (rc=0, CPU-fallback signal) — the large-batch sync GEMM
  path was never in scope for this format. coli_metal_moe_block/
  moe_submit's existing fmt != 1 && fmt != 2 allowlist already refuses
  fmt=7 before any pointer is dereferenced — proven BY TEST with
  deliberately-invalid 0xdeadbeef weight/scale pointers that must never
  be read, not left as an artifact of the allowlist happening not to
  include 7 (yet).
• Both COLI_METAL_RESSET states: make metal-test full green under
  COLI_METAL_RESSET=0 AND COLI_METAL_RESSET=1 (residency-set path
  confirmed active via "[METAL] residency-set: on"), all 11 fp8 GPU
  lines identical in both.

Full case-by-case output (RESSET=0; RESSET=1 is line-identical apart from
the residency-set banner):

Metal fmt=7 native FP8-e4m3 passthrough tests:
  fp8 LUT exactness (256/256 codes via GPU kernel) 0/256 mismatches  ok
  fp8 gate/up-shaped O=2048 I=6144 (spec example) S=1 worst_rel=3.01e-08 (I=6144 O=2048 nblkO=16 nblkI=48 S=1)  ok
  fp8 down-shaped O=6144 I=2048 S=1          worst_rel=3.63e-08 (I=2048 O=6144 nblkO=48 nblkI=16 S=1)  ok
  fp8 gate/up-shaped O=2048 I=6144 S=4       worst_rel=4.05e-08 (I=6144 O=2048 nblkO=16 nblkI=48 S=4)  ok
  fp8 block edges: O,I both non-mult-128     worst_rel=5.95e-08 (I=200 O=130 nblkO=2 nblkI=2 S=2)  ok
  fp8 block edges: O just over 128, I exact  worst_rel=4.65e-08 (I=128 O=129 nblkO=2 nblkI=1 S=1)  ok
  fp8 block edges: O exact, I just over 128  worst_rel=6.93e-08 (I=129 O=128 nblkO=1 nblkI=2 S=1)  ok
  fp8 degenerate 1x1 (single sub-block)      worst_rel=5.87e-08 (I=1 O=1 nblkO=1 nblkI=1 S=1)  ok
  fp8 non-square block grid nblkO=3 nblkI=48 (stride audit) worst_rel=3.33e-08 (I=6144 O=384 nblkO=3 nblkI=48 S=3)  ok
  fp8 GEMM entry explicitly gated off (coli_metal_gemm refuses) rc=0 (expect 0/CPU-fallback)  ok
  fp8 MB_BUILD/moe_submit entry gated off (shared-expert fmt=7 hazard) rc=0 (expect 0/CPU-fallback)  ok

Every pre-existing Metal case (int8/int4/int2/f32, moe_block, gemm S=64,
fused attention, top-8 select) unaffected, in the same run.

Durable vs. current-state (review map)

Per this project's own convention for delineating what's structural from
what has a foreseeable resolution:

	claim	durable / current-state	revisit trigger
1	fmt=7 weight-byte layout (O*I raw e4m3), E4M3 decode LUT, block-scale indexing	durable	none — fixed by the OCP E4M3-FN spec and this format's own definition
2	fmt=1-vs-fmt=7 and fmt=6-vs-fmt=7 collision/refusal logic	durable	none — the ambiguous shapes are structural (byte-arithmetic facts), not tunable
3	UE8M0 scale encoding: recognized, refused, not decoded	current-state	a UE8M0 CPU+Metal decoder landing into this seam (qt_resolve_fmt's fmt=1 candidate branch + a new dispatch case) — invited follow-up, may be us
4	kv_b_proj excluded from fmt=7 selection	current-state	CPU (qt_addrow/qt_matvec_rows) + CUDA MLA-absorb kernels gaining an fmt=7 case
5	"refuse unconditionally" at all three collision shapes (no stamp)	current-state, resolved by the companion PR	fmt7/stamp-registry merging — a stamped, unambiguous container resolves instead of refusing (except the UE8M0 corners, which stay current-state per row 3)
6	Metal validation record below	checkable coordinates	measured 2026-07-22, macOS 26.5.2 (build 25F84), Apple M5 Max (MacBook Pro, Mac17,7), engine base dev 9baae9b

Behavioral contract

This PR was revised against a written contract; every invariant below is a
registered claim (stated before the revision was built), each with its
decisive evidence:

Invariant	Capstone evidence
Every container that loads on dev loads identically here (resolutions, notices, output)	Differential matrix vs. current dev: byte-identical temp-0 generated text on the stock int8-MTP container and an int8-resident mint
No byte pattern readable today changes meaning; ambiguity favors the incumbent	[6144,16384] int8-row (GLM-5.2's real o_proj): refused by the previous head (exit 1, quoted in-thread), loads as fmt=1 here
The writer can never emit what the reader would mis-resolve	Repack refusal demonstrated at a real ambiguous shape ([2048,16384]); round-trip load verified on non-ambiguous output
fmt=7 decode is exact OCP E4M3-FN — passthrough, no approximation	Independently derived oracle (from the OFP8 spec text): bit-exact at all 16,900 positions of a partial-block (130×130) tensor
Failure is loud: every refusal names its condition	New: qt_addrow/qt_matvec_rows refuse unhandled formats by name (previously: SIGSEGV for fmt=7, silent misread for fmt=6 — found by the acceptance suite, fixed here)
Collision family is enumerated, not sampled	In-tree census over every real tensor shape: o_proj is the only ambiguous family (79 instances/container)

Full evidence matrix (instruments and results per requirement)

• Compatibility (I1): old-binary/new-binary differential over both real
  containers, per-run usage-state control, telemetry-excised output
  comparison; plus same-binary control run proving the method's
  discrimination. Result: identical loads, per-tensor resolutions, notices,
  and generated text, all four comparisons.
• Collision policy (E4/I2): reader inversion verified at source and
  behaviorally (collision-family fixtures at 4 real shapes + 1
  non-collision control resolve fmt=1); the three previously-refusing test
  shapes plus four degenerate cases now assert fmt=1; census tool
  regenerable upstream (c/tools/fp8_collision_census.py).
• Exact decode (I4): spec-blind acceptance suite with an independently
  derived E4M3-FN oracle (derivation caught an internal contradiction in
  the OFP8 spec's own bias statement, resolved via its worked examples);
  bit-exact on NaN-free data; NaN weights verified to poison their row
  through the dot product (IEEE-correct propagation, 130/130 rows).
• Wire-path sizing: qt_scale_bytes() wired via shared
  qt_wire_split(); seam test intercepts actual mlock/munlock calls —
  a sites-only regression to O*4 fails 4 named assertions; fmt=6's
  4-byte-tag case corrected while making the helper authoritative.
• Loud failure (I6): refusal tests for fmt=6/7 through both absorb-path
  functions (previously crash/silent); repack tool refuses on zero
  repack-target tensors (previously silent success).
• Gates: make check; METAL=1 zero warnings; metal-test under both
  COLI_METAL_RESSET states; per-commit bisect builds; suppression-
  stripped -Wunused-function build; POSIX-construct check (Windows CI
  pattern) on all new tests.

Deeper artifacts (specific transcripts, census tables, suite design)
available on request.

Commits in this PR

1. fp8: CPU read path + matmul_fp8, as fmt=7 (public ordinal, #524) —
   quant.h's LUT/decode/kernel, colibri.c's byte-accounting and THE DESIGN
   LANDMINE refusal, the UE8M0 recognition seam, qt_resolve_fmt in its
   plain 6-argument (no stamp) form.
2. fp8: Metal mm_gemv kernel for fmt=7 native FP8-e4m3 passthrough — the
   GPU kernel and its full adversarial test suite, including the
   moe_submit gate regression test.
3. fp8: repack tool (Z.ai FP8 shards -> byte-preserved fmt=7 container) —
   the writer side, no metadata stamp.
4. fp8: end-to-end test, real repack tool -> real C loader — the real
   round trip, closing the last coverage gap.

[monotophic]

Working on the Windows / MinGW fail. Will update with a fix shortly.

[monotophic]

Fixed and ready for review

[JustVugg]

I have to block this one, and I'm sorry to do it on work this careful — but the refusal logic fires on a shape GLM-5.2 actually has, so merging it can stop the engine from loading real containers.

The collision includes GLM-5.2's own o_proj

self_attn.o_proj.weight loads as [D, H*v_head] = [6144, 16384] on GLM-5.2. Verified independently:

per-row int8 scale bytes : O*4              = 6144*4      = 24576
per-128x128-block bytes  : ceil(6144/128) * ceil(16384/128) * 4
                         = 48 * 128 * 4                   = 24576   <-- identical

48 * 128 == 6144 == O exactly, so a valid, pre-existing int8-row o_proj satisfies both geometries and hits THE DESIGN LANDMINE — exit(1) at model load. The general family is any int8 tensor where O == ceil(O/128)*ceil(I/128); [6144,16384] is not a contrived corner, it is the shape in production.

This is reachable through this repo's own tooling: --ebits is documented as "default 4; 8 per --mtp/--indexer" and --io-bits defaults to 8, so int8 residents are written by our recommended configuration (the int8-MTP-head setup from #8). Please double-check the exact tensors your local converter emits at those settings — if o_proj lands as int8 in any documented configuration, this merge turns a working snapshot into one that refuses to start.

The reason CI is green: every C fixture uses hidden_size=128, where o_proj is [128,128] and cannot collide. No test in the suite uses a real GLM-5.2 shape.

The fix I'd like, and why it's strictly safer

Invert the collision policy. On is_row && is_blk, resolve to fmt=1 — the incumbent, already-on-disk, decodable format — and move the refusal to the writer: _check_geometry() in repack_fp8_passthrough.py refuses to emit an fmt=7 container at an ambiguous shape. Nothing existing breaks, and nothing this toolchain can produce is ever misread. Today the trade runs the other way: a decodable int8 tensor is refused to protect an encoding the reader can't decode anyway.

Two supporting notes, both verifiable in the branch:

• _check_geometry() has no collision check at all. Its docstring promises "never anything that could coincide with O4 by accident… refuses rather than silently emitting a container the engine might misread"*, but it only asserts the block counts. A [6144,16384] fp8 o_proj is repacked happily, then refused by the loader.
• test_fp8_load.c enshrines the regression: expect_refuse(128,16384,…,512), (256,16384,…,1024), (384,16384,…,1536) are all valid int8-row tensors (ns == O*4) asserted to be rejected. I'd want [6144,16384] added as an explicit non-refusal test.

Three smaller things

1. qt_scale_bytes() has no production callers. The ~25-line comment says the two qt_wire_mmap/qt_unwire_mmap sites were fixed; both still compute scale_b=(int64_t)t->O*4. It only compiles clean because -Wno-unused-function is in CFLAGS. So the claimed fmt=4 expert-pinning fix isn't applied, and an fmt=7 tensor reaching qt_wire_mmap gets weight_b overshooting the allocation whenever nblk > O. Wire it or drop it, and fix the test comment that asserts otherwise.
2. fmt=7 has no live GPU path. The shader branch is correct — I checked the decode against the LUT, including subnormals, the single NaN code and -0.0 — but the only caller of coli_metal_matmul in the tree is a test. Production dispatch (matmul_qt_ex, coli_metal_gemm, moe_submit) all exclude 7, and this PR adds fmt!=7 guards to the fused paths. So adopting fmt=7 on Apple silicon means those tensors run the scalar matmul_fp8 on CPU and disable the fused attention + shared-expert Metal kernel for every layer they appear in. That should be stated plainly, or bind_gemv wired to the new branch.
3. The guard comments say the shader "would silently misread fp8 as f32" — stale now that this same PR adds the fmt==7 branch.

What's right

The E4M3 LUT is correct OCP E4M3-FN (spot-checked 0x01, 0x78, 0x7E, 0x7F/0xFF, 0x80). No OOB in matmul_fp8. qt_bytes fmt=7 prevents the AUTOPIN undercount. The fmt=6 hardening only alters behaviour at I==98 and actually fixes a pre-existing misread of an int8 [1,98]. And the instinct throughout — refuse rather than guess — is the right one; the bug is which side of the ambiguity gets the benefit of the doubt.

Invert the collision policy, add the [6144,16384] regression test, and I'll take this. The ordinal fmt=7 is yours.

[monotophic]

Authored by Fable 5 in Claude Code, analysis in partnership with @monotophic
(we/our = the joint work; me/my = @monotophic alone).

You were right, and the original wasn't up to the standard we want our
contributions to meet — we've tightened our process around exactly the gap
you exposed, and re-verified this from written requirements rather than
implementation intent. Thank you for the catch; the o_proj analysis was
exact (confirmed on real containers on my side, including one where the
pushed head refused to load).

The revision (head d1a6897, rebased onto current dev @ e4b5bf3 —
clean replay over #488 and the Inkling merge):

1. Collision policy inverted as requested. Reader: unstamped
   is_row && is_blk resolves to fmt=1, the incumbent. Writer:
   _check_geometry() refuses to emit fmt=7 at any shape where
   O == ceil(O/128)*ceil(I/128). [6144,16384] is now an explicit
   non-refusal regression test; the three expect_refuse cases you flagged
   (plus four degenerate cases on the same inverted branch) assert fmt=1.
2. The family is enumerated, not sampled: a census tool
   (c/tools/fp8_collision_census.py) sweeps the predicate over every real
   tensor shape; on GLM-5.2, self_attn.o_proj [6144,16384] is the
   only ambiguous family (79 instances/container — every real O here is a
   multiple of 128, so the family reduces to real-O × I=16384).
3. qt_scale_bytes() wired at both qt_wire_mmap/qt_unwire_mmap sites
   via a shared qt_wire_split(), with a seam test that intercepts the
   actual mlock/munlock calls — a sites-only regression to O*4
   fails 4 named assertions. Two adjacent gaps closed while making the
   helper authoritative: an fmt=6 branch (fixed 4-byte tag, not O*4),
   and a loud named refusal in qt_addrow for formats its fall-through
   can't decode (fmt=6/7 there previously walked into the int2 arm —
   the same class your fmt=4 comment in that function documents).
4. GPU-path status stated plainly (PR body): fmt=7 runs scalar
   matmul_fp8 on CPU and fmt=7-carrying layers are excluded from the
   fused Metal kernels by this PR's guards; wiring bind_gemv to the
   (tested) shader branch is a follow-up we want to take on my hardware.
5. Stale guard comments and the FORMATS.md narrative corrected to the
   post-inversion policy.

How it was verified — the part that changed most on our side:

• A spec-blind acceptance suite: built from a written requirements spec
  by a process that never saw this implementation, with an independently
  derived OCP E4M3-FN oracle (from the OFP8 spec text — which, incidentally,
  contradicts itself on E4M3's exponent bias between §4.2 and Table 1;
  resolved via Table 2's worked examples). Against this branch it verified
  bit-exact decode at all 16,900 positions of a partial-block tensor
  (130×130 — shapes no real GLM tensor exercises), NaN row-propagation
  semantics, the collision inversion, and the writer refusal on a real
  ambiguous shape. It also found the qt_addrow crash above — a defect
  every earlier review layer missed.

• Differential container matrix: current-dev binary vs this branch's
  head, over the stock int8-MTP container and an int8-resident mint, with
  per-run usage-state control — identical loads, per-tensor resolutions,
  notices, and byte-identical temp-0 generated text on both containers.
  The before/after in one pair of lines: the previously-pushed head refuses
  the int8-resident container at load (model.layers.0.self_attn. o_proj.weight: [6144,16384] scale array is 24576 bytes — matches BOTH … refusing, exit 1); this head loads it, resolves o_proj as fmt=1, and
  generates the same text as dev.

• Blind validation + an independent hazard audit (their own mutation tests
  and negative controls, including a reproduction of your exact nblk > O
  overshoot scenario), then full gates on the final head: make check,
  METAL=1 zero warnings, metal-test both COLI_METAL_RESSET states,
  per-commit bisect builds, suppression-stripped -Wunused-function build.

  Behavioral contract

This PR was revised against a written contract; every invariant below is a
registered claim (stated before the revision was built), each with its
decisive evidence:

Invariant	Capstone evidence
Every container that loads on dev loads identically here (resolutions, notices, output)	Differential matrix vs. current dev: byte-identical temp-0 generated text on the stock int8-MTP container and an int8-resident mint
No byte pattern readable today changes meaning; ambiguity favors the incumbent	[6144,16384] int8-row (GLM-5.2's real o_proj): refused by the previous head (exit 1, quoted in-thread), loads as fmt=1 here
The writer can never emit what the reader would mis-resolve	Repack refusal demonstrated at a real ambiguous shape ([2048,16384]); round-trip load verified on non-ambiguous output
fmt=7 decode is exact OCP E4M3-FN — passthrough, no approximation	Independently derived oracle (from the OFP8 spec text): bit-exact at all 16,900 positions of a partial-block (130×130) tensor
Failure is loud: every refusal names its condition	New: qt_addrow/qt_matvec_rows refuse unhandled formats by name (previously: SIGSEGV for fmt=7, silent misread for fmt=6 — found by the acceptance suite, fixed here)
Collision family is enumerated, not sampled	In-tree census over every real tensor shape: o_proj is the only ambiguous family (79 instances/container)

Full evidence matrix (instruments and results per requirement)

• Compatibility (I1): old-binary/new-binary differential over both real
  containers, per-run usage-state control, telemetry-excised output
  comparison; plus same-binary control run proving the method's
  discrimination. Result: identical loads, per-tensor resolutions, notices,
  and generated text, all four comparisons.
• Collision policy (E4/I2): reader inversion verified at source and
  behaviorally (collision-family fixtures at 4 real shapes + 1
  non-collision control resolve fmt=1); the three previously-refusing test
  shapes plus four degenerate cases now assert fmt=1; census tool
  regenerable upstream (c/tools/fp8_collision_census.py).
• Exact decode (I4): spec-blind acceptance suite with an independently
  derived E4M3-FN oracle (derivation caught an internal contradiction in
  the OFP8 spec's own bias statement, resolved via its worked examples);
  bit-exact on NaN-free data; NaN weights verified to poison their row
  through the dot product (IEEE-correct propagation, 130/130 rows).
• Wire-path sizing: qt_scale_bytes() wired via shared
  qt_wire_split(); seam test intercepts actual mlock/munlock calls —
  a sites-only regression to O*4 fails 4 named assertions; fmt=6's
  4-byte-tag case corrected while making the helper authoritative.
• Loud failure (I6): refusal tests for fmt=6/7 through both absorb-path
  functions (previously crash/silent); repack tool refuses on zero
  repack-target tensors (previously silent success).
• Gates: make check; METAL=1 zero warnings; metal-test under both
  COLI_METAL_RESSET states; per-commit bisect builds; suppression-
  stripped -Wunused-function build; POSIX-construct check (Windows CI
  pattern) on all new tests.

Deeper artifacts (specific transcripts, census tables, suite design)
available on request.

Thanks again for the review feedback, this was a valuable reset.

[monotophic]

CI caught a Linux-only short-read in our test harness's stderr capture (the refusal fired; the test's read was truncated) — fixed, head now 9f6b4ab.

[monotophic]

Authored by Fable 5 in Claude Code, analysis in partnership with @monotophic
(we/our = the joint work; me/my = @monotophic alone).

Rebased over today's wave: rev3, head 6313f32 on dev @ 75a7374
(TEST_BINS superset resolutions only; product code patch-identical to
9f6b4ab, which your review + our revision comment describe). Full
verification re-run on this head: gates green, the spec-blind acceptance
suite passes (bit-exact fmt=7 decode, collision inversion, writer refusal),
and the differential container matrix vs current dev shows byte-identical
temp-0 output on both real containers.