From 9d3edec790e6f5938ec028aec7316497cba1691b Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 14 Aug 2026 19:03:39 +0000 Subject: [PATCH 1/4] spec(#468): VT-FP8-W8A8-CPU-ARM -- register the static fp8 W8A8 path on CPU Issue #468 records that the model-layer wiring of VT_GDN_FP8_ALPHA_IN_CONV has no test at any tier, and names the structural reason: the fp8 matmul registers only on CUDA, so a CPU-tier test cannot reach the path at all and the CUDA tier needs a GPU. Its "What done would look like" option 1 is a CPU registration for the fp8 matmul -- valuable well beyond that lever, because it makes the whole fp8 model path CPU-testable. This is the spec for option 1, committed before any implementation. It scopes two CPU registrations (kQuantFp8Static, kMatmulFp8Cutlass), the comment repairs they create the need for, and the consequential change to the fusion additivity test, whose three fp8-terminal recipes currently assert the full composite THROWS on CPU. That assertion is a checker's claim, so changing it takes a spec plus red-before and green-after evidence, which is what this file declares and what the implementation commit carries. Two things are stated up front rather than discovered later. First the residual gap: the model-layer entry points key on kMatmulFp8CublasLt, which deliberately stays CUDA-only, so this row makes the OP seam CPU-reachable and does not claim to make MatmulFp8CutlassD execute on a CPU queue. Second the stale comment: the op contract and the CUDA kernel both describe the quant math as x / input_scale while the code multiplies by the reciprocal, which is what upstream ships. The code is right and the comment is wrong, and left alone it invites someone to "correct" a default-ON 35B path into a divide. Also repoints the #468 roadmap row at this spec and its work branch. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5[1m] [claude-code] --- .agents/roadmap_v1.md | 2 +- .agents/specs/vt-fp8-w8a8-cpu-arm.md | 226 +++++++++++++++++++++++++++ 2 files changed, 227 insertions(+), 1 deletion(-) create mode 100644 .agents/specs/vt-fp8-w8a8-cpu-arm.md diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index dd2d6353b..232f05146 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -134,7 +134,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#213](https://github.com/mudler/vllm.cpp/issues/213) | `PERF-27B-LMHEAD-FP4` | Qwen3.6 NVFP4 baselines (27B and 35B-A3B) must reach vLLM speed parity | perf | | [#339](https://github.com/mudler/vllm.cpp/issues/339) | `PERF-27B-LMHEAD-FP4` | 27B c1: every fp8 input projection asks for an f32 output, selecting the slower nvjet template family where vLLM emits bf16 (48 f32-out projections 18.51 ms vs 48 bf16-out 7.05 ms). The merged GDN `in_proj` arm is built DEFAULT OFF as `VT_GDN_FP8_IN_BF16`, spec [`perf-fp8-alpha-fold.md`](specs/perf-fp8-alpha-fold.md) §Attempt 4 — UNMEASURED: no committed gate loads the fp8 tower (`row/GATE-27B-FP8-TOWER-GOLDEN` builds that arm) | perf | | [#402](https://github.com/mudler/vllm.cpp/issues/402) | `PERF-27B-LMHEAD-FP4` | 27B: a STALE COMMENT keeps the dense path off an already-built, default-ON, byte-exact quant fusion — 112 -> 16 QuantFp8Static launches/step (and the epilogue-fusion premise is REFUTED). Its §3 "Lever B" (fold the per-column FP8 alpha into the cuBLASLt epilogue) is spec [`perf-fp8-alpha-fold.md`](specs/perf-fp8-alpha-fold.md), work branch `row/PERF-FP8-ALPHA-FOLD` | perf | -| [#468](https://github.com/mudler/vllm.cpp/issues/468) | `PERF-27B-LMHEAD-FP4` | `VT_GDN_FP8_ALPHA_IN_CONV`'s model-layer wiring has no test at any tier: the fp8 matmul registers only on CUDA (`cuda_matmul.cu:827-829`) so CPU cannot reach it, and `ProjectGdnFp8QkvzForTest` is called with DEFAULT args, so nothing pins the 8 call sites forwarding `mixed_scale`. Lever parked on `row/PERF-ALPHA-IN-CONV-PARKED` | bug | +| [#468](https://github.com/mudler/vllm.cpp/issues/468) | `PERF-27B-LMHEAD-FP4` | `VT_GDN_FP8_ALPHA_IN_CONV`'s model-layer wiring has no test at any tier: the fp8 matmul registers only on CUDA (`cuda_matmul.cu:827-829`) so CPU cannot reach it, and `ProjectGdnFp8QkvzForTest` is called with DEFAULT args, so nothing pins the 8 call sites forwarding `mixed_scale`. Lever parked on `row/PERF-ALPHA-IN-CONV-PARKED`. The issue's option 1 (the CPU registration) is spec [`vt-fp8-w8a8-cpu-arm.md`](specs/vt-fp8-w8a8-cpu-arm.md), work branch `row/VT-FP8-W8A8-CPU-ARM`; it lands the op seam on CPU and leaves the MODEL-layer predicate (which keys on `kMatmulFp8CublasLt`) open | bug | | [#501](https://github.com/mudler/vllm.cpp/issues/501) | `PERF-27B-LMHEAD-FP4` | `AlphaVecBf16TakesTwoLaunch` bounded a COUNT of ulp mismatches instead of their MAGNITUDE, and was RED on its first CUDA run at ~26% — the double-rounding population the bf16-D lever produces by construction. Replaced by a max-ulp bound (`<= 1`, and `<= 0` at a pow2 alpha), measured 0/1-ulp only over 2.17M words on GB10, spec [`perf-fp8-alpha-fold.md`](specs/perf-fp8-alpha-fold.md) §The bf16-vs-f32 divergence is DOUBLE ROUNDING | bug | | [#521](https://github.com/mudler/vllm.cpp/issues/521) | `PERF-27B-LMHEAD-FP4` | [`perf-fp8-alpha-fold.md`](specs/perf-fp8-alpha-fold.md) `:19`/`:211` claim the bf16-D lever "also applies to 35B-A3B" — it is INERT there: `GdnOutDType(dense_model=false)` is F32 on a MoE, contradicting the code's own comment at `qwen3_5.cpp:3617-3619` | bug | | [#391](https://github.com/mudler/vllm.cpp/issues/391) | `PERF-CPU-DECODE-BARRIER` | CPU backend: batch-1 decode is barrier-bound (47% sync), and paged attention branches per element | perf | diff --git a/.agents/specs/vt-fp8-w8a8-cpu-arm.md b/.agents/specs/vt-fp8-w8a8-cpu-arm.md new file mode 100644 index 000000000..f6434ccfb --- /dev/null +++ b/.agents/specs/vt-fp8-w8a8-cpu-arm.md @@ -0,0 +1,226 @@ +# VT-FP8-W8A8-CPU-ARM — register the static FP8 W8A8 path on the CPU backend + +Issue: [#468](https://github.com/mudler/vllm.cpp/issues/468). +Owning row: `PERF-27B-LMHEAD-FP4` (the row #468 is filed against in the +[roadmap issue table](../roadmap_v1.md)). +Related: [`perf-fp8-alpha-fold.md`](perf-fp8-alpha-fold.md) (the lever whose +model-layer wiring #468 found untestable), +[`portable-fusion-framework.md`](portable-fusion-framework.md) §3b/§6 (the +"backend-negotiated fp8 tail" this row retires on CPU). + +#468 lists two ways to close its coverage debt. This row is **option 1**: *"a CPU +registration for the fp8 matmul sufficient to exercise the wiring (which has +value well beyond this lever — it would make the whole fp8 model path +CPU-testable)"*. It does not do option 2, and it does not touch the parked +`VT_GDN_FP8_ALPHA_IN_CONV` lever. + +## Scope + +**IN — three registrations' worth of surface, in two files.** + +1. `src/vt/cpu/cpu_ops.cpp`: register `OpId::kQuantFp8Static` on + `DeviceType::kCPU`. The fp8-e4m3fn codec this needs already lives in that file + (`F32ToFp8`, saturating RNE, whose comment already claims it bit-matches + `vllm::F32ToF8E4M3`; `kFp8Max = 448.0F`), so the kernel is the scale + application and nothing else. +2. `src/vt/cpu/cpu_ops.cpp`: register `OpId::kMatmulFp8Cutlass` on + `DeviceType::kCPU` as a **correctness reference** — f32 accumulate, one folded + `alpha`. It makes no performance claim and is not a lever. +3. `include/vt/ops.h` + `src/vt/cuda/cuda_matmul_fp8_cutlass.cu` + + `include/vt/fused_recipe.h`: comment repairs the first two items create the + need for (below). + +**IN — the consequential test changes.** `tests/vt/test_fused_chain_additivity.cpp` +asserts, for all three fp8-terminal recipes, that the FULL Tier-0 composite +**THROWS** on CPU because `vt::QuantFp8Static` is unregistered there. Registering +it makes those three `CHECK_THROWS` false. They are not deleted: each is +**replaced by the strictly stronger byte-exact assertion** the registration now +makes possible — full composite == standalone-op-sequence golden, fp8 output +included — and the catalog's `cpu_full` flag flips to `true` for those three +rows. That is a change of a checker's claim, which is why this spec exists and +why the RED-before/GREEN-after evidence below is mandatory. + +**OUT.** + +- `OpId::kMatmulFp8CublasLt` and `OpId::kMatmulFp8CublasLtAlphaVec` stay + CUDA-only. A "cuBLASLt" kernel on the CPU backend would be a lie in the name, + and the alpha-vec op's whole claim is about a cuBLASLt epilogue pointer mode. + **This has a consequence that must not be glossed** — see *Residual gap*. +- Any CUDA behavior. Not one line of the CUDA fp8 kernels changes; only a + comment that describes them wrongly. +- `src/vllm/model_executor/models/nemotron_h*.{h,cpp}` and + `tests/vllm/models/test_nemotron_h_*` — owned elsewhere, tracked by #517. +- The parked `VT_GDN_FP8_ALPHA_IN_CONV` lever on + `row/PERF-ALPHA-IN-CONV-PARKED`, and the 8 `mixed_scale` call sites. #468's + option 2 (an argument-forwarding test) is not attempted here. + +### Residual gap — stated up front, not discovered later + +The model-layer entry points `MatmulFp8CutlassD` and +`MatmulFp8CutlassPreQuantD` (`src/vllm/model_executor/models/qwen3_5.cpp`) both +gate on `vt::OpRegistered(vt::OpId::kMatmulFp8CublasLt, d.q.device.type)` and +refuse with *"the fp8 W8A8 path is CUDA-only"*. Registering `kMatmulFp8Cutlass` +and `kQuantFp8Static` on CPU therefore makes the **op seam** CPU-reachable and +the **fusion catalog's fp8 terminal** CPU-reachable; it does **not** on its own +make `MatmulFp8CutlassD` execute on a CPU queue. Closing that last step is a +separate scoped change (either widening the predicate to the op actually +selected, or registering a CPU arm for the cuBLASLt op id) and stays open under +#468. This row does not claim it. + +## Upstream chain + +Pinned oracle `/home/mudler/_git/vllm` @ `5559679229bc961848b121ccdeaa8fa5d79bec98` +(0.26.0.dev0), the pin recorded in [`upstream-sync.md`](../upstream-sync.md). +Verified present at that SHA on 2026-08-14; the line numbers below were re-derived +at HEAD rather than copied from an earlier record. + +| What | Upstream anchor | +|---|---| +| Dispatch to the static fp8 linear method is unconditional | `vllm/model_executor/layers/quantization/modelopt.py:2527-2528` — `if quant_algo == "FP8": return ModelOptFp8LinearMethod(self.fp8_config)`; exclusions are checked first at `:2519-2523` | +| The method is **static, hard-coded** — no dynamic fallback inside it | `modelopt.py:510-513` — `init_fp8_linear_kernel(activation_quant_key=kFp8StaticTensorSym, weight_quant_key=kFp8StaticTensorSym, …)` | +| `input_scale` is created per merged shard and collapsed to a **scalar** | `modelopt.py:502-508` (`PerTensorScaleParameter`, one entry per output partition) then `:528` — `layer.input_scale = Parameter(layer.input_scale.max(), …)` | +| Per-token dynamic is a **different class**, reached by a different `quant_algo` | `ModelOptFp8PcPtLinearMethod`, `modelopt.py:540`, selected at `:393-395` | +| The quant math | `csrc/quantization/w8a8/fp8/common.cuh:58-77` — `x = val * scale` under `is_scale_inverted=true` (`:62`), else `x = val / scale` (`:64`); then `fmaxf(-448, fminf(x, 448))` (`:68`); then the hardware RNE convert (`:71-76`) | +| The reciprocal is formed ONCE, outside the elementwise math | `csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31` and `:38` — `1.0f / scale[…]` | +| Per-tensor selection = one group over the whole tensor | `common.cu:204-210` — `scale.numel() == 1` ⇒ `group_m = num_tokens`, `group_n = hidden_size` | + +So upstream's shipped path is **reciprocal-multiply**, not divide: the `val / scale` +arm at `common.cuh:64` is the `is_scale_inverted=false` branch, and the caller that +feeds the static per-tensor path hands it an already-inverted scale +(`common.cu:31`). Our CUDA kernel mirrors that +(`cuda_matmul_fp8_cutlass.cu:331-333`: `const float inv = 1.0f / input_scale;` +then `LoadIn(x, i) * inv`). The CPU kernel added here does the same. The two are +**not** interchangeable in f32: `x / s` and `x * (1/s)` differ by up to one ulp +before the fp8 round, and near an e4m3 tie that ulp changes the emitted byte. + +**Our one documented deviation is PRESERVED, not revisited.** Upstream applies +`scale_a` and `scale_b` as two epilogue scalars +(`vllm/model_executor/layers/quantization/utils/scaled_mm/cutlass.py` `ScaledEpilogue`); +we fold `alpha = input_scale * weight_scale` into one f32 +(`include/vt/ops.h`, the `MatmulFp8Cutlass` contract). The CPU reference folds it +the same way, so it is a reference for the op we actually ship. + +## Design + +### `QuantFp8StaticKernel` (CPU) + +``` +inv = 1.0f / input_scale // formed ONCE, mirroring common.cu:31 +out[i] = F32ToFp8(LoadF32(x, i) * inv) +``` + +`LoadF32` handles the f32 and bf16 `x` dtypes the op contract admits, matching the +CUDA kernel's `LoadIn` overloads (bf16 is widened to f32 *before* the multiply, so +the two backends round at the same point). `F32ToFp8` supplies the saturation +(`a >= 448 ⇒ 0x7E`, which is the encoding of 448 — clamp-then-convert and +saturating-convert coincide because 448 IS the largest finite e4m3fn value) and +the round-to-nearest-even. Row-chunked through the existing `ForRows` so it +inherits the file's thread pool; the op is elementwise, so chunking cannot move a +value. + +### `MatmulFp8CutlassKernel` (CPU) + +``` +acc = Σ_k Fp8ToF32(a[m,k]) · Fp8ToF32(b[n,k]) // f32 accumulate +out[m,n] = alpha · acc // ONE folded scalar +``` + +Shaped exactly like the neighbouring `MatmulNvfp4Fp4Kernel`: the A row is decoded +once per M and reused across N. **This is a correctness reference, not a +performance path** — it is a naive triple loop, it makes no speed claim, and +nothing routes a production model through it. It is stated as such in the code. + +It is *not* a bit-mirror of the CUDA GEMM and does not claim to be: the CUDA arm +reduces K in tensor-core order and rounds its epilogue through bf16, so the two +agree to fp8/bf16 tolerance, not to the byte. Only the QUANT half is a bit-mirror +claim, and that is what G2 tests. + +### Comment repairs + +Three comments become false the moment these registrations land, and one is +already false today: + +- `include/vt/ops.h` (`QuantFp8Static` contract): *"CUDA only (the 35B W8A8 path + is CUDA-resident)"* → CUDA + CPU. +- `include/vt/ops.h` (`MatmulFp8Cutlass` contract): *"CUDA-only (sm120a)"* → + CUDA (sm120a) + a CPU reference. +- `include/vt/fused_recipe.h` (`kQuantFp8` opcode): *"CUDA-only"*. +- **Already false, and the reason this one matters most:** `include/vt/ops.h` + states the contract as `clamp(x[i] / input_scale, …)` and + `src/vt/cuda/cuda_matmul_fp8_cutlass.cu` repeats it as `fp8_e4m3(clamp(x[i]/input_scale, …))`, + while the code three lines below multiplies by the reciprocal. **The code is + right and the comment is wrong.** Left alone, it invites someone to "correct" + a default-ON 35B path into a divide, which is a token-visible change nothing + currently forbids. The comment is repaired to the reciprocal form with the + upstream anchor beside it. + +## Risks + +| Risk | Mitigation | +|---|---| +| A CPU registration silently becomes a reference-tier fallback on unified-memory accelerators (`MaybeInstallReferenceTier`, `src/vt/op_provider.cpp:204-225`), turning a hard refusal into a slow silent success on e.g. Metal/Vulkan | Not silent by construction: the tier is announced once per (op, device) on stderr as `[vt reference-tier] … running the PORTABLE CPU fallback (correct but slow)` and counted (`op_provider.cpp:516-528`). Recorded here as an intended, visible consequence rather than left to be discovered | +| The three `CHECK_THROWS` in `test_fused_chain_additivity.cpp` are "fixed" by deletion, weakening the additivity proof | They are REPLACED by full-composite byte-exact checks and the `cpu_full` flags flip to `true`; the case's `== 9` catalog count guard is untouched. Assertion count goes UP, not down, and that is recorded in the evidence | +| `F32ToFp8`'s comment claims it bit-matches `vllm::F32ToF8E4M3`, and the claim was never tested against an INDEPENDENT reference | G1 is exactly that test, and its reference is derived from the format definition + upstream's clamp, not from either of our two codecs | +| A CPU-vs-CUDA byte divergence exists and nobody sees it, because this box has no GPU | G2 is declared, and reported PENDING with the reason rather than skipped. It is not counted as satisfied | + +## Tests and gates + +New file `tests/vt/test_ops_fp8_cpu.cpp`, wired into `tests/CMakeLists.txt`. + +**G1 — bitwise, zero tolerance.** `vt::QuantFp8Static` on a CPU queue must +produce bytes identical to an **independently written** reference: +`clamp(x · (1/input_scale), ±448)` then e4m3fn round-to-nearest-even. The +reference is built from the *format*, not from our code: it enumerates all 256 +e4m3fn encodings, decodes each to an exact `double` from the field layout, and +picks the nearest with an even-significand tie-break by scanning. That is a +different algorithm from both `F32ToFp8` (frexp + `nearbyint`) and +`vllm::F32ToF8E4M3`, so agreement is evidence rather than tautology. Not +expressible as `doctest::Approx` — `Approx`'s `scale` term defaults to 1.0 and +gives it a ~1.19e-5 absolute floor, which is meaningless for a byte compare. +Coverage includes both `x` dtypes, values that overflow ±448 in both signs, +subnormal-range values, exact ties, and zeros of both signs. + +**G2 — CPU vs CUDA, bitwise.** The CPU registration must agree byte-for-byte +with the CUDA kernel on the same input. **PENDING on this host: it has no GPU +(`nvidia-smi` absent, no CUDA device).** The case is committed and CUDA-gated so +it binds on the next GPU holder; it is not claimed as satisfied here. Gate hosts +are `dgx.casa` (GB10/sm_121) and `192.168.68.23` (Thor/sm_110). + +**G3 — the GEMM.** `vt::MatmulFp8Cutlass` on CPU against a `double` reference +that reproduces upstream's LOSSY pipeline — clamp, e4m3 RNE, dequant, then exact +accumulation. A reference doing exact arithmetic on the pre-quant floats would +make a WRONG implementation look better than upstream and pass; the tolerance +then bounds only the K-reduction order, which is the only thing that legitimately +differs. + +**Mutations that must prove the gate is armed.** Each applied ALONE to a restored +tree, rebuilt, run, restored; the compiler exit status is recorded beside every +result, because a mutation that fails to BUILD reads as a passing test. + +| # | Mutation | Must | +|---|---|---| +| M1 | `input_scale` ignored (activation left unscaled) | G1 RED | +| M2 | divide instead of reciprocal-multiply | G1 RED — and if it does NOT, that is a reportable finding: the gate would be blind to the exact defect the stale comment invites | +| M3 | saturation removed (no clamp to ±448) | G1 RED on an overflowing input | +| M4 | round-to-nearest-even replaced by truncation | G1 RED | +| M5 | `alpha` reduced to `weight_scale` only | G3 RED | + +`Status:` and the case count are read on every mutation run, not just +`assertions:` — a failing doctest binary can print `0 failed` when a case throws. + +**Full gate:** `scripts/agent-preflight.sh` and the CPU ctest suite, with `df -h` +before and after. `test_op_parity` is RED on `main` at the base SHA +(base-inherited, #755/#672) and is subtracted as a known baseline red. + +## Stop conditions + +- G1 RED against the independent reference on any covered input ⇒ stop and + report. It would mean either the new kernel or the pre-existing `F32ToFp8` + disagrees with upstream's format, and the second is a bug in shipped code that + gets its own issue and fix in-flow. +- A mutation that does not RED ⇒ report the survivor rather than strengthening + the mutation until it dies. A surviving mutant is a finding about the gate. +- Any need to touch `kMatmulFp8CublasLt`, the model layer, or CUDA behavior ⇒ + `NEEDS_DECISION`, not a silent scope widening. +- Disk below ~10 GiB free ⇒ stop and report; an ENOSPC build leaves the PREVIOUS + binary in place and makes checkers emit false policy refusals. From 9087f5ba3779cda94f54558211a7fc5f4a59cdc4 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 14 Aug 2026 19:04:02 +0000 Subject: [PATCH 2/4] feat(#468): register the static fp8 W8A8 path on the CPU backend The fp8 matmul registered only on CUDA, so the fp8 model path was structurally untestable on the machine most of this work happens on: a CPU-tier test could not reach it and the CUDA tier needs a GPU. That is the reason #468 gives for the wiring shipping uncovered, and it is what this closes at the op seam. Two CPU registrations, per .agents/specs/vt-fp8-w8a8-cpu-arm.md: kQuantFp8Static -- the scale application, over the fp8-e4m3fn codec already in cpu_ops.cpp. It multiplies by the RECIPROCAL, formed once outside the loop, which is what upstream ships (common.cuh:62 `x = val * scale` with the inverse formed at common.cu:31) and what our CUDA kernel does. Not a divide: the two differ by up to one f32 ulp before the fp8 round, and near an e4m3 tie that ulp changes the emitted byte. kMatmulFp8Cutlass -- a CORRECTNESS REFERENCE. f32 accumulate, one folded alpha, naive triple loop. It makes no speed claim, nothing routes a production model through it, and it says so in the code. No CUDA behavior changes. kMatmulFp8CublasLt stays CUDA-only, so the MODEL-layer predicate still refuses on CPU; the spec records that residual gap and the new test pins it rather than letting it be assumed closed. Comment repairs, in-flow under #468. Three comments said "CUDA only" and become false here. A fourth was ALREADY false and matters most: include/vt/ops.h and cuda_matmul_fp8_cutlass.cu both describe the quant as `x / input_scale` while the code three lines below multiplies by the reciprocal. The code is right and the comment is wrong; left alone it invites someone to "correct" the code on a default-ON 35B path. Both now carry the upstream anchor and say not to. Gate: tests/vt/test_ops_fp8_cpu.cpp. G1 is BITWISE with zero tolerance against an independently written reference -- it enumerates all 128 finite e4m3fn magnitudes, decodes each from the field layout, and picks nearest-even by scanning, which is a different algorithm from both F32ToFp8 (frexp/nearbyint) and vllm::F32ToF8E4M3, so agreement is evidence rather than tautology. Not expressed as doctest Approx, whose scale term defaults to 1.0 and would give a byte compare a ~1.19e-5 floor. G3 checks the GEMM against a double reference that reproduces upstream's LOSSY pipeline -- clamp, e4m3 RNE, dequant -- because an exact-arithmetic reference would let a wrong implementation look better than upstream and pass. G2 (CPU vs CUDA, byte for byte) is committed and CUDA-gated; it is PENDING, not skipped, on this GPU-less host and says so in its output. test_fused_chain_additivity asserted, for all three fp8-terminal recipes, that the full Tier-0 composite THROWS on CPU. Registering the terminal makes that false. The assertions are not deleted to go green -- each is replaced by the strictly stronger byte-exact check the registration makes available, the prefix check is kept beside it so a future tail regression is localised, and cpu_full flips to true for those three rows. Assertions go 19 -> 25. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5[1m] [claude-code] --- include/vt/fused_recipe.h | 4 +- include/vt/ops.h | 22 +- src/vt/cpu/cpu_ops.cpp | 70 ++++ src/vt/cuda/cuda_matmul_fp8_cutlass.cu | 13 +- tests/CMakeLists.txt | 4 + tests/vt/test_fused_chain_additivity.cpp | 117 ++++-- tests/vt/test_ops_fp8_cpu.cpp | 434 +++++++++++++++++++++++ 7 files changed, 625 insertions(+), 39 deletions(-) create mode 100644 tests/vt/test_ops_fp8_cpu.cpp diff --git a/include/vt/fused_recipe.h b/include/vt/fused_recipe.h index e03689b70..42465a53a 100644 --- a/include/vt/fused_recipe.h +++ b/include/vt/fused_recipe.h @@ -44,7 +44,9 @@ namespace vt { // kSigmoidGate -> vt::SigmoidGateBf16 (a·sigmoid(b)) // kRmsNormGated-> vt::RmsNormGated (gated rms-normalize) // kRope -> vt::RopeFromCache (partial NeoX RoPE from a cos/sin cache) -// kQuantFp8 -> vt::QuantFp8Static (static per-tensor fp8 terminal; CUDA-only) +// kQuantFp8 -> vt::QuantFp8Static (static per-tensor fp8 terminal; CUDA + CPU +// since #468 — the fp8-terminal recipes now +// realize END-TO-END on CPU, not negotiated) // kQuantFp4 -> vt::ScaledFp4Quant (dynamic per-group fp4 terminal) // kAttnQkNormRopeGate -> vt::AttnQkNormRopeGate (fused full-attention preamble) // diff --git a/include/vt/ops.h b/include/vt/ops.h index bf5ef66e4..88d58a629 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -1453,10 +1453,19 @@ void MatmulNvfp4Cutlass(Queue& q, Tensor& out, const Tensor& a_packed, const Ten // per-tensor weight_scale + a f32 per-tensor input_scale (both applied directly: // dequant(w)=f8(w)*weight_scale, dequant(a)=f8(a)*input_scale). -// QuantFp8Static (mirror vLLM static_scaled_fp8_quant, is_scale_inverted=False): -// out_fp8[i] = fp8_e4m3( clamp(x[i] / input_scale, -448, 448) ) // RNE hw cvt +// QuantFp8Static (mirror vLLM static_scaled_fp8_quant): +// inv = 1/input_scale; out_fp8[i] = fp8_e4m3( clamp(x[i] * inv, -448, 448) ) +// RNE convert. The scale is applied as a RECIPROCAL MULTIPLY, not a divide, and +// the reciprocal is formed ONCE outside the elementwise loop — that is what +// upstream ships: `x = val * scale` under is_scale_inverted=true +// (csrc/quantization/w8a8/fp8/common.cuh:62, clamp at :68) with the inverse formed +// by the caller (csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31, +// `1.0f / scale[...]`). DO NOT "correct" the kernels to a divide to match a +// prose formula: `x/s` and `x*(1/s)` differ by up to one f32 ulp, and near an +// e4m3 tie that ulp changes the emitted byte on a default-ON 35B path. // Static per-tensor scale (NOT dynamic/per-token). x [M,K] f32/bf16, out [M,K] -// i8 (raw fp8-e4m3fn bytes). CUDA only (the 35B W8A8 path is CUDA-resident). +// i8 (raw fp8-e4m3fn bytes). CUDA + CPU (the CPU arm is the portable reference +// that makes the fp8 seam testable without a GPU, #468). void QuantFp8Static(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scale); // RmsNormQuantFp8 (fused fp8 RMSNorm -> static per-tensor activation quant). One @@ -1507,7 +1516,12 @@ void RmsNormGatedQuantFp8(Queue& q, Tensor& out_fp8, const Tensor& x, const Tens // sequential scale_a·(scale_b·acc) — within fp8 tolerance, ported deviation). // a_fp8 [M,K] (= QuantFp8Static output), b_fp8 [N,K] the on-disk raw fp8-e4m3fn // weight (K contiguous). out [M,N] bf16 (cutlass epilogue) or f32 (via cast). -// K,N multiples of 16 (128-bit fp8 alignment). CUDA-only (sm120a). +// K,N multiples of 16 (128-bit fp8 alignment). CUDA (sm120a) + a CPU CORRECTNESS +// REFERENCE (f32 accumulate, naive triple loop — no speed claim, and no +// production model routes through it; it exists so the fp8 seam resolves on a +// CPU queue and can be gated without a GPU, #468). The CPU arm agrees with the +// CUDA kernel to fp8/bf16 tolerance, NOT byte-for-byte: the CUDA arm reduces K +// in tensor-core order and rounds its epilogue through bf16. void MatmulFp8Cutlass(Queue& q, Tensor& out, const Tensor& a_fp8, const Tensor& b_fp8, float alpha); diff --git a/src/vt/cpu/cpu_ops.cpp b/src/vt/cpu/cpu_ops.cpp index 5d342c66a..eb416f7cf 100644 --- a/src/vt/cpu/cpu_ops.cpp +++ b/src/vt/cpu/cpu_ops.cpp @@ -532,6 +532,72 @@ uint8_t F32ToFp8(float f) { static_cast(mant)); } +// --- Static per-tensor FP8 W8A8 (VT-FP8-W8A8-CPU-ARM, #468). The CPU arm of the +// path vLLM's ModelOptFp8LinearMethod runs: a static per-tensor activation quant +// followed by a per-tensor fp8 GEMM. It exists so the fp8 seam is reachable, and +// therefore testable, without a GPU. + +// QuantFp8Static CPU kernel — mirror of vLLM's static_scaled_fp8_quant +// (csrc/quantization/w8a8/fp8/common.cuh:58-77 `scaled_fp8_conversion`): +// x = val * scale; r = fmaxf(-448, fminf(x, 448)); hardware RNE convert +// with the RECIPROCAL formed ONCE outside the loop, exactly as upstream forms it +// (csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31 `1.0f / scale[...]`) +// and exactly as our CUDA kernel does (cuda_matmul_fp8_cutlass.cu: `const float +// inv = 1.0f / input_scale;` then `LoadIn(x, i) * inv`). It is a MULTIPLY BY THE +// RECIPROCAL, not a divide: the two differ by up to one f32 ulp before the fp8 +// round, and near an e4m3 tie that ulp changes the emitted byte. +// +// F32ToFp8 supplies both remaining halves: it saturates (|a| >= 448 -> 0x7E, +// which IS the encoding of 448, so clamp-then-convert and saturating-convert +// coincide because 448 is the largest finite e4m3fn value) and it rounds to +// nearest-even. The scale is per-TENSOR — upstream collapses the per-shard +// input_scale to one scalar with `.max()` (modelopt.py:528) and then treats +// `scale.numel() == 1` as a single group spanning the whole tensor +// (common.cu:204-210). LoadF32 widens a bf16 x to f32 BEFORE the multiply, as +// the CUDA kernel's LoadIn overload does, so both backends round at one point. +void QuantFp8StaticKernel(Queue&, Tensor& out_fp8, const Tensor& x, float input_scale) { + const int64_t n = x.shape[0] * x.shape[1]; + const float inv_scale = 1.0F / input_scale; + uint8_t* op = out_fp8.Ptr(); + ForRows(n, [&](int64_t r0, int64_t r1) { + for (int64_t i = r0; i < r1; ++i) op[i] = F32ToFp8(LoadF32(x, i) * inv_scale); + }); +} + +// MatmulFp8Cutlass CPU kernel: out[m,n] = alpha * Sum_k f8val(a[m,k])*f8val(b[n,k]), +// f32 accumulate, ONE folded alpha (= input_scale*weight_scale — our recorded +// deviation from upstream's two epilogue scalars, see include/vt/ops.h). +// +// A CORRECTNESS REFERENCE, NOT A PERFORMANCE PATH. It is a naive triple loop; it +// makes no speed claim and nothing routes a production model through it. Its +// purpose is that the fp8 GEMM seam resolves on a CPU queue so the surrounding +// wiring can be gated without a GPU (#468). +// +// It is deliberately NOT a bit-mirror of the CUDA GEMM and does not claim to be: +// the CUDA arm reduces K in tensor-core order and rounds its epilogue through +// bf16, so the two agree to fp8/bf16 tolerance. Only the QUANT half above carries +// a bit-exactness claim. Shaped like MatmulNvfp4Fp4Kernel: the A row is decoded +// once per M and reused across N. +void MatmulFp8CutlassKernel(Queue&, Tensor& out, const Tensor& a_fp8, const Tensor& b_fp8, + float alpha) { + const int64_t m = a_fp8.shape[0], k = a_fp8.shape[1], n = b_fp8.shape[0]; + const auto* ap = a_fp8.Ptr(); + const auto* bp = b_fp8.Ptr(); + ForRows(m, [&](int64_t r0, int64_t r1) { + std::vector arow(static_cast(k)); + for (int64_t i = r0; i < r1; ++i) { + for (int64_t kk = 0; kk < k; ++kk) + arow[static_cast(kk)] = Fp8ToF32(ap[i * k + kk]); + for (int64_t col = 0; col < n; ++col) { + float acc = 0.0F; + for (int64_t kk = 0; kk < k; ++kk) + acc += arow[static_cast(kk)] * Fp8ToF32(bp[col * k + kk]); + StoreF32(out, i * n + col, alpha * acc); + } + } + }); +} + // Fused fp8 RMSNorm -> static per-tensor quant (mirror vLLM Inductor // fused_add_rms_norm_static_fp8_quant, rms_quant_fusion.py:124). Same reduction // order as RmsNormKernel; the fp8 is taken from the SAME bf16-rounded normed value @@ -3120,6 +3186,10 @@ struct Registrar { reinterpret_cast(static_cast(&RmsNormKernel))); RegisterOp(OpId::kRmsNormQuantFp8, DeviceType::kCPU, reinterpret_cast(static_cast(&RmsNormQuantFp8Kernel))); + RegisterOp(OpId::kQuantFp8Static, DeviceType::kCPU, + reinterpret_cast(static_cast(&QuantFp8StaticKernel))); + RegisterOp(OpId::kMatmulFp8Cutlass, DeviceType::kCPU, + reinterpret_cast(static_cast(&MatmulFp8CutlassKernel))); RegisterOp(OpId::kSiluAndMul, DeviceType::kCPU, reinterpret_cast(static_cast(&SiluAndMulKernel))); RegisterOp(OpId::kGeluAndMul, DeviceType::kCPU, diff --git a/src/vt/cuda/cuda_matmul_fp8_cutlass.cu b/src/vt/cuda/cuda_matmul_fp8_cutlass.cu index 8ce9341a2..c5fbb6b59 100644 --- a/src/vt/cuda/cuda_matmul_fp8_cutlass.cu +++ b/src/vt/cuda/cuda_matmul_fp8_cutlass.cu @@ -317,8 +317,17 @@ void MatmulFp8CutlassKernelCuda(Queue& q, Tensor& out, const Tensor& a_fp8, cons } // ---- Static per-tensor fp8 activation quant (vLLM static_scaled_fp8_quant) --- -// out_fp8[i] = fp8_e4m3(clamp(x[i]/input_scale, -448, 448)). __NV_SATFINITE cvt -// saturates == clamp-then-cvt; RNE == vLLM's hardware cvt. Tin f32/bf16. +// inv = 1/input_scale; out_fp8[i] = fp8_e4m3(clamp(x[i]*inv, -448, 448)). +// A RECIPROCAL MULTIPLY, not a divide, and the reciprocal is hoisted out of the +// loop — that is upstream's shipped form (`x = val * scale` with the inverse +// formed by the caller: csrc/quantization/w8a8/fp8/common.cuh:62 and +// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31). The code below is +// RIGHT; do not "fix" it into `x / input_scale` to match a prose formula. The two +// differ by up to one f32 ulp before the fp8 round, and near an e4m3 tie that +// ulp changes the emitted byte on a default-ON 35B path. +// __NV_SATFINITE cvt saturates == clamp-then-cvt; RNE == vLLM's hardware cvt. +// Tin f32/bf16. The CPU arm (src/vt/cpu/cpu_ops.cpp QuantFp8StaticKernel) is the +// byte-for-byte mirror of this kernel. __device__ __forceinline__ uint8_t F32ToFp8Dev(float f) { return static_cast(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3)); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 97009945a..b3a60b644 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1403,6 +1403,10 @@ if(VLLM_CPP_CUTLASS) "VT_FP4_TEST_PERSISTENT_SAVE=1;VT_FP4_PERSISTENT_CACHE=1;VT_FP4_AUTOTUNE_CACHE_READONLY=0;VT_FP4_AUTOTUNE_CACHE_PATH=${CMAKE_CURRENT_BINARY_DIR}/nvfp4-runtime-save-lifecycle.json;VT_FP4_AUTOTUNE_DELAY_US=0;VT_FP4_PLAN_CACHE=1;VT_FP4_AUTOTUNE=1;VT_FP4_FULL_TACTICS=1") endif() vllm_cpp_add_test(test_ops_fp8_cutlass vt/test_ops_fp8_cutlass.cpp) +# VT-FP8-W8A8-CPU-ARM (#468): the CPU arm of the static fp8 W8A8 path. Runs on a +# box with NO GPU by construction — that is the point of the row — and carries a +# CUDA-gated arm (G2) for CPU-vs-CUDA byte agreement wherever a device exists. +vllm_cpp_add_test(test_ops_fp8_cpu vt/test_ops_fp8_cpu.cpp) # Opt-in arm: run the fp8 plan-cache byte-exact case with the cache ENABLED # (VT_FP8_PLAN_CACHE=1 -> first MatmulFp8CublasLt call builds the plan fresh, # later calls hit the cache). Proves the cached-plan GEMM is BYTE-identical to the diff --git a/tests/vt/test_fused_chain_additivity.cpp b/tests/vt/test_fused_chain_additivity.cpp index 775c49956..809878fb0 100644 --- a/tests/vt/test_fused_chain_additivity.cpp +++ b/tests/vt/test_fused_chain_additivity.cpp @@ -21,10 +21,17 @@ // generic loop, asserts each runs CORRECT on the CPU 'second backend' via the Tier-0 // composite — BYTE-EXACT vs the standalone-op-sequence golden, over the CPU-expressible // scope. For recipes whose quant terminal is CPU-expressible (fp4 + the attn macro + -// plain add-rmsnorm) that is end-to-end; for the CUDA-only static-fp8 quant terminal -// (`vt::QuantFp8Static`, unregistered on CPU per §3b/§6), the composite runs the -// CPU-expressible PREFIX byte-exact, and the test ALSO asserts the full composite THROWS -// on CPU — documenting the backend-negotiated tail rather than silently skipping it. +// plain add-rmsnorm) that is end-to-end; for the three recipes ending in the +// static-fp8 quant terminal (`vt::QuantFp8Static`) the composite ALSO runs end-to-end +// byte-exact, since VT-FP8-W8A8-CPU-ARM (#468) gave that opcode a CPU kernel. Those three +// recipes keep their PREFIX check as well — the prefix is what a backend inherits before +// the terminal exists, and it is what would localise a future regression to the tail. +// +// HISTORY, because the change is a checker's CLAIM and not just a number: until #468 the +// three fp8-terminal rows carried `cpu_full = false` and asserted the full composite +// THROWS on CPU, documenting a backend-negotiated tail (spec §3b/§6). The throw was not +// deleted to go green — it became impossible, and was replaced by the strictly stronger +// byte-exact assertion the registration made available. // // ADDITIVITY EVIDENCE this test pins (catalog grows ⇒ backend does NOT): // * The catalog (`include/vt/recipes.h`) grew 1→6→7 recipes across W0→W1→W3, while @@ -98,10 +105,12 @@ std::vector PackBf16(const std::vector& f) { return out; } -// The recipe with its trailing step (the CUDA-only static-fp8 quant terminal) and its -// trailing operand (the fp8 output slot) dropped — the CPU-expressible PREFIX. Purely -// structural on the POD (no per-recipe knowledge): every fp8-terminal recipe in the -// catalog has the fp8 quant as its last step and the fp8 output as its last operand. +// The recipe with its trailing step (the static-fp8 quant terminal) and its trailing +// operand (the fp8 output slot) dropped — the PREFIX a backend inherits before that +// terminal has a kernel. Purely structural on the POD (no per-recipe knowledge): every +// fp8-terminal recipe in the catalog has the fp8 quant as its last step and the fp8 +// output as its last operand. Still checked alongside the full composite since #468, so +// a future tail regression is localised rather than merely detected. FusedRecipe CpuExpressiblePrefix(const FusedRecipe& r) { FusedRecipe pr = r; pr.n = r.n - 1; @@ -379,9 +388,9 @@ void CheckAttnQkNormRopeGate() { CHECK(gc == gg); } -// --- fp8-terminal recipes: CPU-expressible PREFIX byte-exact + negotiated-tail throw -- +// --- fp8-terminal recipes: PREFIX byte-exact + the fp8 terminal, end-to-end on CPU ---- -// kRmsNormQuantFp8 — prefix: (add residual) + gemma-RMSNorm → bf16; fp8 tail CUDA-only. +// kRmsNormQuantFp8 — prefix: (add residual) + gemma-RMSNorm → bf16; then the fp8 tail. void CheckRmsNormQuantFp8() { const int64_t t = 3, h = 256; const float eps = 1e-6f, scale = 0.125f; @@ -416,17 +425,34 @@ void CheckRmsNormQuantFp8() { CHECK(tmp_c == tmp_g); CHECK(res_c == res_g); - // Backend-negotiated tail (§3b/§6): the full composite reaches the CUDA-only - // vt::QuantFp8Static terminal, which is unregistered on the CPU 'second backend' → - // it THROWS. The CPU backend inherits the portable prefix, negotiates the fp8 tail. - std::vector fp8(static_cast(t * h)); - Tensor tfp8 = MakeTensor(fp8.data(), DType::kI8, {t, h}); - b.op[4] = &tfp8; + // The fp8 TERMINAL, now end-to-end on CPU (VT-FP8-W8A8-CPU-ARM, #468). This + // assertion USED to be `CHECK_THROWS`: vt::QuantFp8Static was unregistered on the + // CPU 'second backend', so the full composite refused and the test could only + // document a backend-negotiated tail. The CPU registration retires that + // negotiation, and the throw is replaced by the strictly STRONGER claim it made + // impossible — the full composite equals the standalone-op-sequence golden BYTE + // for BYTE, fp8 output included. + std::vector res_f = rf; // fresh residual: the prefix run above consumed one + std::vector tmp_f(tmp_g.size()); + std::vector fp8_f(static_cast(t * h)); + Tensor ttmp_f = MakeTensor(tmp_f.data(), DType::kBF16, {t, h}); + Tensor trf = MakeTensor(res_f.data(), DType::kF32, {t, h}); + Tensor tfp8_f = MakeTensor(fp8_f.data(), DType::kI8, {t, h}); + b.op[2] = &trf; + b.op[3] = &ttmp_f; + b.op[4] = &tfp8_f; b.n = 5; - CHECK_THROWS(vt::FusedChainComposite(q, vt::kRmsNormQuantFp8, b, p)); + vt::FusedChainComposite(q, vt::kRmsNormQuantFp8, b, p); + + std::vector fp8_g(fp8_f.size()); + Tensor tfp8_g = MakeTensor(fp8_g.data(), DType::kI8, {t, h}); + vt::QuantFp8Static(q, tfp8_g, ttmp_g, scale); + CHECK(tmp_f == tmp_g); + CHECK(res_f == res_g); + CHECK(fp8_f == fp8_g); } -// kRmsNormGatedQuantFp8 — prefix: gated-RMSNorm → bf16; fp8 tail CUDA-only. +// kRmsNormGatedQuantFp8 — prefix: gated-RMSNorm → bf16; then the fp8 tail. void CheckRmsNormGatedQuantFp8() { const int64_t rows = 3, d = 256; const float eps = 1e-6f, scale = 0.1875f; @@ -457,15 +483,26 @@ void CheckRmsNormGatedQuantFp8() { vt::FusedChainComposite(q, CpuExpressiblePrefix(vt::kRmsNormGatedQuantFp8), b, p); CHECK(tmp_c == tmp_g); - std::vector fp8(static_cast(rows * d)); - Tensor tfp8 = MakeTensor(fp8.data(), DType::kI8, {rows, d}); - b.op[4] = &tfp8; + // fp8 terminal end-to-end on CPU (#468) — was CHECK_THROWS, see the note in + // CheckRmsNormQuantFp8 above. + std::vector tmp_f(tmp_g.size()); + std::vector fp8_f(static_cast(rows * d)); + Tensor ttmp_f = MakeTensor(tmp_f.data(), DType::kBF16, {rows, d}); + Tensor tfp8_f = MakeTensor(fp8_f.data(), DType::kI8, {rows, d}); + b.op[3] = &ttmp_f; + b.op[4] = &tfp8_f; b.n = 5; - CHECK_THROWS(vt::FusedChainComposite(q, vt::kRmsNormGatedQuantFp8, b, p)); + vt::FusedChainComposite(q, vt::kRmsNormGatedQuantFp8, b, p); + + std::vector fp8_g(fp8_f.size()); + Tensor tfp8_g = MakeTensor(fp8_g.data(), DType::kI8, {rows, d}); + vt::QuantFp8Static(q, tfp8_g, ttmp_g, scale); + CHECK(tmp_f == tmp_g); + CHECK(fp8_f == fp8_g); } -// kSiluMulQuantFp8 — the W3 mechanical-sync recipe. Prefix: silu(gate)·up → bf16; fp8 -// tail CUDA-only. This is the executable proof that the CPU 'second backend' inherited +// kSiluMulQuantFp8 — the W3 mechanical-sync recipe. Prefix: silu(gate)·up → bf16; then +// the fp8 tail. This is the executable proof that the CPU 'second backend' inherited // a WHOLE NEW fusion pass (ported in W3) with ZERO backend edits — same generic path. void CheckSiluMulQuantFp8() { const int64_t m = 3, i = 256; @@ -493,17 +530,33 @@ void CheckSiluMulQuantFp8() { vt::FusedChainComposite(q, CpuExpressiblePrefix(vt::kSiluMulQuantFp8), b, p); CHECK(tmp_c == tmp_g); - std::vector fp8(static_cast(m * i)); - Tensor tfp8 = MakeTensor(fp8.data(), DType::kI8, {m, i}); - b.op[3] = &tfp8; + // fp8 terminal end-to-end on CPU (#468) — was CHECK_THROWS. For THIS recipe the + // change is the additivity claim at its sharpest: a whole vLLM fusion pass that + // landed touching 2 shared files and zero backend files is now realized + // END-TO-END by the CPU 'second backend', because one opcode gained a host + // kernel and every recipe using it inherited the reach for free. + std::vector tmp_f(tmp_g.size()); + std::vector fp8_f(static_cast(m * i)); + Tensor ttmp_f = MakeTensor(tmp_f.data(), DType::kBF16, {m, i}); + Tensor tfp8_f = MakeTensor(fp8_f.data(), DType::kI8, {m, i}); + b.op[2] = &ttmp_f; + b.op[3] = &tfp8_f; b.n = 4; - CHECK_THROWS(vt::FusedChainComposite(q, vt::kSiluMulQuantFp8, b, p)); + vt::FusedChainComposite(q, vt::kSiluMulQuantFp8, b, p); + + std::vector fp8_g(fp8_f.size()); + Tensor tfp8_g = MakeTensor(fp8_g.data(), DType::kI8, {m, i}); + vt::QuantFp8Static(q, tfp8_g, ttmp_g, scale); + CHECK(tmp_f == tmp_g); + CHECK(fp8_f == fp8_g); } // The catalog: the SINGLE enumeration of every recipe the framework declares. Each row // pairs a recipe with the CPU-scope driver above. `cpu_full` records whether the CPU // 'second backend' realizes the recipe end-to-end (true) or up to the CPU-expressible -// prefix with the fp8 quant terminal backend-negotiated (false). Adding a recipe adds +// prefix only (false). Every row is `true` since #468 gave the fp8 quant terminal a CPU +// kernel; the column stays because it is the mechanism, and a future opcode with no host +// kernel would set it false again. Adding a recipe adds // ONE row here (a catalog/test concern) — the BACKEND path (composite walker + the one // kFusedChain registration) is untouched. This IS the additivity mechanism, executable. struct CatalogEntry { @@ -520,9 +573,9 @@ const CatalogEntry kCatalog[] = { {&vt::kSiluMulFp4Quant, "kSiluMulFp4Quant", true, &CheckSiluMulFp4Quant}, {&vt::kSigmoidGateFp4Quant, "kSigmoidGateFp4Quant", true, &CheckSigmoidGateFp4Quant}, {&vt::kAttnQkNormRopeGate, "kAttnQkNormRopeGate", true, &CheckAttnQkNormRopeGate}, - {&vt::kRmsNormQuantFp8, "kRmsNormQuantFp8", false, &CheckRmsNormQuantFp8}, - {&vt::kRmsNormGatedQuantFp8, "kRmsNormGatedQuantFp8", false, &CheckRmsNormGatedQuantFp8}, - {&vt::kSiluMulQuantFp8, "kSiluMulQuantFp8", false, &CheckSiluMulQuantFp8}, + {&vt::kRmsNormQuantFp8, "kRmsNormQuantFp8", true, &CheckRmsNormQuantFp8}, + {&vt::kRmsNormGatedQuantFp8, "kRmsNormGatedQuantFp8", true, &CheckRmsNormGatedQuantFp8}, + {&vt::kSiluMulQuantFp8, "kSiluMulQuantFp8", true, &CheckSiluMulQuantFp8}, }; } // namespace diff --git a/tests/vt/test_ops_fp8_cpu.cpp b/tests/vt/test_ops_fp8_cpu.cpp new file mode 100644 index 000000000..041593d43 --- /dev/null +++ b/tests/vt/test_ops_fp8_cpu.cpp @@ -0,0 +1,434 @@ +// vllm.cpp original (vt runtime, inventory deviation §9.1); no upstream mirror +// of the TEST, but the reference math below is transcribed from upstream. +// +// VT-FP8-W8A8-CPU-ARM (.agents/specs/vt-fp8-w8a8-cpu-arm.md), issue #468. +// +// The CPU arm of the static per-tensor FP8 W8A8 path — vt::QuantFp8Static and +// vt::MatmulFp8Cutlass registered on DeviceType::kCPU — is what makes the fp8 +// seam reachable, and therefore gateable, without a GPU. This file is that gate. +// +// G1 BITWISE, ZERO TOLERANCE. vt::QuantFp8Static on CPU must equal an +// INDEPENDENTLY WRITTEN reference quantizer, byte for byte. +// G2 CPU vs CUDA, bitwise, on the identical input. CUDA-gated. +// G3 vt::MatmulFp8Cutlass on CPU against a `double` reference that reproduces +// upstream's LOSSY pipeline (clamp, e4m3 RNE, dequant) before accumulating. +// +// WHY THE G1 REFERENCE IS WRITTEN THE WAY IT IS. It is derived from the FORMAT +// and from upstream's formula, never from a codec in this tree — otherwise it +// would be a tautology dressed as a gate. `RefEncodeRne` enumerates all 128 +// finite e4m3fn magnitudes, decodes each to an exact double from the field +// layout, and picks the nearest with an even-significand tie-break by scanning. +// That is a different ALGORITHM from `F32ToFp8` (frexp + std::nearbyint) and +// from `vllm::F32ToF8E4M3`, so agreement between them is evidence. +// +// Upstream chain (pinned oracle @ 5559679229bc961848b121ccdeaa8fa5d79bec98): +// csrc/quantization/w8a8/fp8/common.cuh:58-77 scaled_fp8_conversion +// :62 x = val * scale (is_scale_inverted == true) +// :68 fmaxf(-448, fminf(x, 448)) +// :71 hardware RNE convert +// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31 1.0f / scale[...] +// — the reciprocal is formed ONCE, outside the elementwise math +// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:204-210 scale.numel()==1 +// ⇒ ONE group over the whole tensor (per-tensor, not per-token) +// vllm/model_executor/layers/quantization/modelopt.py:510-513 / :528 +// — the method is STATIC and hard-coded; input_scale collapses to a scalar +// +// So the scale is applied as `x * (1/s)`, NOT `x / s`. The two differ by up to +// one f32 ulp before the fp8 round, and near an e4m3 tie that ulp changes the +// emitted byte, which is why G1 compares BYTES and not an Approx: doctest's +// Approx carries a `scale` term defaulting to 1.0 and therefore a ~1.19e-5 +// absolute floor, meaningless for a byte compare. +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/ops.h" + +namespace { + +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; +using vt::Tensor; + +Device Cpu() { return Device{DeviceType::kCPU, 0}; } + +bool HasCuda() { + try { + vt::GetBackend(DeviceType::kCUDA); + return true; + } catch (const std::runtime_error&) { + return false; + } +} + +Tensor MakeTensor(void* data, DType dt, Device dev, const std::vector& shape) { + Tensor t; + t.data = data; + t.dtype = dt; + t.device = dev; + t.rank = static_cast(shape.size()); + int64_t stride = 1; + for (int i = t.rank - 1; i >= 0; --i) { + t.shape[i] = shape[static_cast(i)]; + t.stride[i] = stride; + stride *= shape[static_cast(i)]; + } + return t; +} + +// --------------------------------------------------------------------------- +// The INDEPENDENT e4m3fn reference. Written from the format: 1 sign bit, 4 +// exponent bits (bias 7), 3 mantissa bits, NO infinities, and 0x7F/0xFF the only +// NaN encodings (that is what the "fn" in e4m3fn means). Nothing here reads any +// codec in src/. +constexpr float kE4m3MaxFinite = 448.0f; // = 1.75 * 2^8, encoding 0x7E + +// Exact value of one finite e4m3fn magnitude encoding. +double E4m3Exact(unsigned exp_field, unsigned mant) { + if (exp_field == 0) return std::ldexp(static_cast(mant), -9); // mant / 512 + return std::ldexp(1.0 + static_cast(mant) / 8.0, static_cast(exp_field) - 7); +} + +// Round-to-nearest-EVEN encode by exhaustive nearest-value scan over the 128 +// finite magnitudes. `r` must already be clamped to [-448, 448] and finite. +uint8_t RefEncodeRne(float r) { + const auto sign = static_cast(std::signbit(r) ? 0x80u : 0x00u); + const double a = std::fabs(static_cast(r)); + unsigned best_e = 0, best_m = 0; + double best_d = std::numeric_limits::infinity(); + for (unsigned e = 0; e <= 15; ++e) { + for (unsigned m = 0; m <= 7; ++m) { + if (e == 15 && m == 7) continue; // the NaN encoding is not a value + const double d = std::fabs(a - E4m3Exact(e, m)); + // Strictly nearer wins. On an EXACT tie prefer the even significand — the + // tie-break also carries across an exponent step, because mant 7 (odd) at + // exponent e is adjacent to mant 0 (even) at e+1. + if (d < best_d || (d == best_d && (m & 1u) == 0u && (best_m & 1u) != 0u)) { + best_d = d; + best_e = e; + best_m = m; + } + } + } + return static_cast(sign | static_cast(best_e << 3) | + static_cast(best_m)); +} + +// The whole upstream expression, in upstream's order and upstream's f32 width: +// form the reciprocal ONCE, multiply, clamp, convert RNE. +uint8_t RefQuantFp8Static(float x, float input_scale) { + const float inv = 1.0f / input_scale; // common.cu:31 + const float scaled = x * inv; // common.cuh:62 + const float r = std::fmax(-kE4m3MaxFinite, std::fmin(scaled, kE4m3MaxFinite)); // :68 + return RefEncodeRne(r); // :71 RNE cvt +} + +// Dequant for the GEMM reference: the same format decode, sign restored. +double RefDequant(uint8_t byte) { + const double m = E4m3Exact(static_cast(byte >> 3) & 0xFu, + static_cast(byte) & 0x7u); + return (byte & 0x80u) != 0 ? -m : m; +} + +// The G1 input population. Deliberately not just "random in a nice range": +// overflow in BOTH signs, the subnormal ladder, exact ties at several exponents, +// and both zeros — every input class whose handling a mutation can break. +std::vector G1Inputs(float input_scale, uint32_t seed) { + std::vector v; + // Random bulk. + std::mt19937 rng(seed); + std::uniform_real_distribution ux(-2.0f, 2.0f); + for (int i = 0; i < 4096; ++i) v.push_back(ux(rng)); + // Saturation, both signs: these scale past +-448 and MUST clamp. + for (float s : {600.0f, 448.5f, 1e4f, 1e30f}) { + v.push_back(s * input_scale); + v.push_back(-s * input_scale); + } + // Exactly the largest finite value, and just inside it. + v.push_back(448.0f * input_scale); + v.push_back(-448.0f * input_scale); + v.push_back(447.0f * input_scale); + // Exact ties in the NORMAL range: midpoints between adjacent e4m3 values at + // several exponents. RNE must pick the even mantissa. + for (int e = -4; e <= 8; ++e) { + for (int m = 0; m < 7; ++m) { + const double lo = std::ldexp(1.0 + m / 8.0, e); + const double hi = std::ldexp(1.0 + (m + 1) / 8.0, e); + const auto mid = static_cast((lo + hi) / 2.0); + v.push_back(mid * input_scale); + v.push_back(-mid * input_scale); + } + } + // The SUBNORMAL ladder and its midpoints (values m/512 and (2m+1)/1024). + for (int m = 0; m <= 8; ++m) { + v.push_back(static_cast(m / 512.0) * input_scale); + v.push_back(static_cast((2 * m + 1) / 1024.0) * input_scale); + v.push_back(-static_cast((2 * m + 1) / 1024.0) * input_scale); + } + // Both zeros (e4m3fn has a signed zero) and a value that rounds to zero. + v.push_back(0.0f); + v.push_back(-0.0f); + v.push_back(static_cast(1.0 / 4096.0) * input_scale); + return v; +} + +// Runs the CPU op over `x` and compares BYTE FOR BYTE against the reference. +// Returns the number of differing bytes so a caller can report a rate. +size_t RunG1(const std::vector& x, float input_scale, DType x_dtype) { + const auto n = static_cast(x.size()); + Queue q{Cpu(), nullptr}; + + // Materialize x at the requested width. For bf16 the REFERENCE consumes the + // bf16-rounded value too, so this compares the codec and not the store width. + std::vector xf32(x); + std::vector xbf16(x.size()); + for (size_t i = 0; i < x.size(); ++i) xbf16[i] = vt::F32ToBF16(x[i]); + std::vector ref_in(x.size()); + for (size_t i = 0; i < x.size(); ++i) + ref_in[i] = x_dtype == DType::kBF16 ? vt::BF16ToF32(xbf16[i]) : xf32[i]; + + void* xp = x_dtype == DType::kBF16 ? static_cast(xbf16.data()) + : static_cast(xf32.data()); + Tensor tx = MakeTensor(xp, x_dtype, Cpu(), {1, n}); + std::vector got(x.size()); + Tensor tout = MakeTensor(got.data(), DType::kI8, Cpu(), {1, n}); + vt::QuantFp8Static(q, tout, tx, input_scale); + + std::vector want(x.size()); + for (size_t i = 0; i < x.size(); ++i) want[i] = RefQuantFp8Static(ref_in[i], input_scale); + + size_t bad = 0, first_bad = 0; + for (size_t i = 0; i < got.size(); ++i) { + if (got[i] != want[i]) { + if (bad == 0) first_bad = i; + ++bad; + } + } + if (bad != 0) { + CAPTURE(bad); + CAPTURE(first_bad); + CAPTURE(ref_in[first_bad]); + CAPTURE(static_cast(got[first_bad])); + CAPTURE(static_cast(want[first_bad])); + } + CHECK(bad == 0); + // VACUITY GUARD: a kernel that wrote nothing, or a population that is all + // zeros, would compare equal to a reference that also produced zeros. Demand + // that the run actually produced a spread of distinct non-zero encodings. + size_t nonzero = 0; + for (auto b : want) { + if ((b & 0x7Fu) != 0u) ++nonzero; + } + CHECK(nonzero > want.size() / 2); + return bad; +} + +} // namespace + +// =========================================================================== +// G1 — bitwise, zero tolerance. +TEST_CASE("G1: CPU QuantFp8Static is BYTE-identical to an independent e4m3 reference") { + // The registration itself is the thing under test; a missing one must fail + // here rather than surface as a confusing throw inside the helper. + REQUIRE(vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCPU)); + + // Several scales, including 1.0 (where the tie inputs above land EXACTLY on + // e4m3 midpoints, so RNE is what decides every one of them) and a + // non-power-of-two production-shaped scale. + for (float s : {1.0f, 0.5f, 0.035f, 0.0092f, 7.25f}) { + CAPTURE(s); + const auto x = G1Inputs(s, 1234u); + CHECK(RunG1(x, s, DType::kF32) == 0u); + CHECK(RunG1(x, s, DType::kBF16) == 0u); + } +} + +// =========================================================================== +// G2 — the CPU registration must agree with the CUDA kernel BIT for BIT on the +// same input. This is the arm that says the CPU path is a mirror of what ships, +// not merely self-consistent with a host reference. +TEST_CASE("G2: CPU QuantFp8Static == CUDA QuantFp8Static, byte for byte") { + if (!HasCuda()) { + // NOT a silent skip. The CPU registration is still asserted so the case can + // never be vacuous, and the banner names what is owed. + MESSAGE("G2 PENDING: no CUDA device on this host, CPU-vs-CUDA byte agreement " + "was NOT measured (gate hosts: dgx.casa GB10/sm_121, 192.168.68.23 Thor/sm_110)"); + CHECK(vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCPU)); + return; + } + vt::Backend& b = vt::GetBackend(DeviceType::kCUDA); + Queue gq = b.CreateQueue(); + Queue cq{Cpu(), nullptr}; + const Device gpu{DeviceType::kCUDA, 0}; + + for (float s : {1.0f, 0.5f, 0.035f, 0.0092f, 7.25f}) { + CAPTURE(s); + const auto x = G1Inputs(s, 4321u); + const auto n = static_cast(x.size()); + + std::vector cpu_out(x.size()); + std::vector xf(x); + Tensor tx_cpu = MakeTensor(xf.data(), DType::kF32, Cpu(), {1, n}); + Tensor to_cpu = MakeTensor(cpu_out.data(), DType::kI8, Cpu(), {1, n}); + vt::QuantFp8Static(cq, to_cpu, tx_cpu, s); + + void* dx = b.Alloc(xf.size() * sizeof(float)); + void* dout = b.Alloc(xf.size()); + b.Copy(gq, dx, xf.data(), xf.size() * sizeof(float)); + Tensor tx_gpu = MakeTensor(dx, DType::kF32, gpu, {1, n}); + Tensor to_gpu = MakeTensor(dout, DType::kI8, gpu, {1, n}); + vt::QuantFp8Static(gq, to_gpu, tx_gpu, s); + std::vector gpu_out(x.size()); + b.Copy(gq, gpu_out.data(), dout, gpu_out.size()); + b.Synchronize(gq); + b.Free(dx); + b.Free(dout); + + size_t bad = 0, first_bad = 0; + for (size_t i = 0; i < cpu_out.size(); ++i) { + if (cpu_out[i] != gpu_out[i]) { + if (bad == 0) first_bad = i; + ++bad; + } + } + if (bad != 0) { + CAPTURE(bad); + CAPTURE(first_bad); + CAPTURE(x[first_bad]); + CAPTURE(static_cast(cpu_out[first_bad])); + CAPTURE(static_cast(gpu_out[first_bad])); + } + CHECK(bad == 0); + } + b.DestroyQueue(gq); +} + +// =========================================================================== +// G3 — the GEMM, against a reference that is LOSSY in exactly the places +// upstream is lossy. +// +// The reference quantizes the activation through the SAME clamp + e4m3 RNE the +// hardware path uses and then dequantizes both operands back, so it computes +// `alpha * Sum_k f8val(a) * f8val(b)`. A reference that instead accumulated the +// pre-quant f32 activations would be a DIFFERENT, more accurate computation, and +// a wrong implementation could sit closer to it than the correct one does — the +// gate would reward being unlike upstream. Only the accumulation WIDTH differs +// here (double vs the kernel's f32), and the tolerance bounds exactly that: +// `4 * K * FLT_EPSILON * alpha * Sum|terms|` is the standard forward bound on a +// K-term f32 recursive sum, with a factor-4 margin. It scales with the terms, +// not with the (possibly cancelled) result, which is what makes it tight. +namespace { +void RunG3(int M, int N, int K, uint32_t seed, float input_scale, float weight_scale, + DType out_dtype) { + CAPTURE(M); + CAPTURE(N); + CAPTURE(K); + CAPTURE(seed); + const float alpha = input_scale * weight_scale; + std::mt19937 rng(seed); + std::uniform_real_distribution ux(-2.0f, 2.0f); + std::uniform_int_distribution ub(0, 255); + + std::vector x(static_cast(M) * K); + for (auto& v : x) v = ux(rng); + std::vector b_fp8(static_cast(N) * K); + for (auto& v : b_fp8) { + int byte = ub(rng); + if ((byte & 0x7F) == 0x7F) byte &= ~0x7; // avoid the NaN encodings + v = static_cast(byte); + } + + Queue q{Cpu(), nullptr}; + Tensor tx = MakeTensor(x.data(), DType::kF32, Cpu(), {M, K}); + std::vector a_fp8(x.size()); + Tensor ta = MakeTensor(a_fp8.data(), DType::kI8, Cpu(), {M, K}); + vt::QuantFp8Static(q, ta, tx, input_scale); + + Tensor tb = MakeTensor(b_fp8.data(), DType::kI8, Cpu(), {N, K}); + const size_t out_n = static_cast(M) * N; + std::vector out_f32(out_dtype == DType::kF32 ? out_n : 0); + std::vector out_bf16(out_dtype == DType::kBF16 ? out_n : 0); + void* outp = out_dtype == DType::kF32 ? static_cast(out_f32.data()) + : static_cast(out_bf16.data()); + Tensor tout = MakeTensor(outp, out_dtype, Cpu(), {M, N}); + vt::MatmulFp8Cutlass(q, tout, ta, tb, alpha); + + size_t bad = 0, first_bad = 0; + double worst_ratio = 0.0; + size_t nonzero = 0; + for (int m = 0; m < M; ++m) { + for (int n = 0; n < N; ++n) { + double acc = 0.0, abs_sum = 0.0; + for (int k = 0; k < K; ++k) { + const double t = RefDequant(a_fp8[static_cast(m) * K + k]) * + RefDequant(b_fp8[static_cast(n) * K + k]); + acc += t; + abs_sum += std::fabs(t); + } + const double want = static_cast(alpha) * acc; + const size_t i = static_cast(m) * N + n; + const double got = out_dtype == DType::kF32 ? static_cast(out_f32[i]) + : static_cast(vt::BF16ToF32(out_bf16[i])); + // f32 recursive-sum forward bound, x4 margin; a bf16 store adds its own + // half-ulp (2^-9 relative), which is what the second term admits. + const double tol = 4.0 * K * static_cast(std::numeric_limits::epsilon()) * + std::fabs(static_cast(alpha)) * abs_sum + + (out_dtype == DType::kBF16 ? std::fabs(want) * 0.004 : 0.0); + const double diff = std::fabs(got - want); + if (tol > 0.0 && diff / tol > worst_ratio) worst_ratio = diff / tol; + if (!(diff <= tol)) { + if (bad == 0) { + first_bad = i; + CAPTURE(got); + CAPTURE(want); + CAPTURE(tol); + } + ++bad; + } + if (std::fabs(want) > 0.0) ++nonzero; + } + } + if (bad != 0) { + CAPTURE(bad); + CAPTURE(first_bad); + } + CAPTURE(worst_ratio); + CHECK(bad == 0); + // VACUITY GUARD: an all-zero reference would make any implementation pass. + CHECK(nonzero == out_n); +} +} // namespace + +TEST_CASE("G3: CPU MatmulFp8Cutlass matches a LOSSY double W8A8 reference") { + REQUIRE(vt::OpRegistered(vt::OpId::kMatmulFp8Cutlass, DeviceType::kCPU)); + // Decode (M=1) and small-prefill shapes; both output dtypes the op admits. + RunG3(1, 64, 128, 101, 0.035f, 0.017f, DType::kF32); + RunG3(4, 32, 256, 102, 0.035f, 0.017f, DType::kF32); + RunG3(8, 48, 128, 103, 0.041f, 0.0092f, DType::kBF16); + RunG3(3, 16, 64, 104, 1.0f, 1.0f, DType::kF32); +} + +// =========================================================================== +// The seam the whole row exists for: QuantFp8Static -> MatmulFp8Cutlass, the +// exact pair vLLM's ModelOptFp8LinearMethod runs (modelopt.py:510-513), now +// executing end-to-end on a CPU queue. Without both registrations this case +// cannot even be written, which is the coverage debt #468 recorded. +TEST_CASE("the static fp8 W8A8 pair resolves and runs end-to-end on a CPU queue") { + CHECK(vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCPU)); + CHECK(vt::OpRegistered(vt::OpId::kMatmulFp8Cutlass, DeviceType::kCPU)); + // The cuBLASLt fp8 op deliberately stays CUDA-only (a "cuBLASLt" kernel on the + // host would be a lie in the name), which is why the MODEL-layer predicate at + // qwen3_5.cpp `MatmulFp8CutlassD` still refuses on CPU. Pinned here so the + // residual gap recorded in the spec is visible rather than assumed closed. + CHECK_FALSE(vt::OpRegistered(vt::OpId::kMatmulFp8CublasLt, DeviceType::kCPU)); +} From e60a62aff047e0cf5043f5852d52f3c2fbee861f Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 14 Aug 2026 19:11:39 +0000 Subject: [PATCH 3/4] test(#468): G1's M2 kill rested on ONE scale of five -- widen it The mutation series killed every mutant, but M2 (divide instead of reciprocal-multiply) died by 2 assertions, and measuring WHY showed the margin was thinner than the green suggested. `x/s` and `x*(1/s)` agree on almost every input. Over 20000 random values in [-2,2] they never disagree at ANY scale tried. The difference is only visible where an input lands on an e4m3 tie after scaling, which is what G1's constructed tie population exists for -- and even then it is scale-dependent: over that population 10 of 18 candidate scales expose it at all, and of the five G1 shipped with, only 0.0092 did, at 24 of 209 words. So the assertion that keeps the reciprocal form -- the one the repaired comment points at, guarding a default-ON 35B path -- was one scale-list edit away from being silently disarmed. 0.13 (78/209) and 0.77 (82/209) are the strongest detectors measured and are added with that measurement written down beside them, so a later reader prunes the list knowing what it costs. Assertions 44 -> 56, still 4 cases. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5[1m] [claude-code] --- tests/vt/test_ops_fp8_cpu.cpp | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/tests/vt/test_ops_fp8_cpu.cpp b/tests/vt/test_ops_fp8_cpu.cpp index 041593d43..1d9106c51 100644 --- a/tests/vt/test_ops_fp8_cpu.cpp +++ b/tests/vt/test_ops_fp8_cpu.cpp @@ -242,10 +242,22 @@ TEST_CASE("G1: CPU QuantFp8Static is BYTE-identical to an independent e4m3 refer // here rather than surface as a confusing throw inside the helper. REQUIRE(vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCPU)); - // Several scales, including 1.0 (where the tie inputs above land EXACTLY on - // e4m3 midpoints, so RNE is what decides every one of them) and a - // non-power-of-two production-shaped scale. - for (float s : {1.0f, 0.5f, 0.035f, 0.0092f, 7.25f}) { + // THE SCALE SET IS PART OF THE GATE, not decoration. 1.0 and 0.5 make the tie + // inputs above land EXACTLY on e4m3 midpoints, so RNE alone decides every one + // of them. 0.035 / 0.0092 are production-shaped per-tensor scales. + // + // 0.13 and 0.77 are here for ONE measured reason: they are what makes the + // `x/s` vs `x*(1/s)` defect visible. Both forms agree on almost every input — + // over 20000 random values in [-2,2] they NEVER disagree at any scale tried — + // so a gate can only see the difference where an input lands on an e4m3 tie + // after scaling, which is exactly what the tie population above constructs. + // Even then it is scale-dependent: measured over the structured population, + // 10 of 18 candidate scales expose it at all, and of {1.0, 0.5, 0.035, 0.0092, + // 7.25} only 0.0092 does, at 24 of 209 words. 0.13 (78/209) and 0.77 (82/209) + // are the strongest detectors found, so the mutation dies by a wide margin + // rather than by luck. DO NOT prune this list: removing the last detecting + // scale would silently disarm the assertion that keeps the reciprocal form. + for (float s : {1.0f, 0.5f, 0.035f, 0.0092f, 7.25f, 0.13f, 0.77f}) { CAPTURE(s); const auto x = G1Inputs(s, 1234u); CHECK(RunG1(x, s, DType::kF32) == 0u); From 0df7c840c6d863b17175e7c69be9f4bf53e64437 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 14 Aug 2026 19:15:07 +0000 Subject: [PATCH 4/4] spec(#468): record the VT-FP8-W8A8-CPU-ARM gate evidence G1 and G3 PASS; G2 (CPU vs CUDA, byte for byte) is PENDING on this host because it has no GPU, and is recorded as owed rather than counted. Beyond the mutation table, three things the series measured that the pass/fail column does not carry: M1 moves ~99.7% of bytes at every scale but 1.0 (where ignoring the scale is correctly a no-op); M2 is nearly invisible and dies only because the population contains constructed exact ties AND the scale list contains a detector; and F32ToFp8 saturates in two places, so a reviewer mutating only the obvious guard would wrongly read the gate as blind. Also records that M1/M2/M5a first failed to BUILD on -Werror rather than to assert, and were re-expressed before being counted as verdicts. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5[1m] [claude-code] --- .agents/specs/vt-fp8-w8a8-cpu-arm.md | 70 ++++++++++++++++++++++++++++ 1 file changed, 70 insertions(+) diff --git a/.agents/specs/vt-fp8-w8a8-cpu-arm.md b/.agents/specs/vt-fp8-w8a8-cpu-arm.md index f6434ccfb..110174711 100644 --- a/.agents/specs/vt-fp8-w8a8-cpu-arm.md +++ b/.agents/specs/vt-fp8-w8a8-cpu-arm.md @@ -212,6 +212,76 @@ result, because a mutation that fails to BUILD reads as a passing test. before and after. `test_op_parity` is RED on `main` at the base SHA (base-inherited, #755/#672) and is subtracted as a known baseline red. +## Evidence + +Host: `mudler-ubuntu-box`, x86_64, GCC 13.3.0, **no CUDA device** (`nvidia-smi` +absent). Build: `cmake -S . -B build -G Ninja -DVLLM_CPP_BUILD_TESTS=ON +-DVLLM_CPP_BUILD_EXAMPLES=OFF -DVLLM_CPP_SERVER=OFF`, no `CMAKE_BUILD_TYPE` +(so no `NDEBUG`, asserts live), `-Wall -Wextra -Werror`. Date 2026-08-14. +Base SHA `b1cd4d8f6bb7ec5f0bd923a75dcc140becc7fdd8`. + +| Gate | Result | +|---|---| +| G1 bitwise quant | **PASS** | +| G2 CPU vs CUDA bitwise | **PENDING — no GPU on this host.** Committed and CUDA-gated; the case prints the reason and still asserts the CPU registration so it can never be vacuous. Owed on `dgx.casa` (GB10/sm_121) or `192.168.68.23` (Thor/sm_110) | +| G3 fp8 GEMM vs lossy `double` reference | **PASS** | + +`test_ops_fp8_cpu`: 4 cases / 56 assertions / `Status: SUCCESS!`. +`test_fused_chain_additivity`: 1 case / **25** assertions (was 19), `SUCCESS!`. + +### Mutations — all applied ALONE to a restored tree, compiler exit recorded + +| # | Mutation | compile | `test cases` | `assertions` | `Status` | +|---|---|---|---|---|---| +| M0 | both registrations removed (the RED-BEFORE) | 0 | 4 \| 0 passed \| **4 failed** | 6 \| 1 \| **5 failed** | **FAILURE!** | +| M1 | `input_scale` ignored | 0 | 4 \| 3 \| **1 failed** | 56 \| 32 \| **24 failed** | **FAILURE!** | +| M2 | divide instead of reciprocal-multiply | 0 | 4 \| 3 \| **1 failed** | 56 \| 50 \| **6 failed** | **FAILURE!** | +| M3 | saturation removed | 0 | 4 \| 3 \| **1 failed** | 56 \| 28 \| **28 failed** | **FAILURE!** | +| M4 | RNE replaced by truncation | 0 | 4 \| 3 \| **1 failed** | 56 \| 28 \| **28 failed** | **FAILURE!** | +| M5a | kernel ignores `alpha` | 0 | 4 \| 3 \| **1 failed** | 56 \| 53 \| **3 failed** | **FAILURE!** | +| M5b | caller folds `weight_scale` only | 0 | 4 \| 3 \| **1 failed** | 56 \| 53 \| **3 failed** | **FAILURE!** | +| — | restored (green-after) | 0 | 4 \| **4 passed** \| 0 | 56 \| **56** \| 0 | **SUCCESS!** | + +M0's assertion count DROPS to 6 rather than staying at 56 — a changed case count +is signal, and it is why `Status:` is read alongside `assertions:`. + +**M1, M2 and M5a first failed to BUILD, not to assert** (`-Werror=unused-parameter` +/ `-Wunused-variable` on the now-dead `input_scale` / `inv_scale` / `alpha`). +A mutation that fails to build reads as a passing test, so each was re-expressed +with an explicit `(void)` and re-run. Only the `compile_exit=0` rows above are +verdicts. + +### What the mutation series measured, beyond pass/fail + +**M1 (`input_scale` ignored) changes ~99.7% of output bytes** at every scale +except 1.0, where ignoring the scale is correctly a no-op: 4308/4319, 4311/4319, +4311/4319 and 4299/4319 words at scales 0.5, 0.035, 0.0092 and 7.25, and 0/4319 +at 1.0. Overall 17229/21595 = 79.8%. + +**M2 (divide vs reciprocal-multiply) is nearly invisible, and that is the +finding.** Over 20000 random values in [-2, 2] the two forms NEVER disagree, at +any of 14 scales tried. The difference only appears where an input lands on an +e4m3 tie after scaling — which is precisely what G1's constructed tie population +is for — and even there it is scale-dependent: over that population **10 of 18 +candidate scales expose it at all**. Of the five scales G1 originally shipped +with, **only 0.0092 did, at 24 of 209 words**, so the mutant died by 2 +assertions and the assertion protecting the reciprocal form was one scale-list +edit away from being silently disarmed. 0.13 (78/209) and 0.77 (82/209) were +measured as the strongest detectors and added, with the numbers recorded beside +them in the test. M2 now dies by 6 assertions across 3 scales. + +This is the concrete answer to the question the spec's comment-repair section +raised: the gate CAN see the exact defect the stale comment invites, but only +because the input population contains constructed exact ties and the scale list +contains a detector. Neither is decoration. + +**M3 exposed defense-in-depth in the pre-existing codec.** `F32ToFp8` saturates +in TWO places — the early `a >= kFp8Max` return and a late `exp_field > 15` +overflow guard — so removing either alone leaves saturation intact and the +mutant survives. M3 removes both, which is what "saturation removed" has to mean +for this codec, and then it dies by 28 assertions. Recorded because a reviewer +mutating only the obvious guard would wrongly conclude the gate is blind. + ## Stop conditions - G1 RED against the independent reference on any covered input ⇒ stop and