diff --git a/.agents/issue-index.md b/.agents/issue-index.md index ec4e81558..7dae66c06 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,5 +366,6 @@ 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 | | [#1190](https://github.com/mudler/vllm.cpp/issues/1190) | `ENV-AGNOSTIC-CAMPAIGN` | One operator's hosts, share paths and addresses are written into 227 tracked files where a `.env` placeholder belongs, so a second developer who follows the protocol documents is told to reach a box on one home network. Re-derived at `fd64c76ee` with `git grep -cIE 'dgx\.casa\|nas_share\|192\.168\.\|thor:gpu0'`: `dgx.casa` 203 files, `nas_share` 32, `192.168.` 24, `thor:gpu0` 11. The mechanism was already there and almost unused: `.env.example` is tracked, `.env` and `.agents/developer-preferences.md` are ignored, `scripts/agent-onboard.py --env-set` already records one answered value and refuses an undeclared key, and `${VLLM_SOURCE}` and `${GPU_LOCK}` already resolve in 59 and 50 files while `${DEVICE_ARCH}`, `${DEVICE_TOOLKIT_ROOT}` and `${DEVICE_COMPILER}` resolve in none. This SCOPING row lands the rule and the mechanism, not the sweep. The rule is a substitution test: replace the literal with a second developer's value and if the sentence stays true it is CONFIGURATION and becomes `${KEY}`, and if it becomes false it is PROVENANCE and stays literal. That reverses the issue's own ranking, because all 34 hits in the densest guide file `.agents/environment.md` are provenance or named-profile definition and none is configuration, so density does not predict the defect and a blind `sed` would falsify records. Landed: three new keys `GATE_CHECKOUT`, `SHARED_STORAGE_ROOT` and `GATE_DEVICE` derived from the literals that recur and map to no existing key; the create-on-first-use route in `scripts/agent-start.py`, which printed `environment: missing` as a status label and then listed next actions that never mentioned it, so the fallback in practice was a host name copied from a document; the matching obligation in `AGENTS.md`; and one worked example. The worked example `scripts/dgx-bringup.sh` found a live defect rather than a cosmetic one: it defaulted `CUTLASS_DIR` to `$HOME/cutlass_probe` while `.agents/environment.md:389` records `$HOME/cutlass-4.5.0` as mandatory on the same box, and a configure that misses CUTLASS silently drops the sm120a NVFP4 GEMM and FlashAttention-2, which that file measures as moving the SACRED `test_qwen27_paged_engine` from 235/235 to 234/235 with the source untouched, so a stale hard-coded default is a false green. Its new test also caught that `set -a; . ./.env; set +a` over an `.env.example` copy blanks a value the caller exported, so the process-environment contract is now executable. Waves `ENV-AGNOSTIC-W1-TOOLING` through `ENV-AGNOSTIC-W5-LEDGERS` own the sweep, partition all 227 files, and open their own issues. Spec [`env-agnostic.md`](specs/env-agnostic.md) | bug | | [#1193](https://github.com/mudler/vllm.cpp/issues/1193) | `SPEC-DSPARK-QWEN3-ROUTING` | A Qwen3 DSpark draft declaring `architectures=["DSparkDraftModel"]` with `model_type` `qwen3` has no route. The pin forces every DSpark draft that is not `Qwen3DSparkModel` or `Gemma4DSparkModel` onto `model_type` `deepseek_v4` (`vllm/config/speculative.py:934-944` @ `555967922`), and vLLM PR 52197 (merged 2026-08-17 at `7075ddac`) replaced that with a leading branch normalizing the pair to `Qwen3DSparkModel`. We diverge from BOTH: the forced rewrite was never ported, so nothing in `src/vllm/entrypoints/model_loader.cpp` reads a draft config's `architectures` key at all, and `SpeculativeConfig::IsDsparkDraft` (`include/vllm/config/speculative.h:120-136`) has no production caller — every reference outside its header is in `tests/vllm/config/test_speculative_dspark.cpp:132-140`, and `ResolveSpecConfig` branches on `cli.method` alone. The checkpoint is real and gateable here: `RadixArk/Qwen3.8-27B-DSpark` at revision `85ef153be924f17ce4bf62726954eeaa4a73e854` carries exactly that config shape in one 2718576122-byte shard, drafting five layers for a 64-layer Qwen3.8-27B target | bug | diff --git a/.agents/specs/vt-quant-fp8-group.md b/.agents/specs/vt-quant-fp8-group.md new file mode 100644 index 000000000..4350f8121 --- /dev/null +++ b/.agents/specs/vt-quant-fp8-group.md @@ -0,0 +1,277 @@ +# 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** | +| 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 + +`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` | + +**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 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 + 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 + +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); +}