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..c791bf542 --- /dev/null +++ b/.agents/specs/vt-fp8-w8a8-cpu-arm.md @@ -0,0 +1,400 @@ +# 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 takes the **first half of +option 1**, and the distinction matters. #468 asks for *"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)"*. This row delivers +the registration, and the parenthetical benefit is now real at the OP tier. It is +**not sufficient to exercise the wiring**: the GEMM registered here is +`kMatmulFp8Cutlass`, while the model-layer predicate keys on +`kMatmulFp8CublasLt`, so the `mixed_scale` call sites #468 set out to cover remain +unreachable. See §Residual gap. 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 becomes a reference-tier fallback on unified-memory accelerators (`MaybeInstallReferenceTier`, `src/vt/op_provider.cpp:204-225`) | **THIS ROW WAS WRONG. See the refutation below.** | +| 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 | + +### REFUTED by measurement: the reference-tier row above understated the risk + +The risk table originally read: *"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. Recorded here as an intended, visible +consequence."* That reasoned from the ANNOUNCEMENT and never tested the +CONSEQUENCE, and the consequence is not slowness. + +`MaybeInstallReferenceTier` declines only while the CPU provider count is zero +(`op_provider.cpp:213-214`, `if (cpu.count == 0) return false;`). Before this row +that count was 0 for `kQuantFp8Static` and `kMatmulFp8Cutlass`, so a unified-memory +accelerator lacking a native kernel got the refuse-by-name path. **This row's +registrations make the count 1.** Those ops therefore flip from refusing by name to +installing a host kernel that dereferences what may be device pointers. + +Measured, not reasoned: **SIGSEGV on GB10, exit 139** ([#844](https://github.com/mudler/vllm.cpp/issues/844)). +So the real consequence is a crash, or worse a silent wrong answer, on a device +whose memory does not alias the host — and a `correct but slow` banner printed +immediately before it is not mitigation, it is a misleading label. The tier's +safety gate is `ReferenceTierEligible` (`op_provider.cpp:774-780`), not the +announcement. + +This is not fixed here. #844 owns it, it is broader than this row (every CPU +registration widens the same surface), and fixing the tier's device-pointer +contract inside a row about fp8 registration would be exactly the silent scope +widening the stop conditions forbid. Recorded as debt this row CREATES, which is +the honest framing: the row is still worth landing for CPU testability, and it +makes an existing latent defect reachable for two more ops. + +## 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. + +## 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 **21**), `SUCCESS!`. + +**The three replaced assertions are a FORCED consequence, and it is measured, not +argued.** Checking out the BASE version of that test (the one carrying the three +`CHECK_THROWS`) and running it against the new registrations gives +`test cases: 1 | 0 passed | 1 failed`, `assertions: 21 | 18 passed | 3 failed`, +`Status: FAILURE!` — exactly three failures, which are exactly the three +`CHECK_THROWS`, with the other 18 assertions untouched. The registration makes +those three claims false; it does not make them optional. They were replaced by +the stronger byte-exact form, not deleted. + +(The implementation commit message says "19 -> 25". That arithmetic was wrong; +the measured base count is 21. Corrected here rather than by rewriting pushed +history.) + +### 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!** | + +Re-run in full on the post-merge tree (merge of `origin/main` @ `5da1d7f2f`): +identical results, every row, and the restored tree green at 56/56. + +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. + +### Review findings folded in (PR #842) + +A fresh review returned FAIL on three blockers. All are addressed here rather than +argued away. + +- **F2** the branch was `CONFLICTING`. Resolved by merging `origin/main` at + `2f2bce926` and following the intake table, which main RELOCATED from + `roadmap_v1.md` to `issue-index.md`; see that merge commit for the four-way + verification and for why a strict-PREFIX check is vacuous on an in-place row edit. +- **F3** two shipped comments asserted CPU/CUDA equivalence as fact while G2 has + never run. Both downgraded to "declared and owed under G2" + (`cuda_matmul_fp8_cutlass.cu`, `include/vt/ops.h`). This was the row's own + thesis turned on the row: it repaired a comment asserting an unverified + contract, then shipped two more. +- **F4** the reference-tier risk row was refuted by measurement; see above and #844. +- **F5** the "option 1 verbatim" framing is softened wherever it appears. What + landed is the OP seam; the registered GEMM is `kMatmulFp8Cutlass` while the model + predicate keys on `kMatmulFp8CublasLt`, so the `mixed_scale` wiring #468 set out + to cover REMAINS UNREACHABLE. See §Residual gap, which said so from the start. +- **F6** the G2 case name contained a comma. doctest splits `-tc=` on commas, so + the one gate this row still owes was UNSELECTABLE by name: measured + `test cases: 0 | 0 passed | 0 failed | 4 skipped`, `Status: SUCCESS!`, exit 0. + Renamed comma-free; the same selector now returns 1 case. The lesson generalizes + past this file, so it is written next to the case rather than only here: a gate + that selects nothing and prints SUCCESS is the worst failure mode available. + +### Two defects this row found in its OWN evidence, not in the code + +Both are recorded because each made a green result mean less than it looked. + +**The gate ran a REDUCED configuration for most of this row's life.** Every full +run before the final one configured `-DVLLM_CPP_BUILD_EXAMPLES=OFF +-DVLLM_CPP_SERVER=OFF`, chosen while the box sat at 100% disk. CI configures +`cmake -S . -B build -DVLLM_CPP_BUILD_TESTS=ON` and lets both default ON. The +reduced form builds **925 targets and runs 463 tests**; the CI form builds +**1422 and runs 481**. The difference is not cosmetic: `test_minimax_music3_e2e_real` +links `ApiServer` and cannot BUILD without the server, so under the reduced flags +the gate reported `100% tests passed` while a target it never compiled was +sitting in the tree. Same class as F6 -- an instrument reporting on a state it +was not given. The binding numbers in `## Evidence` are the CI-configuration run +only. + +**The relocated intake table changed its MERGE CONTRACT, not just its path.** +`.agents/issue-index.md` is append-only and carries `merge=union` in +`.gitattributes`; its own preamble says "Never edit a row". This row moved its +`#468` edit from `roadmap_v1.md` into that file and carried the keyed-record +discipline with it, which `scripts/check-issue-index-append-only.py` correctly +refused. The review had prescribed "the target's file is a strict PREFIX of the +result", which is precisely the append-only test, and this row substituted a +weaker in-place-diff check on the grounds that a prefix test was vacuous for an +edit-in-place. The prescribed check was right and the adaptation was wrong. The +resolution is to make no edit at all: main's `#468` row already carries the +linkage AGENTS.md requires. + +## 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. 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..20e0926f4 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,16 @@ 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 is EXPECTED to +// agree with the CUDA kernel to fp8/bf16 tolerance and NOT byte-for-byte, because +// the CUDA arm reduces K in tensor-core order and rounds its epilogue through +// bf16 — but that agreement is DECLARED AND OWED, not measured. No committed run +// has compared the two arms: gate G2 of .agents/specs/vt-fp8-w8a8-cpu-arm.md is +// PENDING for want of a GPU. Treat the tolerance above as the claim to be tested, +// not as a result. 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..68598f8ba 100644 --- a/src/vt/cuda/cuda_matmul_fp8_cutlass.cu +++ b/src/vt/cuda/cuda_matmul_fp8_cutlass.cu @@ -317,8 +317,25 @@ 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 INTENDED to be the +// byte-for-byte mirror of this kernel, and that equivalence is DECLARED AND OWED, +// not measured. It is gate G2 of .agents/specs/vt-fp8-w8a8-cpu-arm.md, which is +// PENDING for want of a GPU (#468). What IS measured is weaker and lives on the +// CPU side: G1 proves the CPU kernel matches an independent e4m3 reference derived +// from the format. Two implementations each matching a reference is not the same +// claim as the two matching each other, so do not cite this comment as evidence +// that they agree. Run tests/vt/test_ops_fp8_cpu.cpp on a CUDA host to close it. __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 77d8f01c3..0f97f648e 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1537,6 +1537,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..e3cd4c16c --- /dev/null +++ b/tests/vt/test_ops_fp8_cpu.cpp @@ -0,0 +1,453 @@ +// 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)); + + // 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); + 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. +// NAME THIS CASE WITHOUT A COMMA. doctest splits `-tc=` on commas, so a comma in +// a case name makes the name unselectable: the filter becomes two patterns that +// each match nothing, and the binary then reports +// test cases: 0 | 0 passed | 0 failed | 4 skipped ... Status: SUCCESS! +// with exit 0. That is the worst failure mode available to a gate -- the ONE +// arm this row still owes would have selected nothing and reported success. +// Measured on this file before the rename (#468 review F6). +TEST_CASE("G2: CPU QuantFp8Static equals 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)); +}