From 38336aca20f6bc8cb92fcefdaf3e00f520a82882 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 07:26:51 +0000 Subject: [PATCH 1/2] spec(VT-QUANT-FP8-GROUP): the executing per-group FP8 quant divides twice, and only a byte comparison can see it (#1189) Block-wise FP8 needs an activation quantizer before it needs anything else, and #1189 splits it out as milestone M1. This commit is that milestone's spec, committed before its implementation. The design pass that matters is which upstream kernel to mirror. `per_token_group_quant_fp8` looks like a Triton kernel with a C++ fast path, and it is the other way round: `vllm/model_executor/layers/quantization/utils/fp8_utils.py:635-650` calls `torch.ops._C.per_token_group_fp8_quant` and returns whenever the platform is CUDA-alike and the input is contiguous, which is every case this row cares about. The Triton kernel below it never runs there. The two arms are not interchangeable. The C++ kernel divides, at `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:68` for the scale and `:85` for the value. The Triton kernel multiplies by `(1.0 / fp8_max)` at `fp8_utils.py:145`, under a comment that names the 1-ULP difference this produces. One f32 ULP before an e4m3 round changes the emitted byte near a tie, and upstream's own test cannot see the difference: it compares values at `rtol=0.15` (`tests/kernels/quantization/test_block_fp8.py:112-114`). So the spec declares a bitwise gate against an independently written reference on top of the ported test, because the ported test alone would pass either arm. The spec also records what M1 does not do. It lands unreached, M4 owns the wiring, and the CUDA arm compiles without running because the row takes no GPU lease. Both are listed under `## Owed`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/issue-index.md | 1 + .agents/specs/vt-quant-fp8-group.md | 216 ++++++++++++++++++++++++++++ 2 files changed, 217 insertions(+) create mode 100644 .agents/specs/vt-quant-fp8-group.md diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 83876d612..f6aae10fc 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,3 +366,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | +| [#1189](https://github.com/mudler/vllm.cpp/issues/1189) | `VT-QUANT-FP8-GROUP` | Block-wise (128x128) FP8 so `Qwen/Qwen3.8-27B-FP8` runs instead of being refused. `weight_block_size` appears nowhere in `src/` or `include/`, so `469f38395` refuses the arm by name (#1166). Six independently landable milestones; M1 lands here. M1 is `vt::QuantFp8Group`, the dynamic per-token per-group activation quant, CPU and CUDA. The numerics mirror the kernel that ACTUALLY EXECUTES on a CUDA-alike platform with a contiguous input, which is the C++ custom op at `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:42-96` and NOT the Triton kernel at `fp8_utils.py:95-150`: `fp8_utils.py:635-650` calls the former and returns before the latter. The two arms differ, and the difference is measurable rather than cosmetic. The CUDA kernel divides twice, `local_absmax / max_8bit` at `:68` and `static_cast(src) / y_s` at `:85`; the Triton kernel multiplies by `(1.0 / fp8_max)` at `fp8_utils.py:145` under a comment that names the 1-ULP gap. Upstream's own test tolerates the gap with `rtol=0.15` (`test_block_fp8.py:112-114`), so a value comparison cannot tell the two apart and only a byte comparison against a spelled-out reference can. `eps` is the reduction's INITIAL value (`:47`), not a post-clamp, which is what keeps an all-zero group from dividing by zero. Scope refused here and owed to later milestones: the block-scaled GEMM (M2), `Fp8BlockWeight` and the loader (M3), `Fp8BlockLinearMethod` and the Qwen3.5 wiring (M4), the mainloop-scaled CUTLASS kernel and the column-major/TMA-aligned scale layouts (M5), merged `gate_up`/QKV (M6). M1 lands UNREACHED: no production entry point dispatches `vt::QuantFp8Group` at its merge commit, M4 owns the wiring, and `.agents/specs/vt-quant-fp8-group.md` lists it under `## Owed`. The CUDA arm compiles and its on-hardware leg is owed too, because the row took no GPU lease by design: the CPU arm is the gateable one | feature | diff --git a/.agents/specs/vt-quant-fp8-group.md b/.agents/specs/vt-quant-fp8-group.md new file mode 100644 index 000000000..b91912641 --- /dev/null +++ b/.agents/specs/vt-quant-fp8-group.md @@ -0,0 +1,216 @@ +# VT-QUANT-FP8-GROUP — dynamic per-token, per-group FP8 activation quantization + +Issue: [#1189](https://github.com/mudler/vllm.cpp/issues/1189), milestone **M1**. +Row: `VT-QUANT-FP8-GROUP`. +Pinned oracle: vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` +(`.agents/upstream-sync.md`), HEAD of the local checkout verified before every +`file:line` below was read. + +## Scope + +Add one `vt` op, `vt::QuantFp8Group`, with a CPU kernel and a CUDA kernel: + +```c++ +void QuantFp8Group(Queue& q, Tensor& out_fp8, Tensor& out_scale, + const Tensor& x, int group_size); +``` + +`x` is `[M, K]` f32 or bf16. `out_fp8` is `[M, K]` i8 that carries raw +fp8-e4m3fn bytes. `out_scale` is `[M, K / group_size]` **f32**. The op refuses +`K % group_size != 0` by name. + +This is the activation half of block-wise FP8. It is milestone M1 of #1189 and +it deliberately stops there. + +**Out of scope, each owned by a later milestone of #1189**: the block-scaled +GEMM (M2), `Fp8BlockWeight` and the loader rung (M3), `Fp8BlockLinearMethod` +and the Qwen3.5 wiring (M4), the mainloop-scaled CUTLASS kernel (M5), and the +merged `gate_up` / QKV projections (M6). + +## Upstream anchors + +Read the whole executing chain, not the top-level Python. On a CUDA platform +with a contiguous input, `per_token_group_quant_fp8` does **not** run its Triton +kernel: it calls the C++ custom op and returns +(`vllm/model_executor/layers/quantization/utils/fp8_utils.py:635-650`). The +Triton kernel below it is the fallback for every other platform. + +| What | Where | +|---|---| +| Python entry point, defaults, refusals | `vllm/model_executor/layers/quantization/utils/fp8_utils.py:567-650` | +| the divisibility refusal this op mirrors | `fp8_utils.py:596-599` | +| the contiguity refusal | `fp8_utils.py:600` | +| row-major scale allocation `[M, K/group_size]` f32 | `fp8_utils.py:629-631` | +| CUDA dispatch, taken whenever the platform is CUDA-alike and `x` is contiguous | `fp8_utils.py:635-650` | +| **the executing CUDA kernel**: group absmax and scale | `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:42-74` | +| **the executing CUDA kernel**: divide, clamp, e4m3 store | `per_token_group_quant.cu:76-96` | +| the Triton fallback, for contrast only | `fp8_utils.py:95-150` | +| `fp8_min`, `fp8_max` = `finfo(e4m3fn)` = -448, 448 | `vllm/model_executor/layers/quantization/utils/quant_utils.py:27-35` | +| the native torch reference the upstream test compares against | `tests/kernels/quant_utils.py:157-180` | +| the ported test | `tests/kernels/quantization/test_block_fp8.py:82-118`, grid at `:42-46` | + +## Design + +### Numerics + +Per group of `group_size` contiguous elements of one row: + +```text +amax = max(eps, max |x_f32| over the group) eps = 1e-10 +y_s = amax / 448.0f +q = min(max(x_f32 / y_s, -448.0f), 448.0f) +out = e4m3fn(q), round to nearest even, saturating +``` + +Three details are deliberate, and each is a divergence risk if a later reader +"corrects" it. + +**A divide, not a reciprocal multiply.** `vt::QuantFp8Static` multiplies by a +hoisted reciprocal because that is the form upstream ships for the per-tensor +static path (`csrc/quantization/w8a8/fp8/common.cuh:62`). The per-group path is +the opposite. The executing CUDA kernel writes `float y_s = local_absmax / +max_8bit` (`per_token_group_quant.cu:68`) and `static_cast(src) / y_s` +(`per_token_group_quant.cu:85`). Both are true divisions, and the scale changes +per group, so there is no loop-invariant reciprocal to hoist. The Triton +fallback instead writes `scale_raw = _absmax * (1.0 / fp8_max)` +(`fp8_utils.py:145`) with a comment that names the 1-ULP difference. The two +upstream arms therefore disagree by up to one f32 ULP in `y_s`, and upstream's +own test admits that with `rtol=0.15` on the values +(`test_block_fp8.py:112-114`). We mirror the CUDA arm, because that is the arm +that executes on the target architecture. + +**`eps` is the initial value of the reduction, not a post-clamp.** The kernel +sets `float local_absmax = eps` (`per_token_group_quant.cu:47`) and reduces +`fmaxf` over it. That is identical to `clamp(min=eps)` in the native reference +(`quant_utils.py:176`), and it makes an all-zero group produce +`y_s = 1e-10 / 448`, not a division by zero. + +**The load widens to f32 before the absolute value and before the divide.** +`fabsf(static_cast(src))` (`per_token_group_quant.cu:53`) and +`static_cast(src) / y_s` (`:85`). A bf16 input therefore rounds at +exactly one point, as it does in `vt::QuantFp8Static`. + +### Layout and what it excludes + +`out_scale` is row-major `[M, K/group_size]` f32, upstream's `column_major_scales += False` branch (`fp8_utils.py:629-631`). The column-major and TMA-aligned +layouts (`fp8_utils.py:610-628`) exist for the DeepGEMM and CUTLASS GEMMs. No +consumer in this tree can read them yet, so shipping them now would ship a +parameter no caller passes. They are owed below. + +`use_ue8m0` is excluded, not forgotten. It rounds the scale up to a power of two +for DeepGEMM (`per_token_group_quant.cu:69-71`). Issue #1189 established that +upstream excludes `qwen3_5_text` from DeepGEMM on family 120 +(`vllm/utils/deep_gemm.py:27-46`) and dispatches CUTLASS, so the target path +never sets it. + +`eps` stays a named constant rather than a parameter. Upstream exposes it, and +every upstream call site takes the `1e-10` default (`fp8_utils.py:570`). + +### Structure + +`vt::ScaledFp4Quant` (`include/vt/ops.h:1425-1427`) is the closest existing op: +it is dynamic, per-token, per-group, and it emits a 2-D scale beside the packed +values. `QuantFp8Group` follows its shape. `OpId::kQuantFp8Group` is appended +before `kCount`, the additive convention documented at `include/vt/ops.h:363-368`, +so no existing op id shifts. + +The CUDA arm lives in `src/vt/cuda/cuda_quant_fp8.cu`. That file exists because +`vt::QuantFp8Static` once lived in the cutlass-gated translation unit and was +therefore unregistered on every CUDA architecture outside `VT_CUTLASS_FP8_ARCHS` +(#960, and #844 is the same defect seen from the fallback's end). This kernel +has the same property: it is a divide and a convert, with no cutlass dependency, +and it must resolve on every CUDA architecture. It is added to +`scripts/check-cuda-op-arch-gate.py`'s `REQUIRED` set for that reason. + +## Risks + +| Risk | Control | +|---|---| +| a later reader converts the divide to a hoisted reciprocal, matching the Triton arm and moving emitted bytes near an e4m3 tie | the prose above, plus G1, which compares bytes against a reference that spells the divide out | +| the scale silently widens or narrows | G1 asserts the f32 dtype and the `[M, K/group_size]` shape; the op refuses any other dtype | +| an all-zero group divides by zero | `eps` is the reduction's initial value; G4 feeds an all-zero row | +| a ragged `K` is accepted and reads past the row | the op refuses `K % group_size != 0` by name; G5 asserts the refusal | +| an all-zero output passes every value comparison | G2 and G3 carry a `nonzero == numel` vacuity guard | +| the CPU and CUDA arms drift apart | G6, which is CUDA-gated and currently **owed** | + +## Tests + +`tests/vt/test_ops_quant_fp8_group_cpu.cpp`, registered in +`tests/CMakeLists.txt`. + +- **G1** bitwise, zero tolerance. The CPU op equals an independently written + reference, byte for byte, over a grid that includes saturation in both signs, + the subnormal ladder, exact e4m3 ties, and both zeros. The reference encodes + by an exhaustive nearest-value scan over the 128 finite e4m3fn magnitudes, + which is a different algorithm from the tree's `F32ToFp8`, so agreement is + evidence rather than a tautology. +- **G2** the ported upstream case. `test_block_fp8.py:82-118` with its grid + (`num_tokens` in {7, 2050}, `d` in {512, 4096, 5120, 13824}, `group_size` in + {64, 128, 512}, bf16 and f32), against the native reference transcribed from + `quant_utils.py:157-180`, at upstream's tolerances: `rtol=0.15` on the + dequantized values and `allclose` on the scale. Carries the vacuity guard. +- **G3** the scale is exactly `amax / 448` for a group whose amax is known by + construction, and `1e-10 / 448` for an all-zero group. +- **G4** shape and dtype contract: `out_scale` is f32 `[M, K/group_size]`, and + a wrong scale dtype or shape is refused. +- **G5** the refusals: `K % group_size != 0`, a non-contiguous input, a + device mismatch, each by name. +- **G6** CPU against CUDA, byte for byte, on the identical input. **Owed**: this + host has no GPU and the row took no lease. The case is written and it never + reports a silent skip; without a device it asserts the CPU registration and + prints a banner that names what was not measured, following the G2 precedent + at `tests/vt/test_ops_fp8_cpu.cpp:279`. + +## Gates + +| Gate | Command | +|---|---| +| focused | `ctest -R test_ops_quant_fp8_group_cpu --output-on-failure` | +| op provider totality | `ctest -R test_op_provider` | +| the per-tensor sibling, unchanged | `ctest -R test_ops_fp8_cpu` | +| structural | `python3 scripts/check-cuda-op-arch-gate.py` | +| record | `scripts/agent-preflight.sh --fail-on-skip` | + +The CUDA arm compiles in this change. It does not run in this change. + +## Owed + +- The CUDA arm's on-hardware leg. The kernel is written and it compiles, and + G6 measures nothing without a device. Owed by #1189 milestone M5, which needs + a GPU anyway; the case is already written, so the debt is a run and not a + test. +- **Nothing reaches this op yet.** `vt::QuantFp8Group` is dispatched by no + production entry point at this merge commit: `include/vllm.h` does not expose + it, no loader builds a `Fp8BlockWeight`, and `ModelRegistry::Forward` has no + block-FP8 linear method to call it from. The wiring is owned by #1189 + milestone M4 (`Fp8BlockLinearMethod` and the Qwen3.5 dense forward), which + needs M2 and M3 first. This is the staged-slice exception of + `.agents/reachability.md`, named here, in the commit body, and in the pull + request body. +- The column-major and TMA-aligned scale layouts (`fp8_utils.py:610-628`). + Owed by #1189 milestone M5, which is the first consumer that can read them. +- `use_ue8m0` scale rounding (`per_token_group_quant.cu:69-71`). Not owed by a + milestone: the target architecture never selects DeepGEMM + (`vllm/utils/deep_gemm.py:27-46`). It becomes owed when a DeepGEMM path lands. + +## Stop conditions + +Stop and report `NEEDS_DECISION` if any of the following holds. + +- The pinned oracle's checkout is not at + `5559679229bc961848b121ccdeaa8fa5d79bec98`. Every anchor above was read at + that revision. +- G1 cannot be made bitwise green without widening its tolerance. A tolerance + on a byte comparison is the defect, not the fix. +- The op cannot express upstream's contract without a parameter that no caller + in this tree passes. + +Stop and report `NEEDS_CONTEXT` if the work requires a GPU lease. The row is +scoped so that it does not. + +## Evidence + +Recorded in the pull request body: the RED capture before the implementation +existed, the GREEN capture after, and the per-block gate counts from +`scripts/agent-preflight.sh --fail-on-skip`. From 43220d0ee03728fc2c0d91b97de1b28bf62c4cf1 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Tue, 18 Aug 2026 07:55:49 +0000 Subject: [PATCH 2/2] feat(VT-QUANT-FP8-GROUP): add vt::QuantFp8Group, the dynamic per-token per-group fp8 activation quant, CPU and CUDA (#1189) Block-wise FP8 needs an activation quantizer, and this is it: milestone M1 of #1189. `vt::QuantFp8Group` takes x [M,K] f32 or bf16 and emits the fp8 bytes plus an f32 [M, K/group_size] scale, one scale per contiguous run of `group_size` elements inside a row. WHICH UPSTREAM ARM THIS MIRRORS, because there are two and they disagree. `per_token_group_quant_fp8` reads like a Triton kernel with a C++ fast path and it is the other way round: on a CUDA-alike platform with a contiguous input it calls `torch.ops._C.per_token_group_fp8_quant` and returns (fp8_utils.py:635-650), so the Triton kernel never executes there. The executing kernel is csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu, and it divides twice: `local_absmax / max_8bit` at :68 and `static_cast(src) / y_s` at :85. The Triton fallback instead forms `_absmax * (1.0 / fp8_max)` at fp8_utils.py:145, under an upstream comment that names the 1-ULP difference this opens. This is the opposite polarity from vt::QuantFp8Static, which multiplies by a hoisted reciprocal because that IS upstream's shipped form for the static per-tensor path. Both kernels carry the reason beside the code, because each looks like a defect from the other's point of view. `eps` seeds the reduction rather than clamping it afterwards (per_token_group_quant.cu:47). The two are numerically identical and upstream's form makes it visible that an all-zero group yields 1e-10/448 instead of dividing by zero. WHY THE PORTED TEST IS NOT ENOUGH, measured rather than argued. G2 ports upstream's case with its full grid and upstream's tolerances: values at rtol=0.15, the scale at torch.allclose's rtol=1e-5 (test_block_fp8.py:112-115). Mutating this kernel to the Triton arm's scale form makes G2 fail 6 of its 48 shape checks -- every one of them at num_tokens=2050, none at num_tokens=7, and never the scale check, since a 1-ULP scale difference is about 6e-8 relative. Mutating it to the Triton arm's value form makes G2 pass 50 of 50. So upstream's tolerances catch one of the two forms, on the large shapes, by luck of which element lands on an e4m3 boundary, and miss the other entirely. G1 catches both on every shape because it compares BYTES, against a reference derived from the e4m3fn format by exhaustive nearest-value scan -- a different algorithm from this tree's F32ToFp8 rather than a restatement of it. The CUDA arm gives one thread the whole group instead of upstream's 16-lane shuffle reduction. That cannot change the result: fmaxf is exact and order-independent over finite inputs, unlike a floating-point sum. A lane-parallel rewrite is a performance question for M5, which needs a GPU to measure. `scripts/check-cuda-op-arch-gate.py` now pins kQuantFp8Group as well. The kernel is a max, two divides and a hardware convert, with no cutlass dependency, so a CUDA queue must never fall through to the host reference tier and dereference device pointers -- the defect #960 and #844 record. Its miniature fixture gains the matching registration, because the fixture describes the real TU and the real TU now has two. NOTHING REACHES THIS OP YET, and that is deliberate. No production entry point dispatches vt::QuantFp8Group at this commit: include/vllm.h does not expose it, no loader builds an Fp8BlockWeight, and ModelRegistry::Forward has no block-FP8 linear method to call it from. The wiring is owned by #1189 milestone M4, which needs M2's GEMM and M3's loader rung first. This is the staged-slice exception of .agents/reachability.md and the row's spec lists it under `## Owed`, together with the CUDA arm's on-hardware leg. That kernel compiles in the cuda-fat-build CI job, which builds the vllm target with VLLM_CPP_CUDA=ON for ten architectures on every pull request, and it executes nowhere: no lane runs a CUDA kernel and this host has neither a toolkit nor a device. G6 prints a PENDING banner naming what was not measured instead of skipping silently. The run belongs to M5, which needs a GPU regardless. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code] --- .agents/specs/vt-quant-fp8-group.md | 77 ++- include/vt/ops.h | 61 ++ scripts/check-cuda-op-arch-gate.py | 6 + src/vt/cpu/cpu_ops.cpp | 66 ++ src/vt/cuda/cuda_quant_fp8.cu | 84 +++ src/vt/op_provider.cpp | 2 + src/vt/ops.cpp | 31 + tests/CMakeLists.txt | 4 + tests/scripts/test_check_cuda_op_arch_gate.py | 13 +- tests/vt/test_ops_quant_fp8_group_cpu.cpp | 640 ++++++++++++++++++ 10 files changed, 975 insertions(+), 9 deletions(-) create mode 100644 tests/vt/test_ops_quant_fp8_group_cpu.cpp diff --git a/.agents/specs/vt-quant-fp8-group.md b/.agents/specs/vt-quant-fp8-group.md index b91912641..4350f8121 100644 --- a/.agents/specs/vt-quant-fp8-group.md +++ b/.agents/specs/vt-quant-fp8-group.md @@ -133,6 +133,8 @@ and it must resolve on every CUDA architecture. It is added to | a ragged `K` is accepted and reads past the row | the op refuses `K % group_size != 0` by name; G5 asserts the refusal | | an all-zero output passes every value comparison | G2 and G3 carry a `nonzero == numel` vacuity guard | | the CPU and CUDA arms drift apart | G6, which is CUDA-gated and currently **owed** | +| the CUDA arm does not compile, and this host cannot say so | the `cuda-fat-build` CI job compiles it on the pull request; its verdict is the gate, and this row did not take that measurement itself | +| the CUDA arm compiles and computes the wrong bytes | **not covered here.** G6 is the instrument and it needs a device. Named under `## Owed` | ## Tests @@ -172,14 +174,28 @@ and it must resolve on every CUDA architecture. It is added to | structural | `python3 scripts/check-cuda-op-arch-gate.py` | | record | `scripts/agent-preflight.sh --fail-on-skip` | -The CUDA arm compiles in this change. It does not run in this change. +**Where the CUDA arm is compiled, and where it is not.** The implementing host +has no CUDA toolkit (`nvcc` is absent), so this row did not compile the kernel +locally. Continuous integration does: the `cuda-fat-build` job +(`.github/workflows/ci.yml:669-710`) configures `-DVLLM_CPP_CUDA=ON` for ten +architectures in an `nvidia/cuda` devel container and builds the `vllm` target +on every non-closed `pull_request` event, and `src/vt/cuda/cuda_quant_fp8.cu` is +in the unconditional CUDA source list, so the compile leg is gated there. Its +verdict is the pull request's, not a measurement this row took before opening +it. + +Nothing runs the kernel. That job states it uses no GPU, and no other lane +executes a CUDA kernel, so the arm's first execution belongs to #1189 milestone +M5. `## Owed` records that. ## Owed -- The CUDA arm's on-hardware leg. The kernel is written and it compiles, and - G6 measures nothing without a device. Owed by #1189 milestone M5, which needs - a GPU anyway; the case is already written, so the debt is a run and not a - test. +- **The CUDA arm's on-hardware leg.** The kernel compiles in the + `cuda-fat-build` continuous-integration job and executes nowhere: no lane runs + a CUDA kernel, and this host has neither a toolkit nor a device. G6 is written + and prints a `PENDING` banner that names what was not measured. Owed by #1189 + milestone M5, which needs a GPU anyway. The debt is a run, not a test: the + case selects itself as soon as a device exists. - **Nothing reaches this op yet.** `vt::QuantFp8Group` is dispatched by no production entry point at this merge commit: `include/vllm.h` does not expose it, no loader builds a `Fp8BlockWeight`, and `ModelRegistry::Forward` has no @@ -211,6 +227,51 @@ scoped so that it does not. ## Evidence -Recorded in the pull request body: the RED capture before the implementation -existed, the GREEN capture after, and the per-block gate counts from -`scripts/agent-preflight.sh --fail-on-skip`. +RED before the implementation existed: the focused build failed with +``'QuantFp8Group' is not a member of 'vt'`` and ``'kQuantFp8Group' is not a +member of 'vt::OpId'`` at 12 sites. + +GREEN after: `test_ops_quant_fp8_group_cpu` reports 6 cases, 476 assertions, 0 +failed. `scripts/agent-preflight.sh --fail-on-skip` reports **All gates green** +with no `FAIL` and no `SKIP`. + +### Mutation results + +Every mutation prints `git diff --stat` and `compile_rc`, because a mutation +that never applied and a mutation that failed to build both read as a passing +test. Two of the ten did exactly that and are reported here rather than +dropped. + +| Mutation | `compile_rc` | Result | +|---|---|---| +| the Triton arm's scale form, `amax * (1/448)` | 0 | G1 fails 49/146; G2 fails 6/48 shape checks | +| the Triton arm's value form, `x * (1/y_s)` | 0 | G1 fails 14/146; **G2 passes 50/50** | +| `float amax = 0.0F` | **1** | proves nothing: `-Werror=unused-variable` on `kEps` | +| the same with `kEps` kept live | 0 | G3 fails 260/269; G1 passes | +| per-group scale collapsed to per-row | 0 | G3 fails 7/269; G1 fails 60/146 | +| the divisibility refusal, applied with a shell-escaped pattern | n/a | **never applied**; the harness reported `MUTATION_NOT_UNIQUE count=0` and the case then read `SUCCESS` | +| the divisibility refusal, applied | 0 | G5 fails | +| `out_scale` f32 refusal widened to any float | 0 | G4 fails | +| the CPU registration deleted | **1** | proves nothing: `-Werror=unused-function` | +| the same with the kernel kept live | 0 | G1 fails at its `REQUIRE`; **G5 still passes** | + +Three results are worth keeping. + +**Upstream's tolerances are not a substitute for a byte comparison, and the +measurement is more specific than the argument for it.** The scale assertion at +`rtol=1e-5` never fired for either 1-ULP form: a 1-ULP scale difference is about +6e-8 relative. The value assertion at `rtol=0.15` fired for the scale-form +change in 6 of 24 shapes, every one of them at `num_tokens=2050` and none at +`num_tokens=7`, which is sample size deciding whether any element lands on an +e4m3 boundary. It did not fire at all for the value-divide change. G1 fired for +both, on every shape. + +**A refusal test cannot stand in for a registration test.** With the CPU +registration deleted, G5 still passed: the refusals live in `vt::QuantFp8Group`'s +validation and fire before dispatch. Only G1's `REQUIRE(OpRegistered(...))` +caught it. + +**`-Werror` turns two natural mutations into non-events.** Deleting the eps seed +orphans `kEps`, and deleting the registration orphans the kernel. Both fail to +build, and a harness that reported only the case verdict would have recorded two +passes. Each was re-run in a form that keeps the symbol live. diff --git a/include/vt/ops.h b/include/vt/ops.h index 70e015f9f..7905f4da4 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -423,6 +423,16 @@ enum class OpId : uint8_t { // Appended before kCount so no existing op's id shifts. kConv1d, kConvTranspose1d, + // --- Block-wise FP8 (VT-QUANT-FP8-GROUP, #1189 milestone M1). The DYNAMIC + // per-token, per-group fp8 activation quant that a 128x128 block-scaled FP8 + // GEMM consumes. It is not a parameter of kQuantFp8Static: that op takes ONE + // static per-tensor scale from the checkpoint and emits no scale tensor at + // all, while this one derives a scale per (row, group) at run time and emits + // an f32 [M, K/group_size] second output. The nearest existing shape is + // kScaledFp4Quant, which is dynamic per-token with a 2-D scale as well. + // See vt::QuantFp8Group below for the contract. + // Appended before kCount so no existing op's id shifts. + kQuantFp8Group, kCount }; @@ -922,6 +932,10 @@ using MatmulFp8CublasLtFn = using MatmulFp8CublasLtAlphaVecFn = void (*)(Queue&, Tensor&, const Tensor&, const Tensor&, const Tensor& /*alpha_vec*/, bool); using QuantFp8StaticFn = void (*)(Queue&, Tensor&, const Tensor&, float); +// Two outputs, because the scale is computed rather than supplied: the fp8 bytes +// and the f32 [M, K/group_size] per-group scale. +using QuantFp8GroupFn = void (*)(Queue&, Tensor& /*out_fp8*/, Tensor& /*out_scale*/, + const Tensor& /*x*/, int /*group_size*/); using RmsNormQuantFp8Fn = void (*)(Queue&, Tensor& /*out_fp8*/, Tensor* /*out_bf16*/, const Tensor& /*x*/, const Tensor& /*weight*/, const RmsNormArgs&, Tensor* /*residual*/, float /*input_scale*/); @@ -1521,6 +1535,53 @@ void MatmulNvfp4Cutlass(Queue& q, Tensor& out, const Tensor& a_packed, const Ten // that makes the fp8 seam testable without a GPU, #468). void QuantFp8Static(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scale); +// --- Block-wise FP8 (VT-QUANT-FP8-GROUP, #1189 M1, +// .agents/specs/vt-quant-fp8-group.md). QuantFp8Group is the DYNAMIC per-token, +// per-group sibling of QuantFp8Static: the scale is derived from the data, once +// per contiguous run of `group_size` elements inside a row, and written out. +// +// amax = max(1e-10, max |x_f32| over the group) +// y_s = amax / 448.0f +// out_fp8[i] = fp8_e4m3( min(max(x_f32[i] / y_s, -448.0f), 448.0f) ) +// out_scale[row, group] = y_s +// +// x [M,K] f32/bf16, out_fp8 [M,K] i8 (raw fp8-e4m3fn bytes), out_scale +// [M, K/group_size] F32 — f32 because upstream allocates it f32 and the GEMM +// that consumes it multiplies in f32 (fp8_utils.py:629-631). K must be a +// multiple of group_size; the op refuses any other K by name, as upstream +// asserts at fp8_utils.py:596-599. CPU + CUDA. +// +// WHICH UPSTREAM ARM THIS MIRRORS, because there are two and they DISAGREE. +// `per_token_group_quant_fp8` reads like a Triton kernel with a C++ fast path +// and it is the other way round: on a CUDA-alike platform with a contiguous +// input it calls the C++ custom op and RETURNS (fp8_utils.py:635-650), so the +// Triton kernel never executes there. The executing kernel is +// csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu: +// :47 float local_absmax = eps — eps SEEDS the reduction, so an +// all-zero group cannot divide by 0 +// :53 fmaxf(local_absmax, fabsf((float)src)) +// :68 float y_s = local_absmax / max_8bit — a DIVIDE +// :85 fminf(fmaxf((float)src / y_s, min_8bit), max_8bit) — a DIVIDE +// :86 DST_DTYPE(q) — hardware e4m3 RNE, saturating +// DO NOT "correct" either divide into a hoisted reciprocal multiply to match +// QuantFp8Static's form. That form is right for QuantFp8Static because upstream +// ships it there (common.cuh:62 with the reciprocal formed by the caller); here +// upstream ships a divide, the scale changes per group so nothing is +// loop-invariant, and the Triton fallback's `_absmax * (1.0 / fp8_max)` +// (fp8_utils.py:145) carries an upstream comment naming the 1-ULP gap it opens. +// One f32 ulp before an e4m3 round changes the emitted byte near a tie, and +// upstream's own test cannot see it: it compares values at rtol=0.15 and the +// scale at rtol=1e-5 (test_block_fp8.py:112-115). Only a byte comparison can, +// which is what tests/vt/test_ops_quant_fp8_group_cpu.cpp G1 is. +// +// The column-major and TMA-aligned scale layouts (fp8_utils.py:610-628) and the +// `use_ue8m0` DeepGEMM scale rounding (per_token_group_quant.cu:69-71) are NOT +// implemented. Both are recorded under `## Owed` in the row's spec; no consumer +// in this tree can read either yet, and upstream excludes the target model from +// DeepGEMM on family 120 (vllm/utils/deep_gemm.py:27-46). +void QuantFp8Group(Queue& q, Tensor& out_fp8, Tensor& out_scale, const Tensor& x, + int group_size); + // RmsNormQuantFp8 (fused fp8 RMSNorm -> static per-tensor activation quant). One // HBM pass mirrors vLLM's Inductor `fused_add_rms_norm_static_fp8_quant` // (vllm/compilation/passes/fusion/rms_quant_fusion.py:124) — the RMSNorm producer diff --git a/scripts/check-cuda-op-arch-gate.py b/scripts/check-cuda-op-arch-gate.py index 8cc458195..69bac6b5c 100644 --- a/scripts/check-cuda-op-arch-gate.py +++ b/scripts/check-cuda-op-arch-gate.py @@ -92,6 +92,12 @@ # trapped in the cutlass-fp8 TU; unreachable on every non-cutlass-fp8 CUDA # arch, where it fell to the reference tier and segfaulted (#844). ("kQuantFp8Static", "src/vt/cuda/cuda_quant_fp8.cu", "#960"), + # #1189 M1: the dynamic per-token, per-group fp8 activation quant. Same + # property as its neighbour and the same consequence if it moves -- a max, + # two divides and a hardware e4m3 convert, with no cutlass dependency, so a + # CUDA queue must never fall through to the host reference tier and + # dereference device pointers. It shares the TU deliberately. + ("kQuantFp8Group", "src/vt/cuda/cuda_quant_fp8.cu", "#1189"), ) # Where a stray duplicate registration could hide. Every CUDA-side source. diff --git a/src/vt/cpu/cpu_ops.cpp b/src/vt/cpu/cpu_ops.cpp index eb416f7cf..fcb0afd21 100644 --- a/src/vt/cpu/cpu_ops.cpp +++ b/src/vt/cpu/cpu_ops.cpp @@ -564,6 +564,70 @@ void QuantFp8StaticKernel(Queue&, Tensor& out_fp8, const Tensor& x, float input_ }); } +// --- Block-wise FP8 (VT-QUANT-FP8-GROUP, #1189 M1). QuantFp8Group CPU kernel: +// the DYNAMIC per-token, per-group activation quant. +// +// MIRROR OF THE KERNEL THAT ACTUALLY EXECUTES, which is the C++ custom op and +// not the Triton kernel. `per_token_group_quant_fp8` calls +// `torch.ops._C.per_token_group_fp8_quant` and RETURNS whenever the platform is +// CUDA-alike and the input is contiguous +// (vllm/model_executor/layers/quantization/utils/fp8_utils.py:635-650), so the +// Triton kernel below it never runs there. The executing kernel is +// csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu: +// :47 float local_absmax = eps eps SEEDS the reduction +// :53 fmaxf(local_absmax, fabsf((float)src)) +// :68 float y_s = local_absmax / max_8bit a DIVIDE +// :85 fminf(fmaxf((float)src / y_s, min_8bit), max_8bit) a DIVIDE +// :86 DST_DTYPE(q) hardware e4m3 RNE +// +// TWO DIVIDES, DELIBERATELY, and this is the opposite of QuantFp8StaticKernel +// forty lines above. That kernel multiplies by a hoisted reciprocal because +// upstream ships the reciprocal there (common.cuh:62, with `1.0f / scale` +// formed by the caller at common.cu:31). Here upstream ships a divide, and the +// scale changes per group, so there is no loop-invariant reciprocal to hoist in +// the first place. The Triton fallback's `_absmax * (1.0 / fp8_max)` +// (fp8_utils.py:145) differs by up to one f32 ulp and carries an upstream +// comment saying so. Near an e4m3 tie that ulp changes the emitted byte. +// tests/vt/test_ops_quant_fp8_group_cpu.cpp G1 compares BYTES for this reason; +// upstream's own test compares values at rtol=0.15 and cannot see it. +// +// eps is the reduction's INITIAL value rather than a clamp afterwards. The two +// are numerically identical, and writing it upstream's way makes it visible +// that an all-zero group yields y_s = 1e-10/448 instead of dividing by zero. +// +// LoadF32 widens a bf16 x to f32 before the absolute value and before the +// divide, matching `fabsf(static_cast(src))` at :53 and +// `static_cast(src) / y_s` at :85, so a bf16 input rounds at one point. +// Parallel over ROWS: each row's groups are independent, and the reduction +// order inside a group is fixed and sequential, so the result does not depend +// on the thread count. +void QuantFp8GroupKernel(Queue&, Tensor& out_fp8, Tensor& out_scale, const Tensor& x, + int group_size) { + const int64_t m = x.shape[0], k = x.shape[1]; + const int64_t groups = k / group_size; + constexpr float kEps = 1e-10F; // fp8_utils.py:570, the only value any + // upstream call site passes + constexpr float kFp8MaxV = 448.0F; // quant_utils.py:27-35 finfo(e4m3fn).max + constexpr float kFp8MinV = -448.0F; + uint8_t* op = out_fp8.Ptr(); + float* sp = out_scale.Ptr(); + ForRows(m, [&](int64_t r0, int64_t r1) { + for (int64_t r = r0; r < r1; ++r) { + for (int64_t g = 0; g < groups; ++g) { + const int64_t base = r * k + g * group_size; + float amax = kEps; + for (int64_t i = 0; i < group_size; ++i) + amax = std::fmax(amax, std::fabs(LoadF32(x, base + i))); + const float y_s = amax / kFp8MaxV; + sp[r * groups + g] = y_s; + for (int64_t i = 0; i < group_size; ++i) + op[base + i] = + F32ToFp8(std::fmin(std::fmax(LoadF32(x, base + i) / y_s, kFp8MinV), kFp8MaxV)); + } + } + }); +} + // 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). @@ -3188,6 +3252,8 @@ struct Registrar { reinterpret_cast(static_cast(&RmsNormQuantFp8Kernel))); RegisterOp(OpId::kQuantFp8Static, DeviceType::kCPU, reinterpret_cast(static_cast(&QuantFp8StaticKernel))); + RegisterOp(OpId::kQuantFp8Group, DeviceType::kCPU, + reinterpret_cast(static_cast(&QuantFp8GroupKernel))); RegisterOp(OpId::kMatmulFp8Cutlass, DeviceType::kCPU, reinterpret_cast(static_cast(&MatmulFp8CutlassKernel))); RegisterOp(OpId::kSiluAndMul, DeviceType::kCPU, diff --git a/src/vt/cuda/cuda_quant_fp8.cu b/src/vt/cuda/cuda_quant_fp8.cu index 9de31b81f..98243a4f1 100644 --- a/src/vt/cuda/cuda_quant_fp8.cu +++ b/src/vt/cuda/cuda_quant_fp8.cu @@ -110,6 +110,88 @@ void QuantFp8StaticKernelCuda(Queue& q, Tensor& out_fp8, const Tensor& x, float Check(cudaGetLastError(), "quant_fp8_static launch"); } +// ---- Dynamic per-token, per-GROUP fp8 activation quant ----------------------- +// VT-QUANT-FP8-GROUP (#1189 M1, .agents/specs/vt-quant-fp8-group.md). +// +// Mirror of the kernel that ACTUALLY EXECUTES upstream on this architecture. +// `per_token_group_quant_fp8` calls `torch.ops._C.per_token_group_fp8_quant` +// and returns whenever the platform is CUDA-alike and the input is contiguous +// (vllm/model_executor/layers/quantization/utils/fp8_utils.py:635-650), so its +// Triton kernel never runs here. The executing kernel is +// csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu: +// :47 float local_absmax = eps eps SEEDS the reduction +// :53 fmaxf(local_absmax, fabsf((float)src)) +// :68 float y_s = local_absmax / max_8bit a DIVIDE +// :85 fminf(fmaxf((float)src / y_s, min_8bit), max_8bit) a DIVIDE +// :86 DST_DTYPE(q) hardware e4m3 RNE +// +// THE TWO DIVIDES ARE THE CONTRACT. QuantFp8StaticKernel above multiplies by a +// hoisted reciprocal because that is upstream's shipped form for the static +// per-tensor path; here upstream divides, and the scale changes per group, so +// nothing is loop-invariant. The Triton fallback's `_absmax * (1.0 / fp8_max)` +// (fp8_utils.py:145) differs by up to one f32 ulp and says so in its own +// comment. Near an e4m3 tie that ulp changes the emitted byte. Do not "optimise" +// either divide into a reciprocal: tests/vt/test_ops_quant_fp8_group_cpu.cpp G6 +// compares this kernel with the CPU arm BYTE FOR BYTE, and G1 compares the CPU +// arm with a reference that spells the divide out. +// +// STRUCTURE. Upstream splits one group across 16 lanes and reduces with +// `__shfl_xor_sync` (:21-40, :106). This kernel gives ONE THREAD the whole +// group instead. The two agree bit for bit because `fmaxf` is exact and +// associative over finite inputs, so the reduction ORDER cannot change the +// result, unlike a floating-point sum. That keeps the arm a correctness mirror +// with no shared memory and no intra-group synchronisation; a lane-parallel +// rewrite is a performance question for #1189 M5, which needs a GPU to measure. +// +// This TU is in the UNCONDITIONAL CUDA source list and this kernel, like its +// neighbour, has no cutlass dependency of any kind — it is a max, two divides +// and a hardware convert. `scripts/check-cuda-op-arch-gate.py` pins that, for +// the reason #960 records at the top of this file. +template +__global__ void QuantFp8GroupKernel(uint8_t* out, float* scale, const Tin* x, int group_size, + int64_t num_groups) { + constexpr float kEps = 1e-10f; // fp8_utils.py:570 + constexpr float kFp8MaxV = 448.0f; // quant_utils.py:27-35 finfo(e4m3fn).max + constexpr float kFp8MinV = -448.0f; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t gid = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + gid < num_groups; gid += step) { + const int64_t base = gid * group_size; + float amax = kEps; // :47 — eps SEEDS the reduction, so an all-zero group + // yields 1e-10/448 rather than dividing by zero + for (int i = 0; i < group_size; ++i) amax = fmaxf(amax, fabsf(LoadIn(x, base + i))); // :53 + const float y_s = amax / kFp8MaxV; // :68 + scale[gid] = y_s; + for (int i = 0; i < group_size; ++i) + out[base + i] = + F32ToFp8Dev(fminf(fmaxf(LoadIn(x, base + i) / y_s, kFp8MinV), kFp8MaxV)); // :85 + } +} + +void QuantFp8GroupKernelCuda(Queue& q, Tensor& out_fp8, Tensor& out_scale, const Tensor& x, + int group_size) { + const int64_t m = x.shape[0], k = x.shape[1]; + const int64_t groups_per_row = k / group_size; + const int64_t num_groups = m * groups_per_row; + if (num_groups == 0) return; + cudaStream_t s = AsStream(q); + const int blocks = static_cast(std::min((num_groups + 255) / 256, 65535)); + switch (x.dtype) { + case DType::kF32: + QuantFp8GroupKernel<<>>(out_fp8.Ptr(), + out_scale.Ptr(), x.Ptr(), + group_size, num_groups); + break; + case DType::kBF16: + QuantFp8GroupKernel<__nv_bfloat16><<>>( + out_fp8.Ptr(), out_scale.Ptr(), x.Ptr<__nv_bfloat16>(), group_size, + num_groups); + break; + default: VT_CHECK(false, "cuda quant_fp8_group: unsupported x dtype (f32/bf16 only)"); + } + Check(cudaGetLastError(), "quant_fp8_group launch"); +} + // Table fill only, no CUDA calls (see cuda_ops.cu for the rationale). This // registration must stay at preprocessor-conditional depth 0 in a TU that is // unconditionally compiled for CUDA — that IS the fix for #960. @@ -117,6 +199,8 @@ struct Registrar { Registrar() { RegisterOp(OpId::kQuantFp8Static, DeviceType::kCUDA, reinterpret_cast(static_cast(&QuantFp8StaticKernelCuda))); + RegisterOp(OpId::kQuantFp8Group, DeviceType::kCUDA, + reinterpret_cast(static_cast(&QuantFp8GroupKernelCuda))); } }; Registrar g_registrar; diff --git a/src/vt/op_provider.cpp b/src/vt/op_provider.cpp index 7cb091fcf..bd91bd174 100644 --- a/src/vt/op_provider.cpp +++ b/src/vt/op_provider.cpp @@ -489,6 +489,8 @@ const char* OpNameImpl(OpId op) { return "ConvTranspose1d"; case OpId::kAttentionRelPos: return "AttentionRelPos"; + case OpId::kQuantFp8Group: + return "QuantFp8Group"; case OpId::kCount: break; } diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index b65312462..8a8f6aaa2 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -642,6 +642,37 @@ void QuantFp8Static(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scal reinterpret_cast(GetOp(OpId::kQuantFp8Static, q.device.type))(q, out_fp8, x, input_scale); } +void QuantFp8Group(Queue& q, Tensor& out_fp8, Tensor& out_scale, const Tensor& x, + int group_size) { + VT_CHECK(x.rank == 2 && out_fp8.rank == 2 && out_scale.rank == 2, + "quant_fp8_group: x/out_fp8/out_scale must be rank-2"); + // group_size is validated BEFORE it divides anything: `K % 0` is undefined + // behaviour, so a zero here must refuse rather than trap. + VT_CHECK(group_size > 0, "quant_fp8_group: group_size must be positive"); + const int64_t m = x.shape[0], k = x.shape[1]; + // Mirrors upstream's assert text at + // vllm/model_executor/layers/quantization/utils/fp8_utils.py:596-599. + VT_CHECK(k % group_size == 0, + "quant_fp8_group: the last dimension of x must be divisible by group_size"); + VT_CHECK(out_fp8.shape[0] == m && out_fp8.shape[1] == k, + "quant_fp8_group: out_fp8 must match x shape [M,K]"); + VT_CHECK(out_scale.shape[0] == m && out_scale.shape[1] == k / group_size, + "quant_fp8_group: out_scale must be [M, K/group_size]"); + VT_CHECK(IsFloat(x.dtype), "quant_fp8_group: float x (f32/bf16) required"); + VT_CHECK(out_fp8.dtype == DType::kI8, + "quant_fp8_group: out_fp8 must be i8 (raw fp8-e4m3fn bytes)"); + // f32, not the model dtype: upstream allocates the scale f32 (fp8_utils.py:631) + // and the block-scaled GEMM multiplies it into an f32 accumulator. + VT_CHECK(out_scale.dtype == DType::kF32, "quant_fp8_group: out_scale must be f32"); + // Upstream asserts `x.stride(-1) == 1` (fp8_utils.py:600); a group that is not + // contiguous would read across rows. + VT_CHECK(x.IsContiguous() && out_fp8.IsContiguous() && out_scale.IsContiguous(), + "quant_fp8_group: contiguous tensors required"); + VT_CHECK(x.device == q.device && out_fp8.device == q.device && out_scale.device == q.device, + "quant_fp8_group: device mismatch (x/out_fp8/out_scale/queue)"); + reinterpret_cast(GetOp(OpId::kQuantFp8Group, q.device.type))( + q, out_fp8, out_scale, x, group_size); +} void RmsNormQuantFp8(Queue& q, Tensor& out_fp8, Tensor* out_bf16, const Tensor& x, const Tensor& weight, const RmsNormArgs& args, Tensor* residual, float input_scale) { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 82c362fbe..bad02c87d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1713,6 +1713,10 @@ vllm_cpp_add_test(test_ops_fp8_cutlass vt/test_ops_fp8_cutlass.cpp) # 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) +# VT-QUANT-FP8-GROUP (#1189 M1): the dynamic per-token per-group fp8 activation +# quant. CPU-gateable by construction; its CPU-vs-CUDA byte-identity arm (G6) is +# CUDA-gated and reports PENDING rather than skipping where no device exists. +vllm_cpp_add_test(test_ops_quant_fp8_group_cpu vt/test_ops_quant_fp8_group_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/scripts/test_check_cuda_op_arch_gate.py b/tests/scripts/test_check_cuda_op_arch_gate.py index 3b0d78791..a425c1150 100644 --- a/tests/scripts/test_check_cuda_op_arch_gate.py +++ b/tests/scripts/test_check_cuda_op_arch_gate.py @@ -51,6 +51,16 @@ " reinterpret_cast(" "static_cast(&QuantFp8StaticKernelCuda)));\n" ) +# The miniature describes the REAL tree, and the real TU now registers a second +# op from the same Registrar (#1189 M1, kQuantFp8Group). Without this line the +# baseline miniature is red for a reason that has nothing to do with the +# mutation each case applies, which would hide what those cases measure. Every +# mutation below still targets REGISTRATION, so nothing they assert is widened. +GROUP_REGISTRATION = ( + " RegisterOp(OpId::kQuantFp8Group, DeviceType::kCUDA,\n" + " reinterpret_cast(" + "static_cast(&QuantFp8GroupKernelCuda)));\n" +) class FakeTree: @@ -94,9 +104,10 @@ def __exit__(self, *exc: object) -> None: namespace vt::cuda {{ namespace {{ void QuantFp8StaticKernelCuda(Queue&, Tensor&, const Tensor&, float) {{}} +void QuantFp8GroupKernelCuda(Queue&, Tensor&, Tensor&, const Tensor&, int) {{}} struct Registrar {{ Registrar() {{ -{REGISTRATION} }} +{REGISTRATION}{GROUP_REGISTRATION} }} }}; Registrar g_registrar; }} diff --git a/tests/vt/test_ops_quant_fp8_group_cpu.cpp b/tests/vt/test_ops_quant_fp8_group_cpu.cpp new file mode 100644 index 000000000..c8d478ce5 --- /dev/null +++ b/tests/vt/test_ops_quant_fp8_group_cpu.cpp @@ -0,0 +1,640 @@ +// vllm.cpp — dynamic per-token, per-group FP8 (e4m3fn) activation quant. +// +// VT-QUANT-FP8-GROUP (.agents/specs/vt-quant-fp8-group.md), issue #1189 +// milestone M1. Pinned oracle: vLLM 5559679229bc961848b121ccdeaa8fa5d79bec98. +// +// WHICH UPSTREAM ARM THIS MIRRORS, because there are two and they disagree. +// `per_token_group_quant_fp8` reads like a Triton kernel with a C++ fast path. +// It is the other way round: on a CUDA-alike platform with a contiguous input +// it calls the C++ custom op and RETURNS +// (vllm/model_executor/layers/quantization/utils/fp8_utils.py:635-650), so the +// Triton kernel below it never executes there. The executing kernel is +// csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu +// :47 float local_absmax = eps (eps SEEDS the reduction) +// :53 fmaxf(local_absmax, fabsf((float)src)) +// :68 float y_s = local_absmax / max_8bit <- a DIVIDE +// :85 fminf(fmaxf((float)src / y_s, min_8bit), max_8bit) <- a DIVIDE +// :86 DST_DTYPE(q) (hardware e4m3 RNE, saturating) +// The Triton fallback instead forms `scale_raw = _absmax * (1.0 / fp8_max)` +// (fp8_utils.py:145) under a comment that names the 1-ULP difference. One f32 +// ULP before an e4m3 round changes the emitted byte near a tie. +// +// WHY G1 EXISTS ON TOP OF THE PORTED CASE, and this is MEASURED rather than +// argued. Upstream's own case compares values at rtol=0.15 +// (tests/kernels/quantization/test_block_fp8.py:112-114) and the scale at +// torch.allclose's default rtol=1e-5 (:115). Two mutations of the CPU kernel to +// the Triton arm's form, each built (compile_rc=0) and run: +// * `y_s = amax * (1.0f/448)` instead of `amax / 448` +// G1 fails 49 of 146 assertions. G2 fails 6 of its 48 shape checks, ALL of +// them the VALUE check and ALL at num_tokens=2050; the scale check never +// fires, and no shape at num_tokens=7 fires at all. +// * `x * (1.0f/y_s)` instead of `x / y_s` +// G1 fails 14 of 146. G2 passes ENTIRELY, 50 of 50. +// So upstream's tolerances see one of the two forms, on the large shapes only, +// by luck of which elements land on an e4m3 boundary. They do not see the other +// at all. G1 sees both, on every shape, because it compares BYTES. +// +// G1 BITWISE, ZERO TOLERANCE, against an INDEPENDENTLY WRITTEN reference. +// G2 the ported upstream case (test_block_fp8.py:82-118) with its grid. +// G3 the scale itself: amax/448 by construction, and the all-zero group. +// G4 the shape and dtype contract of out_scale. +// G5 the refusals, each by name. +// G6 CPU vs CUDA byte identity. CUDA-gated; OWED on a host with no GPU. +// +// THE G1 REFERENCE IS DERIVED FROM THE FORMAT, never from a codec in src/, +// 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). G2 needs the same encode 290M times, where an +// exhaustive scan is not affordable, so it uses `FastEncodeRne`, a binary +// search over the same table -- and G1 proves the two agree on every input +// class before G2 is allowed to rely on the fast one. +// +// Byte comparisons, not doctest Approx: Approx carries a `scale` term +// defaulting to 1.0 and therefore a ~1.19e-5 absolute floor, meaningless here. +#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 ("fn"). +constexpr float kFp8Max = 448.0f; // = 1.75 * 2^8, encoding 0x7E +constexpr float kFp8Min = -448.0f; // quant_utils.py:27-35 finfo(e4m3fn) +constexpr float kEps = 1e-10f; // fp8_utils.py:570 default, the only value + // any upstream call site passes + +// 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 carries across an exponent step too, 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 same encode in O(log n), for the 290M-element ported grid where the +// exhaustive scan is not affordable. The 127 finite magnitudes 0x00..0x7E are +// MONOTONIC in the byte value, so a binary search over their midpoints picks the +// nearest, and an exact midpoint hit resolves to the even mantissa. G1 proves +// this agrees with RefEncodeRne on every input class before G2 uses it. +const std::vector& Magnitudes() { + static const std::vector table = [] { + std::vector v; + v.reserve(127); + for (unsigned e = 0; e <= 15; ++e) + for (unsigned m = 0; m <= 7; ++m) { + if (e == 15 && m == 7) continue; + v.push_back(E4m3Exact(e, m)); + } + return v; + }(); + return table; +} + +uint8_t FastEncodeRne(float r) { + const auto sign = static_cast(std::signbit(r) ? 0x80u : 0x00u); + const double a = std::fabs(static_cast(r)); + const std::vector& mag = Magnitudes(); + // First index whose magnitude is >= a. + const auto it = std::lower_bound(mag.begin(), mag.end(), a); + size_t hi = static_cast(it - mag.begin()); + if (hi == 0) return sign; // a == 0 + if (hi >= mag.size()) return static_cast(sign | 0x7Eu); // a == 448 + const size_t lo = hi - 1; + const double dlo = a - mag[lo], dhi = mag[hi] - a; + size_t pick = 0; + if (dlo < dhi) { + pick = lo; + } else if (dhi < dlo) { + pick = hi; + } else { + // Table index == byte value, so `lo & 1` IS the mantissa's low bit. + pick = (lo & 1u) == 0u ? lo : hi; // even significand wins the tie + } + return static_cast(sign | static_cast(pick)); +} + +// Decode a raw e4m3fn byte back to an exact double, for the value comparison +// upstream's ported case makes. Table-driven because G2 decodes ~145M bytes +// twice each and an ldexp per call dominated the case's run time. +const std::vector& DequantTable() { + static const std::vector table = [] { + std::vector v(256); + for (unsigned b = 0; b < 256; ++b) { + const double m = E4m3Exact((b >> 3) & 0xFu, b & 0x7u); + v[b] = (b & 0x80u) != 0 ? -m : m; + } + return v; + }(); + return table; +} + +double RefDequant(uint8_t byte) { return DequantTable()[byte]; } + +// bf16 round-trip, so a bf16 input reaches the reference at the SAME width the +// kernel loads it at. Written here rather than taken from vt/, for the same +// independence reason as RefEncodeRne: round-to-nearest-even on the low 16 bits. +float RoundToBf16(float v) { return vt::BF16ToF32(vt::F32ToBF16(v)); } + +// The whole upstream expression, in upstream's order and upstream's f32 width. +// `enc` selects the exhaustive or the binary-search encoder; the arithmetic +// above it is identical either way. +void RefQuantGroup(const float* x, int64_t n, uint8_t* out, float* out_scale, + uint8_t (*enc)(float)) { + float amax = kEps; // :47 eps SEEDS it + for (int64_t i = 0; i < n; ++i) amax = std::fmax(amax, std::fabs(x[i])); // :53 + const float y_s = amax / kFp8Max; // :68 a DIVIDE + *out_scale = y_s; + for (int64_t i = 0; i < n; ++i) { + const float q = std::fmin(std::fmax(x[i] / y_s, kFp8Min), kFp8Max); // :85 + out[i] = enc(q); // :86 + } +} + +// --- inputs ---------------------------------------------------------------- +// One ROW of G1 input. Every value class whose handling a mutation can break, +// sized so that the group's amax is known and the scaled values land on +// saturation, exact e4m3 ties, the subnormal ladder and both zeros. +std::vector G1Row(int64_t k, uint32_t seed, float span) { + std::vector v; + v.reserve(static_cast(k)); + std::mt19937 rng(seed); + std::uniform_real_distribution ux(-span, span); + // The scaled value x/y_s spans [-448, 448] by construction, so these are the + // ties and ladder points of the e4m3 grid, expressed in x units. + const float unit = span / kFp8Max; + 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); + v.push_back(static_cast((lo + hi) / 2.0) * unit); + v.push_back(-static_cast((lo + hi) / 2.0) * unit); + } + for (int m = 0; m <= 8; ++m) { + v.push_back(static_cast(m / 512.0) * unit); + v.push_back(static_cast((2 * m + 1) / 1024.0) * unit); + v.push_back(-static_cast((2 * m + 1) / 1024.0) * unit); + } + v.push_back(0.0f); + v.push_back(-0.0f); + v.push_back(span); // the amax itself: this element must quantize to 0x7E + v.push_back(-span); + while (static_cast(v.size()) < k) v.push_back(ux(rng)); + v.resize(static_cast(k)); + return v; +} + +// Runs the CPU op and returns the number of differing OUTPUT bytes plus the +// number of differing scale words. `x_dtype` selects the input width; the +// reference consumes the same width, so this compares the codec and not the +// store width. +struct ByteDiff { + size_t bytes = 0; + size_t scales = 0; + size_t nonzero_out = 0; +}; + +ByteDiff RunBitwise(const std::vector& x_row, int64_t m, int64_t k, int group_size, + DType x_dtype, uint8_t (*enc)(float)) { + REQUIRE(static_cast(x_row.size()) == k); + const int64_t groups = k / group_size; + Queue q{Cpu(), nullptr}; + + std::vector xf32(static_cast(m * k)); + std::vector xbf16(static_cast(m * k)); + std::vector ref_in(static_cast(m * k)); + for (int64_t r = 0; r < m; ++r) + for (int64_t c = 0; c < k; ++c) { + // Each row is the same population rotated, so every row has its own amax. + const float v = x_row[static_cast((c + r * 37) % k)] * + (1.0f + 0.25f * static_cast(r)); + xf32[static_cast(r * k + c)] = v; + xbf16[static_cast(r * k + c)] = vt::F32ToBF16(v); + ref_in[static_cast(r * k + c)] = x_dtype == DType::kBF16 ? RoundToBf16(v) : v; + } + + void* xp = x_dtype == DType::kBF16 ? static_cast(xbf16.data()) + : static_cast(xf32.data()); + Tensor tx = MakeTensor(xp, x_dtype, Cpu(), {m, k}); + std::vector got(static_cast(m * k), 0xFFu); + std::vector got_s(static_cast(m * groups), -1.0f); + Tensor tq = MakeTensor(got.data(), DType::kI8, Cpu(), {m, k}); + Tensor ts = MakeTensor(got_s.data(), DType::kF32, Cpu(), {m, groups}); + vt::QuantFp8Group(q, tq, ts, tx, group_size); + + std::vector want(static_cast(m * k)); + std::vector want_s(static_cast(m * groups)); + for (int64_t r = 0; r < m; ++r) + for (int64_t g = 0; g < groups; ++g) + RefQuantGroup(&ref_in[static_cast(r * k + g * group_size)], group_size, + &want[static_cast(r * k + g * group_size)], + &want_s[static_cast(r * groups + g)], enc); + + ByteDiff d; + for (size_t i = 0; i < want.size(); ++i) { + if (got[i] != want[i]) { + if (d.bytes == 0) { + CAPTURE(i); + CAPTURE(static_cast(got[i])); + CAPTURE(static_cast(want[i])); + } + ++d.bytes; + } + if ((want[i] & 0x7Fu) != 0u) ++d.nonzero_out; + } + for (size_t i = 0; i < want_s.size(); ++i) { + // BITWISE on the scale too: an f32 compare here, not an Approx, because a + // 1-ULP scale difference is exactly what distinguishes the two upstream + // arms and an Approx cannot see it. + if (got_s[i] != want_s[i]) { + if (d.scales == 0) { + CAPTURE(i); + CAPTURE(got_s[i]); + CAPTURE(want_s[i]); + } + ++d.scales; + } + } + return d; +} + +} // namespace + +// =========================================================================== +// G1 — bitwise, zero tolerance, against the exhaustive reference. This is the +// only arm that can distinguish the executing CUDA arm's two divides from the +// Triton fallback's reciprocal multiply. +TEST_CASE("G1: CPU QuantFp8Group equals an independent e4m3 reference byte for byte") { + REQUIRE(vt::OpRegistered(vt::OpId::kQuantFp8Group, DeviceType::kCPU)); + // The fast encoder is validated against the exhaustive one here, so G2 may + // rely on it. A disagreement is a defect in the TEST and must not be widened. + size_t enc_mismatch = 0; + for (int e = -12; e <= 9; ++e) + for (int m = 0; m < 16; ++m) { + const auto v = static_cast(std::ldexp(1.0 + m / 16.0, e)); + for (float s : {1.0f, -1.0f}) { + const float r = std::fmin(std::fmax(v * s, kFp8Min), kFp8Max); + if (RefEncodeRne(r) != FastEncodeRne(r)) ++enc_mismatch; + } + } + CHECK(enc_mismatch == 0u); + + // group_size 64/128/512 over K that each divide, both input widths, and M > 1 + // so a per-ROW amax that leaked across rows is visible. + for (int group_size : {64, 128, 512}) { + for (int64_t k : {512, 1024}) { + for (float span : {1.0f, 0.0037f, 91.5f}) { + CAPTURE(group_size); + CAPTURE(k); + CAPTURE(span); + const auto row = G1Row(k, 1234u + static_cast(k), span); + for (DType dt : {DType::kF32, DType::kBF16}) { + const ByteDiff d = RunBitwise(row, 3, k, group_size, dt, &RefEncodeRne); + CHECK(d.bytes == 0u); + CHECK(d.scales == 0u); + // VACUITY GUARD: an all-zero reference would make any implementation + // pass. Most of this population is nonzero by construction. + CHECK(d.nonzero_out > static_cast(3 * k) / 2); + } + } + } + } +} + +// =========================================================================== +// G2 — the ported upstream case. +// +// PORT OF tests/kernels/quantization/test_block_fp8.py:82-118 at vLLM +// 5559679229bc961848b121ccdeaa8fa5d79bec98, with its grid at :42-46: +// NUM_TOKENS = [7, 2050] D = [512, 4096, 5120, 13824] +// GROUP_SIZE = [64, 128, 512] SEEDS = [0] +// and its tolerances: values compared after dequant at rtol=0.15 (:112-114), +// the scale at torch.allclose's defaults rtol=1e-5 atol=1e-8 (:115). +// +// HARNESS ADAPTATIONS, and nothing else changed. (1) Upstream draws `x` from +// torch.rand on the device; there is no bit-compatible host RNG, so this draws +// the same shape from std::mt19937 over the same [0,1) support. (2) Upstream's +// DTYPES list ships only bfloat16, with float32 commented out at :40, and this +// runs bfloat16 only for the same reason. f32 input coverage is not lost: G1 +// compares BOTH widths bitwise, which is the stronger statement. +// (3) COLUMN_MAJOR_SCALES and TMA_ALIGNED_SCALES (:45-46, :117-120) are +// dropped: those layouts are owed to #1189 M5 and this op emits only the +// row-major one. (4) The ROCm 1-ULP branch (:97-110) is not applicable. +TEST_CASE("G2: ported test_per_token_group_quant_fp8 grid matches the native reference") { + REQUIRE(vt::OpRegistered(vt::OpId::kQuantFp8Group, DeviceType::kCPU)); + Queue q{Cpu(), nullptr}; + size_t total_nonzero = 0; + size_t total_elems = 0; + + for (int64_t num_tokens : {7, 2050}) { + for (int64_t d : {512, 4096, 5120, 13824}) { + for (int group_size : {64, 128, 512}) { + for (DType dt : {DType::kBF16}) { // DTYPES = [torch.bfloat16] at :40 + CAPTURE(num_tokens); + CAPTURE(d); + CAPTURE(group_size); + const int64_t groups = d / group_size; + const auto n = static_cast(num_tokens * d); + + std::mt19937 rng(0u); // SEEDS = [0] + std::uniform_real_distribution ux(0.0f, 1.0f); // torch.rand + std::vector xf32(n); + std::vector xbf16(n); + std::vector ref_in(n); + for (size_t i = 0; i < n; ++i) { + const float v = ux(rng); + xf32[i] = v; + xbf16[i] = vt::F32ToBF16(v); + ref_in[i] = dt == DType::kBF16 ? RoundToBf16(v) : v; + } + void* xp = dt == DType::kBF16 ? static_cast(xbf16.data()) + : static_cast(xf32.data()); + Tensor tx = MakeTensor(xp, dt, Cpu(), {num_tokens, d}); + std::vector got(n, 0xFFu); + std::vector got_s(static_cast(num_tokens * groups), -1.0f); + Tensor tq = MakeTensor(got.data(), DType::kI8, Cpu(), {num_tokens, d}); + Tensor ts = MakeTensor(got_s.data(), DType::kF32, Cpu(), {num_tokens, groups}); + vt::QuantFp8Group(q, tq, ts, tx, group_size); + + std::vector want(n); + std::vector want_s(static_cast(num_tokens * groups)); + for (int64_t r = 0; r < num_tokens; ++r) + for (int64_t g = 0; g < groups; ++g) + RefQuantGroup(&ref_in[static_cast(r * d + g * group_size)], group_size, + &want[static_cast(r * d + g * group_size)], + &want_s[static_cast(r * groups + g)], &FastEncodeRne); + + // :112-114 assert allclose(out.float(), ref_out.float(), rtol=0.15) + size_t bad = 0, nonzero = 0; + for (size_t i = 0; i < n; ++i) { + const double a = RefDequant(got[i]), b = RefDequant(want[i]); + if (!(std::fabs(a - b) <= 0.15 * std::fabs(b))) ++bad; + if (b != 0.0) ++nonzero; + } + CHECK(bad == 0u); + // :115 assert allclose(scale, ref_scale) rtol=1e-5 atol=1e-8 + size_t bad_s = 0; + for (size_t i = 0; i < want_s.size(); ++i) + if (!(std::fabs(static_cast(got_s[i]) - want_s[i]) <= + 1e-8 + 1e-5 * std::fabs(static_cast(want_s[i])))) + ++bad_s; + CHECK(bad_s == 0u); + total_nonzero += nonzero; + total_elems += n; + } + } + } + } + // VACUITY GUARD: torch.rand draws from [0,1), so essentially every reference + // element is nonzero. An all-zero reference would make any implementation + // pass every comparison above. + CAPTURE(total_nonzero); + CAPTURE(total_elems); + CHECK(total_nonzero > total_elems - total_elems / 1000); +} + +// =========================================================================== +// G3 — the scale itself. A token gate cannot see a scale that collapsed to +// per-tensor or that lost its eps floor; an exact value can. +TEST_CASE("G3: the group scale is amax/448 and an all-zero group uses the eps floor") { + Queue q{Cpu(), nullptr}; + constexpr int64_t kM = 2, kK = 256; + constexpr int kG = 64; + const int64_t groups = kK / kG; + std::vector x(static_cast(kM * kK), 0.0f); + // Row 0: each group gets its OWN known amax. A per-tensor collapse makes all + // four scales equal, which this detects. + const float amax[4] = {1.0f, 8.0f, 0.125f, 300.0f}; + for (int64_t g = 0; g < groups; ++g) { + for (int64_t i = 0; i < kG; ++i) + x[static_cast(g * kG + i)] = amax[g] * 0.5f * ((i % 3 == 0) ? -1.0f : 1.0f); + x[static_cast(g * kG + 7)] = -amax[g]; // the amax, negative + } + // Row 1 stays all zero: y_s must be eps/448 and every output byte must be 0. + Tensor tx = MakeTensor(x.data(), DType::kF32, Cpu(), {kM, kK}); + std::vector got(static_cast(kM * kK), 0xFFu); + std::vector got_s(static_cast(kM * groups), -1.0f); + Tensor tq = MakeTensor(got.data(), DType::kI8, Cpu(), {kM, kK}); + Tensor ts = MakeTensor(got_s.data(), DType::kF32, Cpu(), {kM, groups}); + vt::QuantFp8Group(q, tq, ts, tx, kG); + + for (int64_t g = 0; g < groups; ++g) { + CAPTURE(g); + CHECK(got_s[static_cast(g)] == amax[g] / kFp8Max); + // The amax element saturates the grid: |x|/y_s == 448 exactly -> 0x7E. + CHECK(got[static_cast(g * kG + 7)] == 0xFEu); // sign bit set + } + // A per-tensor collapse would make these equal. + CHECK(got_s[0] != got_s[1]); + for (int64_t g = 0; g < groups; ++g) { + CAPTURE(g); + CHECK(got_s[static_cast(groups + g)] == kEps / kFp8Max); + for (int64_t i = 0; i < kG; ++i) + CHECK(got[static_cast(kK + g * kG + i)] == 0x00u); + } +} + +// =========================================================================== +// G4 — the scale's dtype and shape are part of the contract. Upstream allocates +// float32 [M, K/group_size] (fp8_utils.py:629-631). A scale that silently +// narrowed to bf16 would still produce plausible tokens. +TEST_CASE("G4: out_scale must be f32 and shaped M by K over group_size") { + Queue q{Cpu(), nullptr}; + constexpr int64_t kM = 4, kK = 256; + constexpr int kG = 128; + std::vector x(static_cast(kM * kK), 0.5f); + Tensor tx = MakeTensor(x.data(), DType::kF32, Cpu(), {kM, kK}); + std::vector outq(static_cast(kM * kK)); + Tensor tq = MakeTensor(outq.data(), DType::kI8, Cpu(), {kM, kK}); + std::vector s(static_cast(kM * (kK / kG))); + + Tensor ok = MakeTensor(s.data(), DType::kF32, Cpu(), {kM, kK / kG}); + CHECK_NOTHROW(vt::QuantFp8Group(q, tq, ok, tx, kG)); + + Tensor narrow = MakeTensor(s.data(), DType::kBF16, Cpu(), {kM, kK / kG}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq, narrow, tx, kG), std::runtime_error); + + Tensor wrong_groups = MakeTensor(s.data(), DType::kF32, Cpu(), {kM, kK / kG + 1}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq, wrong_groups, tx, kG), std::runtime_error); + + Tensor wrong_rows = MakeTensor(s.data(), DType::kF32, Cpu(), {kM - 1, kK / kG}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq, wrong_rows, tx, kG), std::runtime_error); + + Tensor wrong_out = MakeTensor(outq.data(), DType::kF32, Cpu(), {kM, kK}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, wrong_out, ok, tx, kG), std::runtime_error); +} + +// =========================================================================== +// G5 — the refusals. Upstream asserts divisibility at fp8_utils.py:596-599 and +// contiguity at :600. A ragged K accepted silently reads past the row. +TEST_CASE("G5: a K that the group size does not divide is refused by name") { + Queue q{Cpu(), nullptr}; + constexpr int64_t kM = 2, kK = 200; // 200 % 128 != 0 and 200 % 64 != 0 + std::vector x(static_cast(kM * kK), 0.5f); + std::vector outq(static_cast(kM * kK)); + std::vector s(static_cast(kM * 2)); + Tensor tx = MakeTensor(x.data(), DType::kF32, Cpu(), {kM, kK}); + Tensor tq = MakeTensor(outq.data(), DType::kI8, Cpu(), {kM, kK}); + Tensor ts = MakeTensor(s.data(), DType::kF32, Cpu(), {kM, 2}); + CHECK_THROWS_WITH_AS(vt::QuantFp8Group(q, tq, ts, tx, 128), + doctest::Contains("must be divisible"), std::runtime_error); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq, ts, tx, 0), std::runtime_error); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq, ts, tx, -128), std::runtime_error); + + // A non-contiguous x: fp8_utils.py:600 `x.stride(-1) == 1`. + constexpr int64_t kK2 = 256; + std::vector x2(static_cast(kM * kK2), 0.5f); + Tensor gappy = MakeTensor(x2.data(), DType::kF32, Cpu(), {kM, kK2}); + gappy.stride[1] = 2; + std::vector outq2(static_cast(kM * kK2)); + std::vector s2(static_cast(kM * 2)); + Tensor tq2 = MakeTensor(outq2.data(), DType::kI8, Cpu(), {kM, kK2}); + Tensor ts2 = MakeTensor(s2.data(), DType::kF32, Cpu(), {kM, 2}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq2, ts2, gappy, 128), std::runtime_error); + + // A device mismatch between the queue and the tensors. + Tensor onGpu = MakeTensor(x2.data(), DType::kF32, Device{DeviceType::kCUDA, 0}, {kM, kK2}); + CHECK_THROWS_AS(vt::QuantFp8Group(q, tq2, ts2, onGpu, 128), std::runtime_error); +} + +// =========================================================================== +// G6 — CPU vs CUDA, bitwise, on the identical input. The arm that says the CPU +// path mirrors what ships rather than merely being self-consistent. +// NAME THIS CASE WITHOUT A COMMA: doctest splits `-tc=` on commas, so a comma +// makes the name unselectable and the binary reports `0 cases ran ... SUCCESS!` +// with exit 0 (measured on tests/vt/test_ops_fp8_cpu.cpp, #468 review F6). +TEST_CASE("G6: CPU QuantFp8Group equals CUDA QuantFp8Group 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. This is the state + // .agents/specs/vt-quant-fp8-group.md records under `## Owed`: the row took + // no GPU lease by design, so the CUDA arm compiles and does not run. + MESSAGE("G6 PENDING: no CUDA device on this host, CPU-vs-CUDA byte agreement " + "was NOT measured (owed by #1189 milestone M5)"); + CHECK(vt::OpRegistered(vt::OpId::kQuantFp8Group, 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 (int group_size : {64, 128, 512}) { + for (int64_t k : {512, 1024}) { + CAPTURE(group_size); + CAPTURE(k); + constexpr int64_t kM = 3; + const int64_t groups = k / group_size; + const auto row = G1Row(k, 4321u, 1.0f); + std::vector x(static_cast(kM * k)); + for (int64_t r = 0; r < kM; ++r) + for (int64_t c = 0; c < k; ++c) + x[static_cast(r * k + c)] = + row[static_cast((c + r * 37) % k)] * (1.0f + 0.25f * static_cast(r)); + + std::vector cpu_out(static_cast(kM * k)); + std::vector cpu_s(static_cast(kM * groups)); + Tensor tx_cpu = MakeTensor(x.data(), DType::kF32, Cpu(), {kM, k}); + Tensor tq_cpu = MakeTensor(cpu_out.data(), DType::kI8, Cpu(), {kM, k}); + Tensor ts_cpu = MakeTensor(cpu_s.data(), DType::kF32, Cpu(), {kM, groups}); + vt::QuantFp8Group(cq, tq_cpu, ts_cpu, tx_cpu, group_size); + + void* dx = b.Alloc(static_cast(kM * k) * sizeof(float)); + void* dq = b.Alloc(static_cast(kM * k)); + void* ds = b.Alloc(static_cast(kM * groups) * sizeof(float)); + b.Copy(gq, dx, x.data(), static_cast(kM * k) * sizeof(float)); + Tensor tx_gpu = MakeTensor(dx, DType::kF32, gpu, {kM, k}); + Tensor tq_gpu = MakeTensor(dq, DType::kI8, gpu, {kM, k}); + Tensor ts_gpu = MakeTensor(ds, DType::kF32, gpu, {kM, groups}); + vt::QuantFp8Group(gq, tq_gpu, ts_gpu, tx_gpu, group_size); + std::vector gpu_out(static_cast(kM * k)); + std::vector gpu_s(static_cast(kM * groups)); + b.Copy(gq, gpu_out.data(), dq, gpu_out.size()); + b.Copy(gq, gpu_s.data(), ds, gpu_s.size() * sizeof(float)); + b.Synchronize(gq); + + size_t bad = 0, bad_s = 0, nonzero = 0; + for (size_t i = 0; i < cpu_out.size(); ++i) { + if (cpu_out[i] != gpu_out[i]) ++bad; + if ((cpu_out[i] & 0x7Fu) != 0u) ++nonzero; + } + for (size_t i = 0; i < cpu_s.size(); ++i) + if (cpu_s[i] != gpu_s[i]) ++bad_s; + CHECK(bad == 0u); + CHECK(bad_s == 0u); + CHECK(nonzero > cpu_out.size() / 2); // vacuity guard + b.Free(dx); + b.Free(dq); + b.Free(ds); + } + } + b.DestroyQueue(gq); +}