From 2132d85a5215d53fe14d8748abaaa158000f7347 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 07:14:58 +0000 Subject: [PATCH 1/6] spec(VT-FP8-QUANT-ARCH-GATE): pin the fp8 activation quant to an arch-independent TU `vt::QuantFp8Static`'s only CUDA registration lives in `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:376`, and `CMakeLists.txt:1668` compiles that translation unit only when `VT_CUTLASS_FP8_ARCHS` is non-empty. The kernel body has no cutlass dependency at all -- it is `out[i] = e4m3(x[i] * (1/input_scale))` -- so on every CUDA arch outside the cutlass-fp8 cell the op is not registered for `DeviceType::kCUDA`, the resolver installs the portable CPU reference tier for a CUDA queue, and the first call dereferences device pointers and segfaults. Nothing refuses first: the GEMM partner `kMatmulFp8CublasLt` is registered unconditionally, so the model-layer guard that keys on it passes and the crash arrives one call later under a banner reading "correct but slow". This spec is committed before the implementation, per AGENTS.md "Spec before code". It records the scope, the three TUs considered as a new home and why a dedicated file wins, the upstream partition being restored, and the two instruments the change owes: a runtime pin that observes the registration, and a structural checker that can fail at PR time on a host CI actually runs. Issue #960. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/vt-fp8-quant-arch-gate.md | 229 ++++++++++++++++++++++++ 1 file changed, 229 insertions(+) create mode 100644 .agents/specs/vt-fp8-quant-arch-gate.md diff --git a/.agents/specs/vt-fp8-quant-arch-gate.md b/.agents/specs/vt-fp8-quant-arch-gate.md new file mode 100644 index 000000000..c4a9ce577 --- /dev/null +++ b/.agents/specs/vt-fp8-quant-arch-gate.md @@ -0,0 +1,229 @@ +# VT-FP8-QUANT-ARCH-GATE — `QuantFp8Static` is trapped in the cutlass-fp8 build gate + +| | | +|---|---| +| Issue | [#960](https://github.com/mudler/vllm.cpp/issues/960) (with [#844](https://github.com/mudler/vllm.cpp/issues/844), the same defect from the fallback's end) | +| Owning row | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` ([#517](https://github.com/mudler/vllm.cpp/issues/517)), whose A2-Q1 unit ([#810](https://github.com/mudler/vllm.cpp/issues/810)) is the caller this blocks | +| Base | [`vt-fp8-shared-seam.md`](vt-fp8-shared-seam.md) ([#940](https://github.com/mudler/vllm.cpp/issues/940)) — the seam this makes reachable | +| Kind | bug — build gate / op registration | +| Branch | `row/VT-FP8-QUANT-ARCH-GATE-960` | + +## What is wrong + +`vt::QuantFp8Static`'s **only** CUDA registration is +`src/vt/cuda/cuda_matmul_fp8_cutlass.cu:376` @ `0e1bee42f`. `CMakeLists.txt:1668` +adds that translation unit to the `vllm` target **only when +`VT_CUTLASS_FP8_ARCHS` is non-empty**: + +```cmake +set(_FP8_CUTLASS_SOURCES) +if(VT_CUTLASS_FP8_ARCHS) + set(_FP8_CUTLASS_SOURCES src/vt/cuda/cuda_matmul_fp8_cutlass.cu) +endif() +``` + +The kernel body has **no cutlass dependency whatsoever** — zero `cutlass` / +`CUTLASS` tokens in `QuantFp8StaticKernelCuda` (`:353-370`). It is +`out[i] = e4m3_rne_sat(x[i] * (1/input_scale))`, a grid-stride elementwise loop +over a hardware convert intrinsic. It shared a TU with the cutlass sm120 fp8 +GEMM for authorship reasons and inherited that GEMM's arch set. + +On sm_110 (Thor) `VT_CUTLASS_FP8_ARCHS` is empty — `cutlass-fp8: DISABLED (no +requested arch in [110] provides it)` is that arch's **documented normal +profile**, not a misconfiguration — so `OpId::kQuantFp8Static` is not registered +for `DeviceType::kCUDA` at all. That is true of **every** CUDA arch outside the +cutlass-fp8 cell, not only Thor. + +**Nothing refuses first.** The op's GEMM partner `kMatmulFp8CublasLt` **is** +registered unconditionally (`src/vt/cuda/cuda_matmul.cu:920`), and the model-layer +predicate `MatmulFp8CutlassD` keys on *that*, so the guard passes. The missing +quant then resolves through `src/vt/op_provider.cpp:501` to the portable CPU +reference tier — eligible because `CudaBackend::UnifiedMemory()` is true — which +dereferences **device** pointers on the host and takes the process down: + +``` +[vt reference-tier] op=QuantFp8Static device=cuda has NO native kernel; running the PORTABLE CPU fallback (correct but slow) +SIGSEGV +``` + +Nothing silently dequantizes, and nothing refuses either: it crashes one call +later, under a banner that says "correct but slow". + +## Scope + +**In scope.** Relocate the `kQuantFp8Static` CUDA registration into a translation +unit that is unconditionally compiled for CUDA, so every CUDA arch gets the +native kernel; pin that placement with a runtime test and a structural checker. + +**Explicitly not in scope.** + +- The kernel's arithmetic, its dtype dispatch, and every dispatch condition — all + move byte-for-byte. +- #844's class: the reference tier still runs a host kernel over device pointers + for any *other* op that lacks a native CUDA kernel. This change removes one + live **instance** and does not address the class. #844 stays open, and the + refusal-instead-of-crash repair is a separate, larger change to the tier. +- Wiring NemotronH to anything (A2-Q1, #810 / #517). +- `MatmulFp8Cutlass` itself. Its arch gate is **correct**: it is a cutlass sm120 + kernel and on an arch that cannot build it a missing registration is an honest + refusal. The point of this row is precisely that the two ops are different in + this respect and must stop sharing a gate. + +## Upstream anchors + +Pinned oracle `5559679229bc961848b121ccdeaa8fa5d79bec98` +([`upstream-sync.md`](../upstream-sync.md)). + +| Ours | Upstream | +|---|---| +| `QuantFp8StaticKernel` | `csrc/quantization/w8a8/fp8/common.cuh:58-77` `scaled_fp8_conversion` (`:62` `x = val * scale`, `:68` clamp to ±448, `:71` hardware RNE convert) | +| the reciprocal formed once by the caller | `csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31` `1.0f / scale[...]` | +| per-tensor (one group over the tensor) | `csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:204-210` (`scale.numel() == 1`) | +| the method is static, `input_scale` a scalar | `vllm/model_executor/layers/quantization/modelopt.py:510-513`, `:528` | + +Upstream has no analogue of the defect to mirror: vLLM builds +`static_scaled_fp8_quant` from `csrc/quantization/w8a8/fp8/common.cu`, which is +in the unconditional `VLLM_EXT_SRC` list, while its cutlass `scaled_mm` sources +are added under `CUDA_ARCHS` intersections. **The relocation restores upstream's +own partition**, it does not invent one. + +## Design + +A new, unconditionally compiled TU: **`src/vt/cuda/cuda_quant_fp8.cu`**, added to +the `target_sources(vllm PRIVATE ...)` list directly inside `if(VLLM_CPP_CUDA)`. +The kernel, its two helpers and its registration move verbatim; only the local +`Check()` prefix changes from `matmul_fp8_cutlass` to `quant_fp8`, which was +wrong the moment the code moved and is not on any asserted path. + +**Why a new TU and not an existing one.** Three candidates were considered. + +- `cuda_matmul.cu` already hosts the unconditional fp8 GEMM registrations + (`kMatmulFp8CublasLt`, `:920`), which is the strongest argument for it — the + quant's partner already lives there. Against: it is the cuBLAS/cuBLASLt GEMM + wrapper TU, and an elementwise activation quant is not a GEMM. +- `cuda_ops.cu` is unconditional, already includes ``, and hosts + `RmsNormQuantFp8` — literally the *fused* arm of this same math, whose + `RmsNormF32ToFp8Dev` is deliberately the identical convert and whose + bit-identity claim to `RmsNorm(bf16) + QuantFp8Static` depends on it. Against: + it is a 3.6k-line general-kernel TU. +- **A file named for the op.** Chosen. The defect *is* "this kernel's compilation + is governed by a feature it does not use", and both alternatives re-create a + weaker form of it — the kernel's build would again be coupled to an unrelated + file's requirements and includes. A dedicated TU makes the invariant readable + in the CMake diff (the file is in the unconditional list, full stop), is what + the structural checker can assert without inference, is cheap to compile, and + is the natural home for the fp8 activation-quant family as it grows. It also + mirrors AGENTS.md §"Shared seams": new capability arrives as **additive files**. + +The cross-references between the three fp8 sites are written into all three +files, so the relation survives the split. + +## Risks + +| Risk | Handling | +|---|---| +| A behaviour change on GB10, where the op was already registered | Before/after on `dgx.casa` at `121a` with `cutlass-fp8: ENABLED` asserted in the configure log; four fp8 suites, identical case/assertion counts required | +| A duplicate registration masking the move | `RegisterOpProvider` takes first-registration-wins, so a stray second copy would be invisible at run time. Clause (d) of the checker forbids any second `kCUDA` registration of the op | +| Someone "fixes" it back by wrapping the registration in `#ifdef VT_CUTLASS_FP8` | Clause (c): the registration must sit at preprocessor-conditional depth 0. That mutation is a test case | +| The per-source gencode assignment | The new TU is not in `_VT_CUDA_FEATURE_SOURCES`, so `CMakeLists.txt:2231-2236` gives it the full `${VLLM_CPP_CUDA_ARCHITECTURES}` list — which is the point | +| A green checker that parsed nothing | The checker fails if the unconditional source list comes back empty, and `test_empty_source_list_is_not_a_pass` pins that | + +## Tests and gates + +**G4 (runtime pin), `tests/vt/test_ops_fp8_cpu.cpp`.** `OpRegistered(kQuantFp8Static, +kCUDA)` on any CUDA **build** — it needs no CUDA device, because the registration +is a table fill that runs before `main`, and "which build" is exactly the axis the +defect lived on. Its second assertion requires `kMatmulFp8Cutlass` to track +`VT_CUTLASS_FP8` instead, so the case proves the two are now **independent** +rather than merely that one of them is present: on Thor it reads +`CHECK(true) / CHECK_FALSE(false)`, on GB10 `CHECK(true) / CHECK(true)`. + +**G2 (byte gate), same file, pre-existing.** CPU vs CUDA `QuantFp8Static`, byte +for byte, zero tolerance, five scales. It could not be *run* on a non-cutlass-fp8 +CUDA arch before this change — it crashed. Landing this closes the arm that +[`vt-fp8-w8a8-cpu-arm.md`](vt-fp8-w8a8-cpu-arm.md) recorded as owed on that arch. + +**`scripts/check-cuda-op-arch-gate.py` (structural pin)** + its suite +`tests/scripts/test_check_cuda_op_arch_gate.py`. + +*Why both, argued rather than assumed.* G4 is the stronger statement: it observes +the property that matters — the op resolves for CUDA — instead of a proxy for it. +But it can only speak on a host that BUILT the CUDA backend **without** +cutlass-fp8, and **no CI job produces that build**: the GB10 gate host resolves +`cutlass-fp8: ENABLED`, where the defect is unreachable by construction, and every +other job is CPU-only. G4 would not have caught #960 before it landed; it caught +it here only because a human carried the binary to Thor. The checker reads the +build description, runs in the ordinary checker lane on every host including +CPU-only CI, and fails at PR time on the machine of whoever moves the +registration back. Neither instrument subsumes the other: the runtime test is the +claim, the checker is the tripwire. + +The checker asserts four clauses per entry, with no inference about what a kernel +"needs" — HOME (the TU is in the unconditional CUDA source list), REGISTERED +(exactly one live registration in it), UNGUARDED (at preprocessor depth 0), +EXCLUSIVE (no other CUDA source registers the same op for `kCUDA`). It runs the +C++ side through `checker_text.normalize_source`, so a commented-out or `#if 0`-ed +registration reads as absent, which is what it is. + +## Stop conditions + +- Stop and report `NEEDS_DECISION` if relocating is not behaviour-preserving on + some arch — i.e. if any GB10 suite differs before/after. +- Stop if the registration cannot be made unconditional for a demonstrable + reason. It can: the kernel compiles for sm_110 with zero warnings. +- Do not extend the checker's `REQUIRED` table beyond ops whose kernels are + genuinely arch-independent. For a cutlass/Marlin/FA2 kernel the feature gate is + the correct behaviour. + +## Evidence + +Base SHA `0e1bee42f16b5f3fb3ae5a23869f6fd97bfc037d`. + +### Thor (sm_110), CUDA 13.0.88, `-DVLLM_CPP_CUDA_ARCHITECTURES=110`, no cutlass + +Configure on all three builds: `CUDA target architectures: 110`, +`CUDA feature cutlass-fp8: DISABLED (no requested arch in [110] provides it)`, +`CUTLASS not found`. `BUILD_EXIT=0`, `warnings: 0`, `enospc: 0` each time. +Disk 319 G free before and after. + +| Tree | binary sha256 | `test_ops_fp8_cpu` | +|---|---|---| +| base `0e1bee42f` | `6b4d4df071a6…` | `test cases: 2 \| 1 passed \| 1 failed \| 2 skipped` · `assertions: 43 \| 43 passed \| 0 failed` · `Status: FAILURE!` · **exit 139 (SIGSEGV)** | +| base + G4 only (RED-first) | `63b7940e8609…` | G4 isolated: `test cases: 1 \| 0 passed \| 1 failed \| 4 skipped` · `assertions: 2 \| 1 passed \| 1 failed` · `Status: FAILURE!` · exit 1 | +| base + G4 + fix | `690bf71448ea…` | `test cases: 5 \| 5 passed \| 0 failed \| 0 skipped` · `assertions: 62 \| 62 passed \| 0 failed` · `Status: SUCCESS!` · exit 0 | + +The base run reproduces #960 verbatim, including the trap it names: +`assertions: 43 | 43 passed | 0 failed` printed beside `Status: FAILURE!`, so +anything grepping the assertions line alone reads a crash as green. + +The RED-first run is the important one. G4's **first** assertion failed +(`CHECK( vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCUDA) )` → +`values: CHECK( false )`) while its **second** passed +(`CHECK_FALSE(...kMatmulFp8Cutlass...)` → `CHECK_FALSE( false )`), so the case was +not vacuous and the arch genuinely lacks the cutlass GEMM. Non-zero case count in +both directions, and the case name carries no comma (`-tc` splits on commas: a +comma would have selected nothing and reported `SUCCESS!` with exit 0). + +The green run is where **G2** — CPU vs CUDA, byte for byte — executes on sm_110 +for the first time and passes: `bad == 0` at every one of five scales over 4096 +elements each, and no `[vt reference-tier]` banner is printed at all. 62 = the 60 +assertions this suite reports on GB10, plus G4's 2. + +### GB10 (sm_121a), `-DVLLM_CPP_CUDA_ARCHITECTURES=121a`, CUTLASS `$HOME/cutlass` + + + +### Local (CPU-only) + +`check-cuda-op-arch-gate --report`, its 14-case suite, `check-device-leakage` +(DSR 32 == baseline 32), and the full preflight. + +## Outcome + + + +## Now + +Landed on `main` via the row branch. `QuantFp8Static` is registered for CUDA on +every arch; the FP8 W8A8 arm is unblocked on non-cutlass-fp8 CUDA archs, which is +the base #810/#517 A2-Q1 needs. #844's class remains open. From 0d6fbe881267c600bb3d463056fd5b2156386c07 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 07:37:38 +0000 Subject: [PATCH 2/6] fix(VT-FP8-QUANT-ARCH-GATE): register QuantFp8Static from an unconditional CUDA TU `vt::QuantFp8Static` was registered for `DeviceType::kCUDA` from `src/vt/cuda/cuda_matmul_fp8_cutlass.cu`, a translation unit CMake adds to the target only when `VT_CUTLASS_FP8_ARCHS` resolves non-empty. The kernel needs no cutlass -- it is a grid-stride `x * (1/scale)` followed by a hardware e4m3 convert -- so it inherited an arch gate that has nothing to do with it. On sm_110, where `cutlass-fp8: DISABLED` is the documented normal profile, the op was simply absent for CUDA, the resolver installed the portable CPU reference tier for a CUDA queue, and the first call dereferenced device pointers: SIGSEGV, exit 139. The kernel, its two device helpers and its registration move verbatim into a new `src/vt/cuda/cuda_quant_fp8.cu`, listed in the unconditional `if(VLLM_CPP_CUDA)` source list. Only the local `Check()` message prefix changes, from `matmul_fp8_cutlass` to `quant_fp8`, which stopped being true the moment the code moved. No arithmetic, no dtype dispatch and no dispatch condition is touched. A dedicated file rather than `cuda_matmul.cu` (which already hosts the unconditional fp8 GEMM registrations) or `cuda_ops.cu` (which hosts the FUSED arm of this same math): the defect is that a kernel's compilation was governed by a feature it does not use, and both alternatives re-create a weaker form of that coupling by binding it to an unrelated file's includes and requirements. This also restores upstream's own partition -- vLLM builds `static_scaled_fp8_quant` from the unconditional `VLLM_EXT_SRC` list and gates only its cutlass `scaled_mm` sources. Two instruments, because neither covers the other. `test_ops_fp8_cpu` G4 observes the property that matters, `OpRegistered(kQuantFp8Static, kCUDA)`, and requires `kMatmulFp8Cutlass` to track `VT_CUTLASS_FP8` instead, so it proves the two are independent rather than that one is present. But it can only speak on a CUDA build WITHOUT cutlass-fp8, and no CI job produces one, so it would not have caught this before it landed. `scripts/check-cuda-op-arch-gate.py` reads the build description, runs on every host including CPU-only CI, and fails at PR time: HOME (the TU is in the unconditional list), REGISTERED (exactly one live registration), UNGUARDED (preprocessor depth 0, so an `#ifdef VT_CUTLASS_FP8` wrapper is not a pass), EXCLUSIVE (no second kCUDA registration elsewhere). Its own suite binds the checker as a MODULE rather than pulling three names out of it at import time, so `check-pr-size`'s creation-mutation stub fails every case instead of producing an ImportError the evidence contract reads as "executed no tests". `docs/USAGE.md` gains the operator-facing half: `cutlass-fp8: DISABLED` removes a GEMM, not FP8, and a `[vt reference-tier]` banner naming an op on a `cuda` device is a missing kernel to report rather than a slow path to accept. This removes one live INSTANCE of #844 and does not address its class: the reference tier still runs a host kernel over device pointers for any other op with no native CUDA kernel, and still calls it "correct but slow". #844 stays open. Also fixes #989 in flow, because it had to: registering the new checker's creation mutation means editing `scripts/check-pr-size.py`, and that file's evidence contract requires its own suite green at HEAD. That suite has been red on `main` since #888 added `.agents/reachability.md` without a path class, and red silently -- it runs in no CI job, so the only thing that loads it is the checker-evidence contract, which fires only when someone edits a checker. Third instance of the same class after #856 and #668, fixed the same way both were. Issue #960. Spec `.agents/specs/vt-fp8-quant-arch-gate.md`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/issue-index.md | 2 + .github/workflows/ci.yml | 4 + CMakeLists.txt | 6 + docs/FEATURES.md | 1 + docs/USAGE.md | 25 ++ scripts/agent-preflight.sh | 2 + scripts/check-cuda-op-arch-gate.py | 301 ++++++++++++++++++ scripts/check-pr-size.py | 13 + src/vt/cuda/cuda_matmul_fp8_cutlass.cu | 74 +---- src/vt/cuda/cuda_quant_fp8.cu | 125 ++++++++ tests/scripts/test_check_cuda_op_arch_gate.py | 233 ++++++++++++++ tests/scripts/test_check_pr_size.py | 3 + tests/vt/test_ops_fp8_cpu.cpp | 34 ++ 13 files changed, 766 insertions(+), 57 deletions(-) create mode 100644 scripts/check-cuda-op-arch-gate.py create mode 100644 src/vt/cuda/cuda_quant_fp8.cu create mode 100644 tests/scripts/test_check_cuda_op_arch_gate.py diff --git a/.agents/issue-index.md b/.agents/issue-index.md index e0ba51af7..af03c5c49 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -265,3 +265,5 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#975](https://github.com/mudler/vllm.cpp/issues/975) | `LTX25-IC-LORA` | The reference-image / reference-video arm of `ltx-2.5` is still refused, and BOTH reasons the refusal ever gave are now false: the IC-LoRA metadata ([#923](https://github.com/mudler/vllm.cpp/issues/923) reads it at load, `iclora_utils.py:30-49`) and the token-APPEND machinery ([#930](https://github.com/mudler/vllm.cpp/issues/930) built it in `c7cb59fbb`; the LAST-frame keyframe is SERVED on it). Two causes remain. (1) The reference CLIP has no pixel path: upstream reads it at `height // scale` by `width // scale` (`iclora_utils.py:116-117`), refuses a target the factor does not divide (`:112-115`), keeps frame 0 then every Nth frame (`temporal_subsample`, `:87-89`, called at `:144`) and encodes the clip (`:145-148`), while this engine's only pixel-to-latent route encodes ONE frame at the phase's own resolution and nothing reads `ref_video_dir` (`src/vllm/multimodal/video_engine.cpp:375`). `Ltx2ConvVideoEncode` already takes a `frame_count`, so the encoder is not the gap. (2) The reference item is a STAGE-1 item and stage 2 must run UNFUSED: `ic_lora.py:108` gives stage 1 `loras=tuple(loras)` and the reference conditioning (`:269-278`, `:377-402`), `:119` gives stage 2 `loras=()` and `:314-321` gives it `combined_image_conditionings` with no reference item — and this engine holds ONE DiT, fused at load, that every phase runs. `Ltx2LatentState` having no attention-mask field is NOT the reason either: the default arm builds no mask (`iclora_utils.py:159-160`, `:168-169`). Listed under `## Owed` in [`ltx25-ic-lora.md`](specs/ltx25-ic-lora.md) | feature | | [#940](https://github.com/mudler/vllm.cpp/issues/940) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | The FP8 W8A8 linear path is not a shared seam: `ResidentFp8`, `MatmulFp8CutlassD` and `MatmulFp8CutlassPreQuantD` lived in the anonymous namespace of `src/vllm/model_executor/models/qwen3_5.cpp` (`:1458`, `:1495`, `:1517` @ `c7cb59fbb`), so a second model could reach them only by copying them — the hand-rolled parallel path AGENTS.md §"Shared seams" forbids. NVFP4 already had `dense_nvfp4_gemm.h` + `compressed_tensors/schemes/nvfp4.h`; FP8 had neither half. Forced by `MODEL-NEMOTRON-H` ([#517](https://github.com/mudler/vllm.cpp/issues/517)), whose 46 FP8 W8A8 mamba `in_proj`/`out_proj` projections are 36.6% of decode bytes and 27.6% of GEMM FLOPs, and whose `in_proj` produces the fused `zxbcdt` the conv and the SSD scan consume — so that block cannot be split and has no device path at all without the seam. Extracted to `dense_fp8_gemm.h` + `quantization/fp8.h` with Qwen3.5 byte-identity as the gate; spec [`vt-fp8-shared-seam.md`](specs/vt-fp8-shared-seam.md) | bug | | [#974](https://github.com/mudler/vllm.cpp/issues/974) | — | The FP8 W8A8 resident helpers move weight bytes host->device without `vllm::load_stats::AddDeviceUpload` and without the post-upload `AdoptDeviceBytesAsHost`, while every other resident-weight helper in the same file performs both: `ResidentWeight` (`src/vllm/model_executor/models/qwen3_5.cpp:1009,1016 @ c7cb59fbb`) and `ResidentNvfp4` (`:1106,1111,1116,1121`), whose own comment states the obligation against [#150](https://github.com/mudler/vllm.cpp/issues/150). Affects `ResidentFp8` (`:1458`, now `dense_fp8_gemm.h`), `ResidentFp8Qkv` (`:1555`) and `ResidentFp8Qkvz` (`:3440`). Two unmeasured consequences: load accounting is short by the whole fp8 tower, and its device pages are never re-tagged, which is the shape of the GB10 weight-residency penalty on the 27B decode gap's largest attributed bucket. Found while extracting those entry points in [#940](https://github.com/mudler/vllm.cpp/issues/940) and deliberately NOT fixed there: a byte-identity gate cannot see a device -behaviour change hidden in a move. Listed under `## Owed` in [`vt-fp8-shared-seam.md`](specs/vt-fp8-shared-seam.md) | bug | +| [#960](https://github.com/mudler/vllm.cpp/issues/960) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `vt::QuantFp8Static`'s ONLY CUDA registration lived at `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:376` (@ `0e1bee42f`), and `CMakeLists.txt:1668` compiles that translation unit only when `VT_CUTLASS_FP8_ARCHS` is non-empty — yet the kernel body has ZERO cutlass tokens (`:353-370`): it is `out[i] = e4m3(x[i] * (1/input_scale))`, a grid-stride elementwise convert. So on every CUDA arch outside the cutlass-fp8 cell — sm_110/Thor is the measured one, and `cutlass-fp8: DISABLED for [110]` is that arch's DOCUMENTED NORMAL PROFILE, not a misconfiguration — `OpId::kQuantFp8Static` was not registered for `DeviceType::kCUDA` at all. Nothing refused first: the GEMM partner `kMatmulFp8CublasLt` IS registered unconditionally (`src/vt/cuda/cuda_matmul.cu:920`), so `MatmulFp8CutlassD`'s guard passed, and the missing quant then resolved through `src/vt/op_provider.cpp:501` to the portable CPU reference tier — eligible because `CudaBackend::UnifiedMemory()` is true — which dereferenced DEVICE pointers on the host and SIGSEGV'd one call later under a banner reading "correct but slow". Fixed by relocating the registration to a new unconditionally-compiled TU `src/vt/cuda/cuda_quant_fp8.cu`, which restores upstream's own partition (vLLM builds `static_scaled_fp8_quant` from the unconditional `VLLM_EXT_SRC` list and gates only its cutlass `scaled_mm` sources). This removes one live INSTANCE of [#844](https://github.com/mudler/vllm.cpp/issues/844) and does not address its class, which stays open. Unblocks the FP8 W8A8 arm on every non-cutlass CUDA arch — the base [#810](https://github.com/mudler/vllm.cpp/issues/810)/[#517](https://github.com/mudler/vllm.cpp/issues/517) A2-Q1 needs, where 46 FP8 mamba projections are 36.6% of decode bytes. Spec [`vt-fp8-quant-arch-gate.md`](specs/vt-fp8-quant-arch-gate.md) | bug | +| [#989](https://github.com/mudler/vllm.cpp/issues/989) | `VT-FP8-QUANT-ARCH-GATE` | `scripts/check-pr-size.py`'s `classify_path` has no entry for `.agents/reachability.md` (added by `POLICY-NOTHING-LANDS-DEAD`, [#888](https://github.com/mudler/vllm.cpp/issues/888) @ `8f49ac3be`), and it FAILS CLOSED, so `pr-size` — a REQUIRED check — refuses every pull request that touches that guide, and `tests/scripts/test_check_pr_size.py` has been red on `main` ever since. Red SILENTLY: that suite is wired into no CI job and is not in `agent-preflight.sh`'s `SUITES`, so the only thing that ever loads it is `check-pr-size`'s own executable-evidence contract, which fires only when a PR edits a checker — the red is reachable exclusively by the next person who must touch that file, and presents to them as their own breakage (the [#584](https://github.com/mudler/vllm.cpp/issues/584)/[#965](https://github.com/mudler/vllm.cpp/issues/965) shape). Third instance of the class after [#856](https://github.com/mudler/vllm.cpp/issues/856) (`issue-index.md` + the style guides) and [#668](https://github.com/mudler/vllm.cpp/issues/668) (`.agents/oracles/*`), both fixed in flow by the row that tripped over them. FIXED IN FLOW while landing [#960](https://github.com/mudler/vllm.cpp/issues/960), which could not register its new checker's creation mutation without touching `check-pr-size.py` at all. NOT fixed: wiring that suite into CI, which is its own change and would red `main` until this landed | bug | diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 0b8ccec10..b32f07b6f 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -186,6 +186,10 @@ jobs: run: | python3 scripts/check-fp4-resident-consistency.py python3 tests/scripts/test_check_fp4_resident_consistency.py + - name: An arch-independent CUDA op is registered from an unconditional TU + run: | + python3 scripts/check-cuda-op-arch-gate.py --report + python3 tests/scripts/test_check_cuda_op_arch_gate.py - name: Model decode is born on the runner (device-resident logits) run: | python3 scripts/check-runner-routing-consistency.py diff --git a/CMakeLists.txt b/CMakeLists.txt index a7763f60d..533d03863 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1584,6 +1584,12 @@ if(VLLM_CPP_CUDA) src/vt/cuda/cuda_dropin.cu src/vt/cuda/cuda_matmul.cu src/vt/cuda/cuda_quant_dot.cu + # UNCONDITIONAL ON PURPOSE (issue #960). The static per-tensor fp8 activation + # quant has no cutlass dependency; it lived in the cutlass-fp8 TU below and so + # went unregistered for kCUDA on every arch outside VT_CUTLASS_FP8_ARCHS, + # where the resolver then ran the portable CPU tier over device pointers and + # segfaulted. scripts/check-cuda-op-arch-gate.py fails if it leaves this list. + src/vt/cuda/cuda_quant_fp8.cu src/vt/cuda/cuda_matmul_nvfp4.cu src/vt/cuda/cuda_ops.cu src/vt/cuda/cuda_gdn.cu diff --git a/docs/FEATURES.md b/docs/FEATURES.md index d2a27237f..81974d45c 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -77,6 +77,7 @@ are our reading of their documented behavior, not measurements. | MXFP4 compressed-tensors | ◐ W4A16 Marlin, mem 2.63x less. gate_up FUSION + decode-graph default-ON; #44 3/3, 32B 6/6. **`VT_MARLIN_DENSE` DEFAULT-ON** (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA, byte-faithful, beats MoE (c8 0.969) | ✅ | ✅ | ☐ | | fp8 weights | ✅ | ✅ | ✅ | ☐ | | Per-tensor FP8 W8A8 linear is a shared seam any model can bind | ✅ `models/dense_fp8_gemm.h` + `layers::Fp8W8A8LinearMethod` (#940), bound via `layers::MakeLinearMethod`. One definition, CUDA only ([spec](../.agents/specs/vt-fp8-shared-seam.md)) | ✅ `Fp8LinearMethod` | ✅ | ☐ | +| FP8 W8A8 works on a CUDA arch without `cutlass-fp8` | ✅ `vt::QuantFp8Static` registers from an unconditional TU (#960); sm_110 measured ([spec](../.agents/specs/vt-fp8-quant-arch-gate.md)) | ✅ | ✅ | ☐ | | fp8-tower GDN `in_proj` emits bf16, unlocking packed GDN decode | ◐ `VT_GDN_FP8_IN_BF16` + `VT_GDN_PACKED_DECODE_FP8_TOWER` (inert alone), both default **OFF**, ungated (#339) ([spec](../.agents/specs/perf-fp8-alpha-fold.md)) | ✅ bf16 `out_dtype` | ☐ | ☐ | | Merged fp8 projection folds per-column alpha in the GEMM epilogue | ◐ `VT_FP8_ALPHA_VEC_EPILOGUE`, CUDA only, default off, ungated; refuses split-K under a bf16-D equivalence claim (`claims_splitk1_premise`, default off) | n/a | n/a | n/a | | `vt::MulColVecF32` carries a bf16 store width | ✅ f32 arm byte-identical; bf16 arm rounds once; CPU + CUDA | n/a | ☐ | ☐ | diff --git a/docs/USAGE.md b/docs/USAGE.md index 8994ea535..761b3d4c4 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -114,6 +114,31 @@ requested value — but it sent a contributor looking in the wrong place ([#168](https://github.com/mudler/vllm.cpp/issues/168)). The `build.ninja` gencode line remains the ground truth if you want to double-check. +### A DISABLED feature removes its kernels, not the ops that do not need it + +`cutlass-fp8: DISABLED` means this build has no CUTLASS sm120 FP8 **GEMM**. It +does not mean the build has no FP8. The static per-tensor activation quant +`vt::QuantFp8Static` is a hardware `e4m3` convert with no CUTLASS dependency, so +it is compiled and registered on **every** CUDA architecture +(`src/vt/cuda/cuda_quant_fp8.cu`), and the cuBLASLt FP8 GEMM it feeds is +registered unconditionally too. FP8 W8A8 checkpoints therefore load and run on a +CUDA build with no CUTLASS at all: `-DVLLM_CPP_CUTLASS_DIR` and +`-DVLLM_CPP_CUTLASS_FETCH` are not required for that path. + +Until [#960](https://github.com/mudler/vllm.cpp/issues/960) the quant shared a +translation unit with that CUTLASS GEMM, so it inherited the GEMM's architecture +set and was simply absent on `110`. The engine then ran the portable CPU fallback +over device pointers and the process died with `SIGSEGV` after printing + +```text +[vt reference-tier] op=QuantFp8Static device=cuda has NO native kernel; running the PORTABLE CPU fallback (correct but slow) +``` + +If you ever see that banner naming an op on a `cuda` device, this build is +missing a kernel it needs. Report it — it is not a slow path, and the message's +"correct but slow" is not true when the device is not the CPU +([#844](https://github.com/mudler/vllm.cpp/issues/844)). + ## Using more than one engine in a process Constructing a `LoadedEngine`, destroying it, and constructing another in the diff --git a/scripts/agent-preflight.sh b/scripts/agent-preflight.sh index 836d77c36..80dd7a4e6 100755 --- a/scripts/agent-preflight.sh +++ b/scripts/agent-preflight.sh @@ -71,6 +71,7 @@ CHECKERS=( check-env-doc check-fusion-consistency check-fp4-resident-consistency + check-cuda-op-arch-gate check-runner-routing-consistency check-surface-coverage check-test-registration @@ -113,6 +114,7 @@ SUITES=( test_checker_text test_check_fusion_consistency test_check_fp4_resident_consistency + test_check_cuda_op_arch_gate test_check_runner_routing_consistency test_check_surface_coverage test_check_test_registration diff --git a/scripts/check-cuda-op-arch-gate.py b/scripts/check-cuda-op-arch-gate.py new file mode 100644 index 000000000..8cc458195 --- /dev/null +++ b/scripts/check-cuda-op-arch-gate.py @@ -0,0 +1,301 @@ +#!/usr/bin/env python3 +"""Fail if a CUDA op that must exist on EVERY CUDA arch is registered from a +feature-gated translation unit. + +THE DEFECT THIS EXISTS TO PREVENT ALREADY HAPPENED (issue #960, and #844 is the +same defect seen from the fallback's end). + +`vt::QuantFp8Static`'s CUDA kernel is `out[i] = e4m3(x[i] * (1/scale))` — a plain +elementwise convert with no cutlass dependency of any kind. It lived in +`src/vt/cuda/cuda_matmul_fp8_cutlass.cu`, whose ONLY build gate is +`VT_CUTLASS_FP8_ARCHS`: CMake adds that TU to `_FP8_CUTLASS_SOURCES`, and hence +to the `vllm` target, only when the requested arch list intersects the +cutlass-fp8 feature cell. On sm_110 (Thor) that intersection is EMPTY, which is +the documented normal profile for the arch and not a misconfiguration. So the +kernel was simply not compiled, `OpId::kQuantFp8Static` was never registered for +`DeviceType::kCUDA`, and a CUDA queue asking for it resolved to the portable CPU +reference tier — which dereferenced device pointers and killed the process. + +NOTHING REFUSED FIRST, which is what made it expensive. The op's GEMM partner +`kMatmulFp8CublasLt` is registered unconditionally in `cuda_matmul.cu`, so the +model-layer predicate that keys on it passed, the build looked complete, and the +crash arrived one call later inside a "correct but slow" fallback banner. + +WHY A STRUCTURAL CHECK AND NOT ONLY A TEST. `tests/vt/test_ops_fp8_cpu.cpp` G4 +asserts the registration at run time and is the stronger statement — it observes +the property that actually matters rather than a proxy for it. But it can only +speak on a host that BUILT the CUDA backend without cutlass-fp8, and no CI job +does: the GB10 gate host resolves `cutlass-fp8: ENABLED`, where the defect is +unreachable by construction, and every other job is CPU-only. This file runs in +the ordinary checker lane on every host, reads the build description rather than +a binary, and therefore fails at PR time on the machine of whoever moved the +registration back. Neither instrument subsumes the other; the runtime test is the +claim, this is the tripwire. + +WHAT IS ASSERTED, per entry in `REQUIRED`, with no inference about what a kernel +"needs": + + (a) HOME — the named source file is listed in the UNCONDITIONAL CUDA + source list: the `target_sources(vllm PRIVATE ...)` whose + enclosing `if()` stack is exactly `[VLLM_CPP_CUDA]`. A file + moved under any further condition (`if(VLLM_CPP_CUTLASS)`, + `if(VT_CUTLASS_FP8_ARCHS)`, an `else()` branch, a `foreach`) + is NOT in that list and fails here. + (b) REGISTERED — that file contains exactly ONE + `RegisterOp(OpId::, DeviceType::kCUDA, ...)`. + (c) UNGUARDED — that registration sits at preprocessor-conditional depth 0 in + the file, so wrapping it in `#ifdef VT_CUTLASS_FP8` — which + would restore the exact original behaviour while leaving the + TU in the unconditional list — is not a pass. + (d) EXCLUSIVE — no OTHER CUDA source registers the same OpId for kCUDA. A + second, gated registration would let a runtime test pass on a + cutlass host for the wrong reason and would make (a)-(c) + describe a copy that is not the one being selected. + +TEXT THE COMPILER NEVER SEES IS NOT A PASS. The C++ side runs through +`checker_text.normalize_source`, so a commented-out or `#if 0`-ed registration +reads as absent, which is what it is. CMake `#` comments are stripped the same +way before the source list is parsed. + +WHAT THIS DOES NOT DO, stated plainly. It does not decide which ops BELONG in the +required set — that is a judgement recorded in `REQUIRED` with its issue, one +line per op, and adding a feature-gated op to it would be wrong. It does not +check that the kernel is correct, that the arch can run it, or that any other op +is registered anywhere. And it reads CMake lexically: a source list assembled +through a variable or a `foreach` is invisible to it and would be reported as a +MISSING home rather than silently accepted. + +Usage: + python3 scripts/check-cuda-op-arch-gate.py [--report] +""" + +from __future__ import annotations + +import argparse +import re +import sys +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) + +from checker_text import blank_out, normalize_source # noqa: E402 + +REPO = Path(__file__).resolve().parent.parent + +# The ops whose CUDA registration must not depend on a CUDA FEATURE. One line per +# op: the OpId, the TU that owns it, and the issue that put it here. Keep this +# short and argued — an op with a genuinely arch-specific kernel (the cutlass +# GEMMs, Marlin, FA2) does NOT belong in it, because for those the feature gate is +# the correct behaviour and a missing registration is an honest refusal. +REQUIRED = ( + # #960: the static per-tensor fp8 activation quant. No cutlass dependency; + # 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"), +) + +# Where a stray duplicate registration could hide. Every CUDA-side source. +CUDA_SOURCE_GLOBS = ("src/vt/cuda/*.cu", "src/vt/cuda/*.cpp", "src/vllm/platforms/cuda.cpp") + +_CMAKE_COMMENT = re.compile(r"#[^\n]*") +_CMAKE_TOKEN = re.compile(r"^[ \t]*([A-Za-z_][A-Za-z0-9_]*)[ \t]*\(", re.M) +_CPP_DIRECTIVE = re.compile(r"^[ \t]*#[ \t]*(if|ifdef|ifndef|elif|else|endif)\b[^\n]*", re.M) + + +def strip_cmake_comments(text: str) -> str: + """Blank `#` comments, position-preserving (CMake has no block comment form + we use). Bracket comments `#[[ ]]` are not used in this tree.""" + out = text + for m in reversed(list(_CMAKE_COMMENT.finditer(text))): + out = out[: m.start()] + blank_out(m.group(0)) + out[m.end() :] + return out + + +def _command_span(text: str, open_paren: int) -> int: + """Offset just past the `)` that closes the paren at `open_paren`.""" + depth = 0 + i = open_paren + while i < len(text): + if text[i] == "(": + depth += 1 + elif text[i] == ")": + depth -= 1 + if depth == 0: + return i + 1 + i += 1 + return len(text) + + +def unconditional_cuda_sources(cmake_text: str) -> set[str]: + """Sources given to `target_sources(vllm PRIVATE ...)` under an `if()` stack of + exactly `[VLLM_CPP_CUDA]` and no `foreach`/`function`/`macro` scope. + + The stack is tracked over the WHOLE file rather than by searching for a single + block, so moving the call one level deeper — into `if(VLLM_CPP_CUTLASS)`, into + an `else()`, into a `foreach()` — changes the answer, which is the motion this + checker exists to notice. + """ + text = strip_cmake_comments(cmake_text) + if_stack: list[str] = [] + other_depth = 0 + sources: set[str] = set() + + for m in _CMAKE_TOKEN.finditer(text): + cmd = m.group(1).lower() + open_paren = m.end() - 1 + end = _command_span(text, open_paren) + args = text[open_paren + 1 : end - 1] + + if cmd == "if": + if_stack.append(args.strip()) + elif cmd == "elseif": + if if_stack: + if_stack[-1] = args.strip() + elif cmd == "else": + # An `else()` branch is NOT the `if()` condition. Marking it with a + # sentinel keeps the stack depth right and keeps the branch out of the + # unconditional set. + if if_stack: + if_stack[-1] = "!" + if_stack[-1] + elif cmd == "endif": + if if_stack: + if_stack.pop() + elif cmd in ("foreach", "while", "function", "macro"): + other_depth += 1 + elif cmd in ("endforeach", "endwhile", "endfunction", "endmacro"): + other_depth = max(0, other_depth - 1) + elif cmd == "target_sources": + if if_stack != ["VLLM_CPP_CUDA"] or other_depth != 0: + continue + words = args.split() + if len(words) < 2 or words[0] != "vllm": + continue + for w in words[1:]: + if w in ("PRIVATE", "PUBLIC", "INTERFACE"): + continue + if "$" in w: # a variable expansion: not a literal home + continue + sources.add(w) + return sources + + +def cuda_registrations(op: str, root: Path = REPO) -> dict[str, list[tuple[int, int]]]: + """Every live `RegisterOp(OpId::, DeviceType::kCUDA` in the CUDA sources. + + Returns {repo-relative path: [(line, preprocessor depth), ...]}. + """ + pattern = re.compile( + r"RegisterOp\s*\(\s*OpId::" + re.escape(op) + r"\s*,\s*DeviceType::kCUDA\b" + ) + found: dict[str, list[tuple[int, int]]] = {} + for glob in CUDA_SOURCE_GLOBS: + for path in sorted(root.glob(glob)): + raw = path.read_text(encoding="utf-8", errors="replace") + text = normalize_source(raw) + hits = list(pattern.finditer(text)) + if not hits: + continue + # Preprocessor depth at each hit. Directive lines survive normalization + # (`strip_preprocessor_disabled` blanks bodies, not directives), so a + # `#if 0`-ed registration has already vanished from `text` above and a + # real `#ifdef` is still counted here. + depths: list[tuple[int, int]] = [] + marks = [(d.start(), d.group(1)) for d in _CPP_DIRECTIVE.finditer(text)] + for h in hits: + depth = 0 + for pos, kw in marks: + if pos >= h.start(): + break + if kw in ("if", "ifdef", "ifndef"): + depth += 1 + elif kw == "endif": + depth = max(0, depth - 1) + depths.append((text.count("\n", 0, h.start()) + 1, depth)) + found[str(path.relative_to(root))] = depths + return found + + +def check(report: bool = False, root: Path = REPO) -> list[str]: + errors: list[str] = [] + cmake = (root / "CMakeLists.txt").read_text(encoding="utf-8") + unconditional = unconditional_cuda_sources(cmake) + if not unconditional: + return [ + "CMakeLists.txt: found NO unconditional CUDA target_sources list. This " + "checker cannot pass vacuously -- either the build description moved or " + "the parser is broken; fix one of them." + ] + if report: + print(f"unconditional CUDA sources under if(VLLM_CPP_CUDA): {len(unconditional)}") + + for op, home, issue in REQUIRED: + regs = cuda_registrations(op, root) + if report: + print(f"{op} ({issue}) -> home {home}; registrations: {regs or '{}'}") + + # (a) HOME + if home not in unconditional: + errors.append( + f"{op} ({issue}): {home} is NOT in the unconditional CUDA source list " + f"(target_sources(vllm PRIVATE ...) directly under if(VLLM_CPP_CUDA)). " + f"A CUDA arch outside the feature set now gets no kernel for this op " + f"and falls to the portable CPU reference tier over DEVICE pointers." + ) + + # (b) REGISTERED and (c) UNGUARDED + here = regs.get(home, []) + if len(here) != 1: + errors.append( + f"{op} ({issue}): expected exactly ONE live " + f"RegisterOp(OpId::{op}, DeviceType::kCUDA, ...) in {home}, found " + f"{len(here)}. A commented-out or #if 0 registration reads as absent " + f"here, which is what it is to the compiler." + ) + else: + line, depth = here[0] + if depth != 0: + errors.append( + f"{home}:{line}: {op} ({issue}): the CUDA registration sits at " + f"preprocessor-conditional depth {depth}. Guarding it re-creates " + f"the arch gate the unconditional TU was meant to remove." + ) + + # (d) EXCLUSIVE + for path, hits in sorted(regs.items()): + if path == home: + continue + lines = ", ".join(str(ln) for ln, _ in hits) + errors.append( + f"{path}:{lines}: {op} ({issue}) is ALSO registered for kCUDA here. " + f"Its single home is {home}; a second registration lets a runtime " + f"check pass on a feature-enabled host for the wrong reason." + ) + + return errors + + +def main() -> int: + ap = argparse.ArgumentParser(description=__doc__.splitlines()[0]) + ap.add_argument("--report", action="store_true", help="print what was examined") + ap.add_argument( + "--root", + default=str(REPO), + help="repository root to examine (the checkout containing CMakeLists.txt)", + ) + args = ap.parse_args() + + errors = check(report=args.report, root=Path(args.root)) + if errors: + for e in errors: + print(f"check-cuda-op-arch-gate: {e}", file=sys.stderr) + print( + f"\ncheck-cuda-op-arch-gate: FAIL ({len(errors)} error(s))", + file=sys.stderr, + ) + return 1 + print(f"check-cuda-op-arch-gate: OK ({len(REQUIRED)} op(s) pinned to an unconditional CUDA TU)") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/check-pr-size.py b/scripts/check-pr-size.py index d7b4cdc9c..5af519b36 100755 --- a/scripts/check-pr-size.py +++ b/scripts/check-pr-size.py @@ -110,6 +110,14 @@ "CLAUDE.md", ".agents/workflow.md", ".agents/verification.md", + # The "nothing lands dead" task guide (#888, `8f49ac3be`). Never + # classified, and classify_path FAILS CLOSED, so pr-size refused every + # PR that touched it and this file's own suite has been red on main ever + # since -- silently, because that suite runs in no CI job and only the + # next checker change loads it (#989). Third instance of the same class + # after #856 and #668; listed explicitly rather than letting .agents/ + # become a blanket exemption. + ".agents/reachability.md", ".agents/porting.md", # The per-model coverage checklist that porting.md points at (#318). Same # procedure class as its sibling guides; listed explicitly rather than @@ -309,6 +317,11 @@ # reduced one. "scripts/check-container-matrix.py": DISABLED_CREATION_CHECKER, "scripts/check-container-workflow.py": DISABLED_CREATION_CHECKER, + # 2026-08-16: the CUDA arch-gate registration guard (#960). Created in the + # same PR, so there is no BASE version to mutate; its own suite loads the + # checker as a module and calls into it, so the disabled stub fails at import + # rather than quietly passing a reduced set of cases. + "scripts/check-cuda-op-arch-gate.py": DISABLED_CREATION_CHECKER, } SELF_CHECKER = "scripts/check-pr-size.py" EVIDENCE_TIMEOUT_SECONDS = 120 diff --git a/src/vt/cuda/cuda_matmul_fp8_cutlass.cu b/src/vt/cuda/cuda_matmul_fp8_cutlass.cu index 68598f8ba..1a22416ab 100644 --- a/src/vt/cuda/cuda_matmul_fp8_cutlass.cu +++ b/src/vt/cuda/cuda_matmul_fp8_cutlass.cu @@ -18,9 +18,12 @@ // tolerance of the two-stage form; the checkpoint scales ARE per-tensor. // // Isolated TU (heavy cutlass templates) — built only for sm_12{0,1}a. Pairs with -// QuantFp8Static (below), the static per-tensor activation quant that mirrors -// vLLM's static_scaled_fp8_quant (is_scale_inverted=False: x/input_scale, clamp, -// RNE hardware cvt). See .agents/specs/cutlass-dropin-feasibility.md. +// QuantFp8Static, the static per-tensor activation quant that mirrors vLLM's +// static_scaled_fp8_quant (is_scale_inverted=False: x/input_scale, clamp, RNE +// hardware cvt) — which lives in `src/vt/cuda/cuda_quant_fp8.cu` and is compiled +// UNCONDITIONALLY for CUDA, because it needs no cutlass and this TU's arch gate +// was silently withholding it from every other CUDA arch (issue #960). +// See .agents/specs/cutlass-dropin-feasibility.md. #include #include #include @@ -316,65 +319,22 @@ void MatmulFp8CutlassKernelCuda(Queue& q, Tensor& out, const Tensor& a_fp8, cons Check(cudaGetLastError(), "matmul_fp8_cutlass launch"); } -// ---- Static per-tensor fp8 activation quant (vLLM static_scaled_fp8_quant) --- -// inv = 1/input_scale; out_fp8[i] = fp8_e4m3(clamp(x[i]*inv, -448, 448)). -// A RECIPROCAL MULTIPLY, not a divide, and the reciprocal is hoisted out of the -// loop — that is upstream's shipped form (`x = val * scale` with the inverse -// formed by the caller: csrc/quantization/w8a8/fp8/common.cuh:62 and -// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31). The code below is -// RIGHT; do not "fix" it into `x / input_scale` to match a prose formula. The two -// differ by up to one f32 ulp before the fp8 round, and near an e4m3 tie that -// ulp changes the emitted byte on a default-ON 35B path. -// __NV_SATFINITE cvt saturates == clamp-then-cvt; RNE == vLLM's hardware cvt. -// Tin f32/bf16. -// -// The CPU arm (src/vt/cpu/cpu_ops.cpp QuantFp8StaticKernel) is INTENDED to be the -// byte-for-byte mirror of this kernel, and that equivalence is DECLARED AND OWED, -// not measured. It is gate G2 of .agents/specs/vt-fp8-w8a8-cpu-arm.md, which is -// PENDING for want of a GPU (#468). What IS measured is weaker and lives on the -// CPU side: G1 proves the CPU kernel matches an independent e4m3 reference derived -// from the format. Two implementations each matching a reference is not the same -// claim as the two matching each other, so do not cite this comment as evidence -// that they agree. Run tests/vt/test_ops_fp8_cpu.cpp on a CUDA host to close it. -__device__ __forceinline__ uint8_t F32ToFp8Dev(float f) { - return static_cast(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3)); -} -__device__ inline float LoadIn(const float* p, int64_t i) { return p[i]; } -__device__ inline float LoadIn(const __nv_bfloat16* p, int64_t i) { return __bfloat162float(p[i]); } - -template -__global__ void QuantFp8StaticKernel(uint8_t* out, const Tin* x, float input_scale, int64_t n) { - const int64_t step = static_cast(gridDim.x) * blockDim.x; - const float inv = 1.0f / input_scale; - for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; i < n; i += step) - out[i] = F32ToFp8Dev(LoadIn(x, i) * inv); -} - -void QuantFp8StaticKernelCuda(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scale) { - const int64_t n = x.shape[0] * x.shape[1]; - if (n == 0) return; - cudaStream_t s = AsStream(q); - const int blocks = static_cast(std::min((n + 255) / 256, 65535)); - switch (x.dtype) { - case DType::kF32: - QuantFp8StaticKernel<<>>(out_fp8.Ptr(), x.Ptr(), - input_scale, n); - break; - case DType::kBF16: - QuantFp8StaticKernel<__nv_bfloat16><<>>( - out_fp8.Ptr(), x.Ptr<__nv_bfloat16>(), input_scale, n); - break; - default: VT_CHECK(false, "cuda quant_fp8_static: unsupported x dtype (f32/bf16 only)"); - } - Check(cudaGetLastError(), "quant_fp8_static launch"); -} +// ---- Static per-tensor fp8 activation quant: NOT HERE ANY MORE (issue #960) -- +// `QuantFp8Static`'s CUDA kernel used to live below this line, and that was the +// defect. This TU is compiled ONLY when `VT_CUTLASS_FP8_ARCHS` is non-empty, so +// a kernel with no cutlass dependency whatsoever inherited cutlass's arch set +// and `OpId::kQuantFp8Static` went UNREGISTERED for `DeviceType::kCUDA` on every +// other CUDA arch — where the resolver then fell through to the portable CPU +// reference tier and dereferenced device pointers (SIGSEGV; #844 is the same +// defect from the fallback's end). It now lives in +// `src/vt/cuda/cuda_quant_fp8.cu`, which is in the unconditional +// `if(VLLM_CPP_CUDA)` source list. Do not move it back: +// `scripts/check-cuda-op-arch-gate.py` fails if you do. struct Registrar { Registrar() { RegisterOp(OpId::kMatmulFp8Cutlass, DeviceType::kCUDA, reinterpret_cast(static_cast(&MatmulFp8CutlassKernelCuda))); - RegisterOp(OpId::kQuantFp8Static, DeviceType::kCUDA, - reinterpret_cast(static_cast(&QuantFp8StaticKernelCuda))); } }; Registrar g_registrar; diff --git a/src/vt/cuda/cuda_quant_fp8.cu b/src/vt/cuda/cuda_quant_fp8.cu new file mode 100644 index 000000000..9de31b81f --- /dev/null +++ b/src/vt/cuda/cuda_quant_fp8.cu @@ -0,0 +1,125 @@ +// vllm.cpp — static per-tensor FP8 (e4m3) activation quant, CUDA arm. +// +// Mirror of vLLM's `static_scaled_fp8_quant` +// (csrc/quantization/w8a8/fp8/common.cuh:58-77 `scaled_fp8_conversion` and +// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31/:204-210, pinned +// oracle @ 5559679229bc961848b121ccdeaa8fa5d79bec98). `is_scale_inverted == +// false` at the call site, so the reciprocal is formed once by the caller and +// the elementwise math is a MULTIPLY. +// +// WHY THIS FILE EXISTS AT ALL — issue #960, and read it before moving anything +// back. This kernel used to live in `cuda_matmul_fp8_cutlass.cu`, whose sole +// build gate is `VT_CUTLASS_FP8_ARCHS` (CMakeLists.txt: the TU is added to +// `_FP8_CUTLASS_SOURCES` only when that variable is non-empty). The kernel has +// NO cutlass dependency of any kind — it is `x * (1/s)` followed by a hardware +// e4m3 convert — but sharing the translation unit made its REGISTRATION +// inherit cutlass's arch set. On every CUDA arch outside that set (sm_110/Thor +// is the measured one; it is not a Thor quirk) `OpId::kQuantFp8Static` was +// therefore not registered for `DeviceType::kCUDA` at all, so a CUDA queue +// asking for it fell through to the portable CPU reference tier, which +// dereferenced device pointers and took the process down with SIGSEGV (#844 is +// the same defect seen from the fallback's end). Its GEMM partner +// `kMatmulFp8CublasLt` is registered unconditionally in `cuda_matmul.cu`, so +// nothing upstream of the quant refused: the build looked complete and crashed +// one call later. +// +// So this TU is listed in the UNCONDITIONAL `if(VLLM_CPP_CUDA)` source list and +// carries no feature-gated include. Keep it that way: a kernel whose +// compilation is governed by a feature it does not use is the defect, and +// co-locating it with either the cutlass GEMM or the general op grab-bag would +// re-create a weaker form of the same coupling. +// `scripts/check-cuda-op-arch-gate.py` pins the invariant structurally; +// `tests/vt/test_ops_fp8_cpu.cpp` G4 pins it at run time. +#include +#include +#include + +#include +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::cuda { +namespace { + +void Check(cudaError_t err, const char* what) { + if (err != cudaSuccess) { + throw std::runtime_error(std::string("vt cuda: quant_fp8: ") + what + ": " + + cudaGetErrorString(err)); + } +} + +cudaStream_t AsStream(const Queue& q) { return static_cast(q.handle); } + +// ---- Static per-tensor fp8 activation quant (vLLM static_scaled_fp8_quant) --- +// inv = 1/input_scale; out_fp8[i] = fp8_e4m3(clamp(x[i]*inv, -448, 448)). +// A RECIPROCAL MULTIPLY, not a divide, and the reciprocal is hoisted out of the +// loop — that is upstream's shipped form (`x = val * scale` with the inverse +// formed by the caller: csrc/quantization/w8a8/fp8/common.cuh:62 and +// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31). The code below is +// RIGHT; do not "fix" it into `x / input_scale` to match a prose formula. The two +// differ by up to one f32 ulp before the fp8 round, and near an e4m3 tie that +// ulp changes the emitted byte on a default-ON 35B path. +// __NV_SATFINITE cvt saturates == clamp-then-cvt; RNE == vLLM's hardware cvt. +// Tin f32/bf16. +// +// The CPU arm (src/vt/cpu/cpu_ops.cpp QuantFp8StaticKernel) is the byte-for-byte +// mirror of this kernel. That equivalence is gate G2 of +// .agents/specs/vt-fp8-w8a8-cpu-arm.md; it is MEASURED on sm_110 and sm_121a +// (see .agents/specs/vt-fp8-quant-arch-gate.md — it could not be measured before +// #960 because this kernel was not registered on a non-cutlass-fp8 CUDA arch). +// The independent evidence on the CPU side is weaker and separate: G1 proves the +// CPU kernel matches an e4m3 reference derived from the format. +// +// The FUSED arm of this same math is `RmsNormQuantFp8` in cuda_ops.cu, whose +// `RmsNormF32ToFp8Dev` is deliberately the identical convert: that op's +// bit-identity claim to `RmsNorm(bf16) + QuantFp8Static` depends on it. Change +// one and you have silently changed the other's contract. +__device__ __forceinline__ uint8_t F32ToFp8Dev(float f) { + return static_cast(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3)); +} +__device__ inline float LoadIn(const float* p, int64_t i) { return p[i]; } +__device__ inline float LoadIn(const __nv_bfloat16* p, int64_t i) { return __bfloat162float(p[i]); } + +template +__global__ void QuantFp8StaticKernel(uint8_t* out, const Tin* x, float input_scale, int64_t n) { + const int64_t step = static_cast(gridDim.x) * blockDim.x; + const float inv = 1.0f / input_scale; + for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; i < n; i += step) + out[i] = F32ToFp8Dev(LoadIn(x, i) * inv); +} + +void QuantFp8StaticKernelCuda(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scale) { + const int64_t n = x.shape[0] * x.shape[1]; + if (n == 0) return; + cudaStream_t s = AsStream(q); + const int blocks = static_cast(std::min((n + 255) / 256, 65535)); + switch (x.dtype) { + case DType::kF32: + QuantFp8StaticKernel<<>>(out_fp8.Ptr(), x.Ptr(), + input_scale, n); + break; + case DType::kBF16: + QuantFp8StaticKernel<__nv_bfloat16><<>>( + out_fp8.Ptr(), x.Ptr<__nv_bfloat16>(), input_scale, n); + break; + default: VT_CHECK(false, "cuda quant_fp8_static: unsupported x dtype (f32/bf16 only)"); + } + Check(cudaGetLastError(), "quant_fp8_static 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. +struct Registrar { + Registrar() { + RegisterOp(OpId::kQuantFp8Static, DeviceType::kCUDA, + reinterpret_cast(static_cast(&QuantFp8StaticKernelCuda))); + } +}; +Registrar g_registrar; + +} // namespace +} // namespace vt::cuda diff --git a/tests/scripts/test_check_cuda_op_arch_gate.py b/tests/scripts/test_check_cuda_op_arch_gate.py new file mode 100644 index 000000000..3b0d78791 --- /dev/null +++ b/tests/scripts/test_check_cuda_op_arch_gate.py @@ -0,0 +1,233 @@ +#!/usr/bin/env python3 +"""Unit and mutation checks for scripts/check-cuda-op-arch-gate.py. + +The checker is the only mechanism that bites at PR time when a CUDA op that must +exist on every CUDA arch is moved back into a feature-gated translation unit +(issue #960). The runtime pin — `tests/vt/test_ops_fp8_cpu.cpp` G4 — is the +stronger claim but can only speak on a CUDA build WITHOUT cutlass-fp8, which no +CI job produces. So these cases prove the structural gate detects each way the +invariant can be broken, and then run it against the LIVE tree so a refactor that +makes the invariant unreachable cannot pass silently. + +EVERY MUTATION BREAKS EXACTLY ONE CLAUSE. A mutation that breaks two proves only +that the union fires. The four clauses (HOME / REGISTERED / UNGUARDED / +EXCLUSIVE) each get their own case, plus the two "text the compiler never sees" +disguises that the earlier generation of checkers in this tree passed on. + +THE VACUITY CASE MATTERS MOST. A checker that reports OK because it parsed +nothing is worse than no checker: it is a green light attached to no measurement. +`test_empty_source_list_is_not_a_pass` pins that. +""" + +from __future__ import annotations + +import importlib.util +import shutil +import subprocess +import sys +import tempfile +import unittest +from pathlib import Path + +ROOT = Path(__file__).resolve().parents[2] +CHECKER = ROOT / "scripts/check-cuda-op-arch-gate.py" +SPEC = importlib.util.spec_from_file_location("check_cuda_op_arch_gate", CHECKER) +assert SPEC is not None and SPEC.loader is not None +checker = importlib.util.module_from_spec(SPEC) +sys.modules[SPEC.name] = checker +SPEC.loader.exec_module(checker) + +# BOUND AS A MODULE, NOT AS THREE NAMES. `scripts/check-pr-size.py` proves a new +# checker red-before by replacing it with a disabled stub and re-running THIS +# module; pulling the functions out at import time would make that an ImportError +# rather than a run of failing cases, and the evidence contract reads an +# unimportable module as "executed no tests" instead of as a red. Every reference +# below goes through `checker.` so the stub fails each case on its own. + +HOME = "src/vt/cuda/cuda_quant_fp8.cu" +GATED = "src/vt/cuda/cuda_matmul_fp8_cutlass.cu" +REGISTRATION = ( + " RegisterOp(OpId::kQuantFp8Static, DeviceType::kCUDA,\n" + " reinterpret_cast(" + "static_cast(&QuantFp8StaticKernelCuda)));\n" +) + + +class FakeTree: + """A miniature checkout: CMakeLists.txt plus the two CUDA TUs, arranged + exactly as the real tree is, so a mutation can be applied to one of them.""" + + def __init__(self, cmake: str, home_src: str, gated_src: str) -> None: + self.dir = Path(tempfile.mkdtemp(prefix="cuda-op-arch-gate-")) + (self.dir / "CMakeLists.txt").write_text(cmake, encoding="utf-8") + cuda = self.dir / "src/vt/cuda" + cuda.mkdir(parents=True) + (cuda / "cuda_quant_fp8.cu").write_text(home_src, encoding="utf-8") + (cuda / "cuda_matmul_fp8_cutlass.cu").write_text(gated_src, encoding="utf-8") + + def __enter__(self) -> Path: + return self.dir + + def __exit__(self, *exc: object) -> None: + shutil.rmtree(self.dir, ignore_errors=True) + + +BASE_CMAKE = """\ +project(mini) +if(VLLM_CPP_CUDA) + target_sources(vllm PRIVATE + src/vt/cuda/cuda_matmul.cu + src/vt/cuda/cuda_quant_fp8.cu + src/vt/cuda/cuda_ops.cu) + if(VLLM_CPP_CUTLASS) + set(_FP8_CUTLASS_SOURCES) + if(VT_CUTLASS_FP8_ARCHS) + set(_FP8_CUTLASS_SOURCES src/vt/cuda/cuda_matmul_fp8_cutlass.cu) + endif() + target_sources(vllm PRIVATE ${_FP8_CUTLASS_SOURCES}) + endif() +endif() +""" + +BASE_HOME = f"""\ +#include "vt/ops.h" +namespace vt::cuda {{ +namespace {{ +void QuantFp8StaticKernelCuda(Queue&, Tensor&, const Tensor&, float) {{}} +struct Registrar {{ + Registrar() {{ +{REGISTRATION} }} +}}; +Registrar g_registrar; +}} +}} +""" + +BASE_GATED = """\ +#include "vt/ops.h" +namespace vt::cuda { +namespace { +struct Registrar { + Registrar() { + RegisterOp(OpId::kMatmulFp8Cutlass, DeviceType::kCUDA, + reinterpret_cast(&MatmulFp8CutlassKernelCuda)); + } +}; +Registrar g_registrar; +} +} +""" + + +def run(cmake: str = BASE_CMAKE, home: str = BASE_HOME, gated: str = BASE_GATED) -> list[str]: + with FakeTree(cmake, home, gated) as root: + return checker.check(root=root) + + +class TestCmakeParse(unittest.TestCase): + def test_reads_the_unconditional_list_only(self) -> None: + srcs = checker.unconditional_cuda_sources(BASE_CMAKE) + self.assertIn("src/vt/cuda/cuda_quant_fp8.cu", srcs) + # The cutlass TU is added under a NESTED if(), never at depth [VLLM_CPP_CUDA]. + self.assertNotIn("src/vt/cuda/cuda_matmul_fp8_cutlass.cu", srcs) + # ...and a variable expansion is not a literal home. + self.assertNotIn("${_FP8_CUTLASS_SOURCES}", srcs) + + def test_else_branch_is_not_unconditional(self) -> None: + cmake = BASE_CMAKE.replace( + "if(VLLM_CPP_CUDA)\n target_sources", + "if(VLLM_CPP_HIP)\nelse()\n target_sources", + ) + self.assertNotIn("src/vt/cuda/cuda_quant_fp8.cu", checker.unconditional_cuda_sources(cmake)) + + def test_cmake_comment_is_not_a_source(self) -> None: + cmake = BASE_CMAKE.replace( + " src/vt/cuda/cuda_quant_fp8.cu\n", + " # src/vt/cuda/cuda_quant_fp8.cu\n", + ) + self.assertNotIn("src/vt/cuda/cuda_quant_fp8.cu", checker.unconditional_cuda_sources(cmake)) + + +class TestMutations(unittest.TestCase): + def test_baseline_miniature_is_green(self) -> None: + # Non-vacuity for every case below: they must differ from a passing state. + self.assertEqual(run(), []) + + def test_HOME_moving_the_TU_under_the_cutlass_gate_goes_red(self) -> None: + # THE ORIGINAL DEFECT, reproduced: the TU is compiled only when the + # cutlass-fp8 arch set is non-empty. + cmake = BASE_CMAKE.replace(" src/vt/cuda/cuda_quant_fp8.cu\n", "").replace( + "set(_FP8_CUTLASS_SOURCES src/vt/cuda/cuda_matmul_fp8_cutlass.cu)", + "set(_FP8_CUTLASS_SOURCES src/vt/cuda/cuda_matmul_fp8_cutlass.cu" + " src/vt/cuda/cuda_quant_fp8.cu)", + ) + problems = run(cmake=cmake) + self.assertTrue(any("unconditional CUDA source list" in p for p in problems), problems) + + def test_REGISTERED_deleting_the_registration_goes_red(self) -> None: + problems = run(home=BASE_HOME.replace(REGISTRATION, "")) + self.assertTrue(any("expected exactly ONE" in p for p in problems), problems) + + def test_UNGUARDED_wrapping_the_registration_in_ifdef_goes_red(self) -> None: + # The subtle regression: the TU stays in the unconditional list, so clause + # (a) is satisfied, and the registration is still textually present, so a + # naive grep passes -- but the arch gate is exactly back. + guarded = BASE_HOME.replace( + REGISTRATION, "#ifdef VT_CUTLASS_FP8\n" + REGISTRATION + "#endif\n" + ) + problems = run(home=guarded) + self.assertTrue(any("conditional depth" in p for p in problems), problems) + + def test_EXCLUSIVE_a_second_gated_registration_goes_red(self) -> None: + problems = run(gated=BASE_GATED.replace( + " RegisterOp(OpId::kMatmulFp8Cutlass", REGISTRATION + " RegisterOp(OpId::kMatmulFp8Cutlass" + )) + self.assertTrue(any("ALSO registered for kCUDA" in p for p in problems), problems) + + def test_disguised_deletion_by_comment_goes_red(self) -> None: + # A commented-out registration is a deletion to the compiler. It must read + # as one here too -- the failure mode this tree has paid for before. + commented = BASE_HOME.replace( + REGISTRATION, + "".join("//" + ln + "\n" for ln in REGISTRATION.splitlines()), + ) + problems = run(home=commented) + self.assertTrue(any("expected exactly ONE" in p for p in problems), problems) + + def test_disguised_deletion_by_if_zero_goes_red(self) -> None: + disabled = BASE_HOME.replace(REGISTRATION, "#if 0\n" + REGISTRATION + "#endif\n") + problems = run(home=disabled) + self.assertTrue(any("expected exactly ONE" in p for p in problems), problems) + + def test_empty_source_list_is_not_a_pass(self) -> None: + # A parser that matches nothing must FAIL, not report OK. A green light + # attached to no measurement is the worst outcome available to a gate. + problems = run(cmake="project(mini)\n") + self.assertTrue(any("found NO unconditional" in p for p in problems), problems) + + +class TestLiveTree(unittest.TestCase): + def test_live_tree_passes(self) -> None: + self.assertEqual(checker.check(root=ROOT), []) + + def test_live_registration_is_where_the_checker_says(self) -> None: + # Pins the checker to the REAL file rather than only to miniatures: if the + # kernel is renamed or the TU disappears, this fails rather than drifting. + regs = checker.cuda_registrations("kQuantFp8Static", ROOT) + self.assertEqual(list(regs), [HOME], regs) + self.assertEqual([depth for _, depth in regs[HOME]], [0], regs) + self.assertFalse((ROOT / GATED).read_text(encoding="utf-8").count("kQuantFp8Static,")) + + def test_checker_cli_exits_zero_on_the_live_tree(self) -> None: + proc = subprocess.run( + [sys.executable, str(CHECKER), "--report"], + capture_output=True, + text=True, + check=False, + ) + self.assertEqual(proc.returncode, 0, proc.stderr) + self.assertIn("check-cuda-op-arch-gate: OK", proc.stdout) + + +if __name__ == "__main__": + unittest.main() diff --git a/tests/scripts/test_check_pr_size.py b/tests/scripts/test_check_pr_size.py index 225a87d5b..b05f92c41 100755 --- a/tests/scripts/test_check_pr_size.py +++ b/tests/scripts/test_check_pr_size.py @@ -490,6 +490,9 @@ def test_every_created_checker_has_closed_bootstrap_evidence(self) -> None: # every case rather than quietly passing a reduced one. "scripts/check-container-matrix.py", "scripts/check-container-workflow.py", + # 2026-08-16: the CUDA arch-gate registration guard (#960). Its suite + # reaches into the checker's parser, so the disabled stub cannot load. + "scripts/check-cuda-op-arch-gate.py", } self.assertEqual(set(checker.CREATION_MUTATIONS), expected) for path, mutation in checker.CREATION_MUTATIONS.items(): diff --git a/tests/vt/test_ops_fp8_cpu.cpp b/tests/vt/test_ops_fp8_cpu.cpp index e3cd4c16c..a3fde0a6a 100644 --- a/tests/vt/test_ops_fp8_cpu.cpp +++ b/tests/vt/test_ops_fp8_cpu.cpp @@ -451,3 +451,37 @@ TEST_CASE("the static fp8 W8A8 pair resolves and runs end-to-end on a CPU queue" // residual gap recorded in the spec is visible rather than assumed closed. CHECK_FALSE(vt::OpRegistered(vt::OpId::kMatmulFp8CublasLt, DeviceType::kCPU)); } + +// =========================================================================== +// G4 (issue #960) — THE REGISTRATION ITSELF, on every CUDA build. +// +// `QuantFp8Static`'s CUDA kernel is `x * (1/s)` plus a hardware e4m3 convert and +// has no cutlass dependency of any kind, but it USED TO SHARE a translation unit +// with the cutlass sm120 fp8 GEMM — and that TU is compiled only when +// `VT_CUTLASS_FP8_ARCHS` resolves non-empty. So on every CUDA arch outside that +// set (sm_110 is the measured one) `OpId::kQuantFp8Static` was not registered for +// `DeviceType::kCUDA` at all, the resolver installed the portable CPU reference +// tier for a CUDA queue, and the first real call dereferenced device pointers and +// took the process down. G2 above cannot state that: on a host WITHOUT the native +// kernel it crashes before it can report, and on a host WITH it the condition +// never arises. This case is the one that reads the same on both. +// +// It deliberately does NOT need a CUDA DEVICE — `OpRegistered` is a table lookup +// over registrars that ran before main, so it answers on any CUDA BUILD, which is +// exactly the axis the defect lived on. +#if defined(VLLM_CPP_CUDA) +TEST_CASE("G4: QuantFp8Static is registered for CUDA independent of cutlass-fp8") { + CHECK(vt::OpRegistered(vt::OpId::kQuantFp8Static, DeviceType::kCUDA)); + // NON-VACUITY, and the actual claim: the quant registration is INDEPENDENT of + // the cutlass one. Asserting only the line above would pass on a cutlass-fp8 + // host for the old reason as well as the new one. Here the cutlass GEMM is + // required to track its own feature macro, so on a build where it is ABSENT + // (Thor/sm_110) this case still proves the quant survived the arch gate, and on + // a build where it is PRESENT (GB10/sm_121a) it proves nothing regressed. +#if defined(VT_CUTLASS_FP8) + CHECK(vt::OpRegistered(vt::OpId::kMatmulFp8Cutlass, DeviceType::kCUDA)); +#else + CHECK_FALSE(vt::OpRegistered(vt::OpId::kMatmulFp8Cutlass, DeviceType::kCUDA)); +#endif +} +#endif // VLLM_CPP_CUDA From b05fd38a79ae2bbf5a0278b752fbc123d2985304 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 07:25:17 +0000 Subject: [PATCH 3/6] docs(VT-FP8-QUANT-ARCH-GATE): record the measured evidence and the outcome Fills the spec's evidence sections with the runs themselves rather than a summary of them: the Thor red-before at the base SHA (exit 139, and the `assertions: 43 | 43 passed | 0 failed` printed beside `Status: FAILURE!` that makes a crash read as green to anything grepping that line), the RED-first G4 result with its non-vacuity arm passing, the green-after, and a fresh-tree clean build at HEAD because three incremental rebuilds of one directory do not prove the committed tree builds. The GB10 before/after is the behaviour-preservation claim and is stated as such: four fp8 suites, identical case and assertion counts, distinct binary shas in both columns, `cutlass-fp8: ENABLED for [121a]` asserted in both configure logs. `test_linear_method` fails IDENTICALLY in both, which is how the run proves that failure is #907 and not ours -- named case, file:line, and assertion. Also recorded: the first GB10 after-run reported the suite at 4 cases when the binary should have had 5. `tar` restored the test file with an mtime older than the object compiled during the base build, so ninja skipped the compile and relinked only because libvllm.a had changed -- a suite reporting on a binary that contained the fix but not the test, with a sha that had legitimately moved. Verify the case COUNT, not the sha. Issue #960. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/vt-fp8-quant-arch-gate.md | 103 +++++++++++++++++++++++- 1 file changed, 99 insertions(+), 4 deletions(-) diff --git a/.agents/specs/vt-fp8-quant-arch-gate.md b/.agents/specs/vt-fp8-quant-arch-gate.md index c4a9ce577..cc58dcee3 100644 --- a/.agents/specs/vt-fp8-quant-arch-gate.md +++ b/.agents/specs/vt-fp8-quant-arch-gate.md @@ -191,6 +191,14 @@ Disk 319 G free before and after. | base `0e1bee42f` | `6b4d4df071a6…` | `test cases: 2 \| 1 passed \| 1 failed \| 2 skipped` · `assertions: 43 \| 43 passed \| 0 failed` · `Status: FAILURE!` · **exit 139 (SIGSEGV)** | | base + G4 only (RED-first) | `63b7940e8609…` | G4 isolated: `test cases: 1 \| 0 passed \| 1 failed \| 4 skipped` · `assertions: 2 \| 1 passed \| 1 failed` · `Status: FAILURE!` · exit 1 | | base + G4 + fix | `690bf71448ea…` | `test cases: 5 \| 5 passed \| 0 failed \| 0 skipped` · `assertions: 62 \| 62 passed \| 0 failed` · `Status: SUCCESS!` · exit 0 | +| HEAD, FRESH tree + clean build | `dc83e683f4fd…` | same: `5 \| 5 passed \| 0 failed \| 0 skipped` · `62 \| 62 passed \| 0 failed` · `Status: SUCCESS!` · exit 0 | + +The last row exists because the three above it are incremental rebuilds of one +tree, and an incremental build is not proof that the committed tree builds. It is +a fresh `git archive` of the branch into a new directory with no build state, +configured and built from scratch. It is at `b9a99ba11`; the only later commit +touches `tests/scripts/test_check_cuda_op_arch_gate.py`, which is in no C++ +target (`git diff --name-only b9a99ba11 HEAD`). The base run reproduces #960 verbatim, including the trap it names: `assertions: 43 | 43 passed | 0 failed` printed beside `Status: FAILURE!`, so @@ -211,16 +219,103 @@ assertions this suite reports on GB10, plus G4's 2. ### GB10 (sm_121a), `-DVLLM_CPP_CUDA_ARCHITECTURES=121a`, CUTLASS `$HOME/cutlass` - +Configure asserted on BOTH builds: `CUDA target architectures: 121a` and +`CUDA feature cutlass-fp8: ENABLED for [121a]`. A `DISABLED` line here would void +the result — it would measure the bug rather than the fix. `BUILD_EXIT=0`, +`warnings: 0`, `enospc: 0` both times, `-j 4`, disk 2.7 T free throughout. Every +binary sha differs between the two columns, so neither column is a stale artifact. + +| suite | BEFORE (`0e1bee42f`) | AFTER (branch) | +|---|---|---| +| `test_ops_fp8_cpu` | `4 \| 4 passed \| 0 failed` · `60 \| 60 passed \| 0 failed` · `SUCCESS!` | `5 \| 5 passed \| 0 failed` · `62 \| 62 passed \| 0 failed` · `SUCCESS!` | +| `test_ops_fp8_cutlass` | `8 \| 8 passed \| 0 failed` · `86 \| 86 passed \| 0 failed` · `SUCCESS!` | identical | +| `test_linear_method` | `10 \| 9 passed \| 1 failed` · `97 \| 95 passed \| 2 failed` · `FAILURE!` | identical | +| `test_ops_fused_chain` | `10 \| 10 passed \| 0 failed` · `583 \| 583 passed \| 0 failed` · `SUCCESS!` | identical | + +**GB10 is unchanged.** The only delta is `test_ops_fp8_cpu` gaining exactly G4: ++1 case, +2 assertions. The four pre-existing cases and their 60 assertions are +untouched, which is the claim — the registration moved translation units on a +host that already had it, and nothing about its behaviour moved with it. + +`test_linear_method` fails **identically in both columns**, which is how this run +proves it is not ours: `linear_method: MXFP4 fused gate_up ~= split (numerically) ++ fused path ran`, `test_linear_method.cpp:247`, `CHECK( after == before + 1 )` — +a Marlin dispatch counter, twice, on a suite with no fp8 arm. That is #907's +`test_linear_method` row, at the cutlass-enabled build's shape (10/97 rather than +the 8/85 #907 recorded, because `VT_MARLIN_NVFP4` adds cases). It is red at the +BASE SHA on this box, measured, not assumed. + +**A trap worth recording, because it nearly produced a false result.** The first +AFTER build reported `test_ops_fp8_cpu` at 4 cases / 60 assertions — G4 missing +from a binary whose sha had changed. `tar` had restored the test file with its +LOCAL mtime (06:46 UTC), older than the object compiled during the BASE build +(07:10 UTC), so ninja skipped the compile and only relinked because `libvllm.a` +had changed. The suite reported on a binary containing the fix but not the test. +The numbers above are from a rebuild after `touch`. Verify the case COUNT, never +just the sha. ### Local (CPU-only) -`check-cuda-op-arch-gate --report`, its 14-case suite, `check-device-leakage` -(DSR 32 == baseline 32), and the full preflight. +Release-equivalent CPU build (`-DVLLM_CPP_BUILD_TESTS=ON`), `BUILD_EXIT=0`, +**0 warnings**: `ctest` **489/489 passed, 0 failed** (2 skipped for absent +checkpoints: `test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). +`test_ops_fp8_cpu` reads 4 cases / 56 assertions here — G4 compiles out on a +non-CUDA build, which is correct and is why it is `#if defined(VLLM_CPP_CUDA)`. + +`scripts/agent-preflight.sh`: every gate green except +`test_cpu_x86_llamacpp_floor`, which exited `NO_QUIET_WINDOW` (4) at loadavg +77.49 while this box was building — the harness refusing to measure under +contention, #618, inherited. + +`check-cuda-op-arch-gate --report` OK, its 14-case suite OK, `check-device-leakage` +DSR 32 == baseline 32, `check-agent-record`, `check-public-doc-tables`, +`check-test-registration`, `check-surface-coverage`, `check-fusion-consistency`, +`check-fp4-resident-consistency`, `check-now-current`, `check-env-doc` all OK. + +### The checker's own mutation evidence, executed rather than described + +`scripts/check-pr-size.py --base --head HEAD` **passes**, and that +is not a formality: its evidence contract checks out the branch into a scratch +worktree, runs `tests/scripts/test_check_cuda_op_arch_gate.py` at HEAD (must +pass, non-zero case count), then overwrites the checker with a disabled stub and +re-runs the same module (must fail, non-zero case count). Both halves are +required and both are machine-verified. The same contract runs for +`scripts/check-pr-size.py` itself against `tests/scripts/test_check_pr_size.py`. + +That contract also caught a defect in the first version of the suite: it bound +`unconditional_cuda_sources` / `cuda_registrations` / `check` at import time, so +the stub produced an ImportError instead of failing cases, and the contract +reported `semantic evidence did not execute tests` — neither red nor green, the +instrument declining to say. Fixed by binding the module and going through +`checker.` inside each case, which is what the container-matrix and +container-workflow suites already do for the same reason. ## Outcome - +**What was measured.** The relocation is behaviour-preserving on a host that +already had the op (GB10: four fp8 suites, identical case and assertion counts, +identical pre-existing failure) and is the difference between a crash and a pass +on a host that did not (Thor: `exit 139` → `5 | 5 passed`, `62 | 62 passed`). + +**What was rejected.** `cuda_matmul.cu` and `cuda_ops.cu` as homes, argued under +[Design](#design): both would re-couple a cutlass-free kernel's compilation to an +unrelated file. Also rejected: fixing this by widening `MatmulFp8CutlassD`'s +guard or by making the reference tier refuse — the first hides the missing kernel +behind a refusal on a path that does work, and the second is #844's class, which +is a larger change to the tier and is deliberately still open. + +**Why the checker is written the way it is.** It asserts placement, not need. A +version that tried to infer "does this kernel require cutlass" by scanning the +body for `cutlass` tokens was rejected before it was written: it is transitive +through helpers, so it would be both false-positive and false-negative, and a +fuzzy gate is worse than none. The `REQUIRED` table is an argued list of two-line +entries instead, and adding a genuinely arch-specific kernel to it would be wrong. + +**What this did NOT close.** #844. The portable reference tier still accepts +`DeviceType::kCUDA` tensors and still calls itself "correct but slow" while +dereferencing device pointers. This row removed the one instance that was live on +a shipping path; the next feature-gated op to lose its native kernel will +reproduce it exactly. ## Now From 078d539e778eff4f8d921d8e2169388e9f45016e Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 08:40:39 +0000 Subject: [PATCH 4/6] docs(VT-FP8-QUANT-ARCH-GATE): state the row's real position in its `## Now` The `## Now` line said "landed on `main`", which it is not: it is open as #991, measured on both gate hosts, awaiting a fresh review and an operator merge. A `## Now` that anticipates its own landing is exactly the record that cannot be trusted afterwards -- the whole point of the line is that a reader learns the position without asking anyone. It also names the only red lanes and why they are not this row's: `windows-msvc-cpu`/`windows-msvc-vulkan` fail at the identical step on #988 and #982, which share no code with this change, so the arm is matched rather than assumed; the cause is #968 under #584's PR-only lane and a fix is already open as #983. Issue #960. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/vt-fp8-quant-arch-gate.md | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/.agents/specs/vt-fp8-quant-arch-gate.md b/.agents/specs/vt-fp8-quant-arch-gate.md index cc58dcee3..f78ac5a7f 100644 --- a/.agents/specs/vt-fp8-quant-arch-gate.md +++ b/.agents/specs/vt-fp8-quant-arch-gate.md @@ -319,6 +319,13 @@ reproduce it exactly. ## Now -Landed on `main` via the row branch. `QuantFp8Static` is registered for CUDA on -every arch; the FP8 W8A8 arm is unblocked on non-cutlass-fp8 CUDA archs, which is -the base #810/#517 A2-Q1 needs. #844's class remains open. +Open as [#991](https://github.com/mudler/vllm.cpp/pull/991), on +`row/VT-FP8-QUANT-ARCH-GATE-960-V2`, awaiting a fresh review and an operator +merge. Measured on both gate hosts and green; the only red CI lanes are +`windows-msvc-cpu` / `windows-msvc-vulkan`, which fail at the identical step on +unrelated open PRs (#988, #982) and are #968 under #584's PR-only lane, with a +fix already open as #983. + +`QuantFp8Static` is registered for CUDA on every arch, so the FP8 W8A8 arm is +reachable on a non-cutlass-fp8 CUDA arch — the base #810/#517 A2-Q1 needs. +#844's class remains open, and nothing here narrows it. From aa0eb92c41f8e16eb88bc885af70af580f7fe2cb Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 09:46:58 +0000 Subject: [PATCH 5/6] record(VT-FP8-QUANT-ARCH-GATE): re-append the two index rows after main's The merge of `3ce1cf7c7` produced a CLEAN but WRONG `.agents/issue-index.md`: every row from both sides survived, and `check-issue-index-append-only.py` passed, but main's newest row (#987) ended up AFTER this branch's #960 and #989 rather than before them. `merge=union` interleaves by hunk, not by arrival, so "no conflict" says nothing about order -- and the file's whole contract is that it is append-only, which is an ORDER claim. Detected by the check the union driver cannot make: is `origin/main`'s file a strict PREFIX of ours? It was not. Repaired by taking main's version wholesale and re-appending only this branch's own two rows, then verifying the shared prefix is byte-for-byte identical (251 rows) and that exactly #960 and #989 follow it. No row was edited and none was deleted. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/issue-index.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 5bff90f4e..660c8a5d6 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -266,6 +266,6 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#979](https://github.com/mudler/vllm.cpp/issues/979) | `BACKEND-GATE-CUDA-SGLANG` | Four-way GB10 benchmark on Qwen3.8-27B: state a common quantization PER PAIR, or record the pair as not-comparable. At the four recorded pins NO single quantization is common to all four engines, and two of the assumptions this campaign started from are wrong. vLLM at `555967922` has NO in-tree GGUF at all (`6635279d8` moved it to an unpinned out-of-tree `vllm-gguf-plugin`; `model_loader/__init__.py:33-49` has no `gguf` load format), so the vLLM-vs-llama.cpp pair is NOT COMPARABLE. SGLang's `qwen3_5` GGUF failure is NOT a two-line alias table: `loader.py:2129-2142` maps only `cohere` and `qwen3_moe`, but adding the alias would still hit three further blockers. Where the silence lives is the LOAD PATH, not `conv1d`: `loader.py:2149-2153`, `weight_utils.py:1280,1321` and `qwen3_5.py:1359-1412` carry no completeness guard, so a load that prints no error proves nothing. `conv1d` itself fails LOUDLY, because `qwen3_5.py:199` passes `quant_config=None` to that `ColumnParallelLinear` and `linear.py:176-179` then assigns `UnquantizedLinearMethod`, so `GGUFConfig.get_quant_method` is never reached for it. Reconciles two stale records with evidence: `BACKEND-GATE-CUDA-SGLANG` was `BLOCKED` on `SERVE-ASYNC-LLM`, which is discharged, and the SGLang oracle recorded `gateable = no` for two and a half weeks after it had run here. Also files the missing `BACKEND-GATE-CUDA-LLAMACPP` row. SGLang's published 38.28 tok/s is NVFP4 plus DSpark, a DRAFTED number, and no `dspark` speculator ships in `python/sglang/srt/speculative/` at `f63458b5`. DSpark is NOT absent from the pin: `docs_new/index.mdx:86,107,108,127` links the 2026-07-06 lmsys blog, three days before the pinned tree's own 2026-07-09 date, and `speculative/spec_info.py:60-70` registers out-of-tree algorithms at runtime, so a drafted SGLang arm needs a pin advance or a pinned plugin rather than an absence claim. **This row was corrected IN PLACE on 2026-08-16, before it landed**, which the append-only rule permits because the rule protects against a union merge duplicating a line that exists at the merge base and `origin/main` carries no `#979` row. Spec [`bench-qwen38-27b-four-way.md`](specs/bench-qwen38-27b-four-way.md) | perf | | [#940](https://github.com/mudler/vllm.cpp/issues/940) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | The FP8 W8A8 linear path is not a shared seam: `ResidentFp8`, `MatmulFp8CutlassD` and `MatmulFp8CutlassPreQuantD` lived in the anonymous namespace of `src/vllm/model_executor/models/qwen3_5.cpp` (`:1458`, `:1495`, `:1517` @ `c7cb59fbb`), so a second model could reach them only by copying them — the hand-rolled parallel path AGENTS.md §"Shared seams" forbids. NVFP4 already had `dense_nvfp4_gemm.h` + `compressed_tensors/schemes/nvfp4.h`; FP8 had neither half. Forced by `MODEL-NEMOTRON-H` ([#517](https://github.com/mudler/vllm.cpp/issues/517)), whose 46 FP8 W8A8 mamba `in_proj`/`out_proj` projections are 36.6% of decode bytes and 27.6% of GEMM FLOPs, and whose `in_proj` produces the fused `zxbcdt` the conv and the SSD scan consume — so that block cannot be split and has no device path at all without the seam. Extracted to `dense_fp8_gemm.h` + `quantization/fp8.h` with Qwen3.5 byte-identity as the gate; spec [`vt-fp8-shared-seam.md`](specs/vt-fp8-shared-seam.md) | bug | | [#974](https://github.com/mudler/vllm.cpp/issues/974) | — | The FP8 W8A8 resident helpers move weight bytes host->device without `vllm::load_stats::AddDeviceUpload` and without the post-upload `AdoptDeviceBytesAsHost`, while every other resident-weight helper in the same file performs both: `ResidentWeight` (`src/vllm/model_executor/models/qwen3_5.cpp:1009,1016 @ c7cb59fbb`) and `ResidentNvfp4` (`:1106,1111,1116,1121`), whose own comment states the obligation against [#150](https://github.com/mudler/vllm.cpp/issues/150). Affects `ResidentFp8` (`:1458`, now `dense_fp8_gemm.h`), `ResidentFp8Qkv` (`:1555`) and `ResidentFp8Qkvz` (`:3440`). Two unmeasured consequences: load accounting is short by the whole fp8 tower, and its device pages are never re-tagged, which is the shape of the GB10 weight-residency penalty on the 27B decode gap's largest attributed bucket. Found while extracting those entry points in [#940](https://github.com/mudler/vllm.cpp/issues/940) and deliberately NOT fixed there: a byte-identity gate cannot see a device -behaviour change hidden in a move. Listed under `## Owed` in [`vt-fp8-shared-seam.md`](specs/vt-fp8-shared-seam.md) | bug | +| [#987](https://github.com/mudler/vllm.cpp/issues/987) | `LTX25-RETAKE` | Two `ltx-2.5` refusal messages state reasons that are no longer true. (a) `src/vllm/multimodal/ltx2_video.cpp:1608 @ 0e1bee42f` says "nothing reads `ref_video_dir` at all", and MiniMax-H3 has always consumed the directory in full (`ReadReferenceClipChw`, `src/vllm/multimodal/minimax_h3_video.cpp:135 @ 0e1bee42f`, called at `:650`); [#975](https://github.com/mudler/vllm.cpp/issues/975) inherited the wider claim from this message. The claim that holds is narrower: the LTX-2.5 engine never reads the directory's CONTENTS. (b) `ltx2_video.cpp:1636-1638 @ 0e1bee42f` says "there is no AUDIO_VAE_ENCODER key filter", and `c2019b0e3` landed `Ltx2AudioVaeEncoderKeyRules()` (`include/vllm/model_executor/models/ltx2_audio_input.h:73 @ 0e1bee42f`) with a live call through `Ltx2EncodeAudioToLatent`. Both rewritten in the `WHAT IS *NOT* THE REASON` shape in the same flow, with one assertion tied to the LOCAL fact that the LTX side now reads the directory | bug | | [#960](https://github.com/mudler/vllm.cpp/issues/960) | `MODEL-TEXT-nemotron-h-nemotron-hfor-causal-lm` | `vt::QuantFp8Static`'s ONLY CUDA registration lived at `src/vt/cuda/cuda_matmul_fp8_cutlass.cu:376` (@ `0e1bee42f`), and `CMakeLists.txt:1668` compiles that translation unit only when `VT_CUTLASS_FP8_ARCHS` is non-empty — yet the kernel body has ZERO cutlass tokens (`:353-370`): it is `out[i] = e4m3(x[i] * (1/input_scale))`, a grid-stride elementwise convert. So on every CUDA arch outside the cutlass-fp8 cell — sm_110/Thor is the measured one, and `cutlass-fp8: DISABLED for [110]` is that arch's DOCUMENTED NORMAL PROFILE, not a misconfiguration — `OpId::kQuantFp8Static` was not registered for `DeviceType::kCUDA` at all. Nothing refused first: the GEMM partner `kMatmulFp8CublasLt` IS registered unconditionally (`src/vt/cuda/cuda_matmul.cu:920`), so `MatmulFp8CutlassD`'s guard passed, and the missing quant then resolved through `src/vt/op_provider.cpp:501` to the portable CPU reference tier — eligible because `CudaBackend::UnifiedMemory()` is true — which dereferenced DEVICE pointers on the host and SIGSEGV'd one call later under a banner reading "correct but slow". Fixed by relocating the registration to a new unconditionally-compiled TU `src/vt/cuda/cuda_quant_fp8.cu`, which restores upstream's own partition (vLLM builds `static_scaled_fp8_quant` from the unconditional `VLLM_EXT_SRC` list and gates only its cutlass `scaled_mm` sources). This removes one live INSTANCE of [#844](https://github.com/mudler/vllm.cpp/issues/844) and does not address its class, which stays open. Unblocks the FP8 W8A8 arm on every non-cutlass CUDA arch — the base [#810](https://github.com/mudler/vllm.cpp/issues/810)/[#517](https://github.com/mudler/vllm.cpp/issues/517) A2-Q1 needs, where 46 FP8 mamba projections are 36.6% of decode bytes. Spec [`vt-fp8-quant-arch-gate.md`](specs/vt-fp8-quant-arch-gate.md) | bug | | [#989](https://github.com/mudler/vllm.cpp/issues/989) | `VT-FP8-QUANT-ARCH-GATE` | `scripts/check-pr-size.py`'s `classify_path` has no entry for `.agents/reachability.md` (added by `POLICY-NOTHING-LANDS-DEAD`, [#888](https://github.com/mudler/vllm.cpp/issues/888) @ `8f49ac3be`), and it FAILS CLOSED, so `pr-size` — a REQUIRED check — refuses every pull request that touches that guide, and `tests/scripts/test_check_pr_size.py` has been red on `main` ever since. Red SILENTLY: that suite is wired into no CI job and is not in `agent-preflight.sh`'s `SUITES`, so the only thing that ever loads it is `check-pr-size`'s own executable-evidence contract, which fires only when a PR edits a checker — the red is reachable exclusively by the next person who must touch that file, and presents to them as their own breakage (the [#584](https://github.com/mudler/vllm.cpp/issues/584)/[#965](https://github.com/mudler/vllm.cpp/issues/965) shape). Third instance of the class after [#856](https://github.com/mudler/vllm.cpp/issues/856) (`issue-index.md` + the style guides) and [#668](https://github.com/mudler/vllm.cpp/issues/668) (`.agents/oracles/*`), both fixed in flow by the row that tripped over them. FIXED IN FLOW while landing [#960](https://github.com/mudler/vllm.cpp/issues/960), which could not register its new checker's creation mutation without touching `check-pr-size.py` at all. NOT fixed: wiring that suite into CI, which is its own change and would red `main` until this landed | bug | -| [#987](https://github.com/mudler/vllm.cpp/issues/987) | `LTX25-RETAKE` | Two `ltx-2.5` refusal messages state reasons that are no longer true. (a) `src/vllm/multimodal/ltx2_video.cpp:1608 @ 0e1bee42f` says "nothing reads `ref_video_dir` at all", and MiniMax-H3 has always consumed the directory in full (`ReadReferenceClipChw`, `src/vllm/multimodal/minimax_h3_video.cpp:135 @ 0e1bee42f`, called at `:650`); [#975](https://github.com/mudler/vllm.cpp/issues/975) inherited the wider claim from this message. The claim that holds is narrower: the LTX-2.5 engine never reads the directory's CONTENTS. (b) `ltx2_video.cpp:1636-1638 @ 0e1bee42f` says "there is no AUDIO_VAE_ENCODER key filter", and `c2019b0e3` landed `Ltx2AudioVaeEncoderKeyRules()` (`include/vllm/model_executor/models/ltx2_audio_input.h:73 @ 0e1bee42f`) with a live call through `Ltx2EncodeAudioToLatent`. Both rewritten in the `WHAT IS *NOT* THE REASON` shape in the same flow, with one assertion tied to the LOCAL fact that the LTX side now reads the directory | bug | From 142fcba7fcd88741abd862ab385ea050cb27613e Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sun, 16 Aug 2026 10:04:18 +0000 Subject: [PATCH 6/6] docs(VT-FP8-QUANT-ARCH-GATE): name the inherited MSVC red by its diagnostic "Known-red" is not a measurement, and this file was carrying it as one. The two `windows-msvc-*` jobs now carry the diagnostic that failed them: `warning C4244: '=': conversion from 'const double' to 'float'`, raised inside MSVC's own `` while compiling `src/vllm/multimodal/ltx2_video.cpp`, promoted to `error C2220` by `/WX`. That is #968, its fix is already open as #983, and `main` carries no baseline for those jobs because #584 makes them PR-only. The matched arm is recorded too: #988 and #982 fail at the identical step and share no file with this row. Also records the post-merge full gate: 491/491 rather than 489/489, because `3ce1cf7c7` adds two suites. Quoting the old number after a merge that changed the denominator would be a stale measurement presented as a current one. Issue #960. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/vt-fp8-quant-arch-gate.md | 19 +++++++++++++------ 1 file changed, 13 insertions(+), 6 deletions(-) diff --git a/.agents/specs/vt-fp8-quant-arch-gate.md b/.agents/specs/vt-fp8-quant-arch-gate.md index f78ac5a7f..29c891235 100644 --- a/.agents/specs/vt-fp8-quant-arch-gate.md +++ b/.agents/specs/vt-fp8-quant-arch-gate.md @@ -257,8 +257,10 @@ just the sha. ### Local (CPU-only) Release-equivalent CPU build (`-DVLLM_CPP_BUILD_TESTS=ON`), `BUILD_EXIT=0`, -**0 warnings**: `ctest` **489/489 passed, 0 failed** (2 skipped for absent -checkpoints: `test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). +**0 warnings**: `ctest` **489/489 passed, 0 failed** at the first base, and +**491/491 passed, 0 failed** after merging `3ce1cf7c7` (which adds two suites), +both with 2 skipped for absent checkpoints +(`test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). `test_ops_fp8_cpu` reads 4 cases / 56 assertions here — G4 compiles out on a non-CUDA build, which is correct and is why it is `#if defined(VLLM_CPP_CUDA)`. @@ -321,10 +323,15 @@ reproduce it exactly. Open as [#991](https://github.com/mudler/vllm.cpp/pull/991), on `row/VT-FP8-QUANT-ARCH-GATE-960-V2`, awaiting a fresh review and an operator -merge. Measured on both gate hosts and green; the only red CI lanes are -`windows-msvc-cpu` / `windows-msvc-vulkan`, which fail at the identical step on -unrelated open PRs (#988, #982) and are #968 under #584's PR-only lane, with a -fix already open as #983. +merge. Measured on both gate hosts. CI at `078d539e7`'s parent ran every job +green except `windows-msvc-cpu` / `windows-msvc-vulkan`, and those two are +inherited: MSVC raises `warning C4244: '=': conversion from 'const double' to +'float'` inside its own `` while compiling +`src/vllm/multimodal/ltx2_video.cpp`, promoted to `error C2220` by `/WX`. That is +#968, under #584's PR-only lane so `main` carries no baseline for it, a fix is +already open as #983, and the matched arm is measured rather than assumed — +#988 and #982 fail at the byte-identical step name and neither shares any file +with this row. `QuantFp8Static` is registered for CUDA on every arch, so the FP8 W8A8 arm is reachable on a non-cutlass-fp8 CUDA arch — the base #810/#517 A2-Q1 needs.