feat(KV-FP8): W3 -- half-sized KV blocks the runner actually serves, and a default scale that knows it was declared (#1593) - #1641
Merged
Merged
Conversation
… W1 guard that refused CUDA before the provider table (#1593) `KV-FP8` W1 landed the CPU half in 2026-07 and left the CUDA arm as a named later brick. It is now the critical path of benchmark campaign #1574, whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` in `hf_quant_config.json`, so no cell of that three-way can be served correctly without it. W1 is the ORACLE for this wave. Every gate here compares CUDA to the landed CPU kernels; nothing re-ports the numerics. THE STORE is a 1:1 port of the fp8 branch of vLLM's `reshape_and_cache_flash_kernel` (`csrc/libtorch_stable/cache_kernels.cu:314-401`) plus `CopyWithScaleOp` (`:241-252`) at pin `555967922`, restricted to upstream's `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — the only arm the op's wrapper admits, because the vt cache is the NHD unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. The converter is upstream's own `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)` (`quant_utils.cuh:497-503`), a true DIVIDE rather than the activation path's hoisted reciprocal multiply, and its byte-for-byte equality to the CPU software codec `vt::F32ToF8E4M3` is already MEASURED at zero tolerance on sm_110 and sm_121a (`.agents/specs/vt-fp8-quant-arch-gate.md` G2). THE READ adds `LoadKv(ptr, i, scale)` beside `Load`. It is INERT on the f32 and bf16 arms — they forward to `Load` unchanged, so every existing caller reads the same bytes in the same order — and on `uint8_t` it is upstream's `scaled_vec_conversion<float, uint8_t>` (`quant_utils.cuh:302-308`), written as the SAME ARITHMETIC as `vt::F8E4M3ToF32` so that CUDA==CPU on the read is a property of the source rather than of a measurement. Only the two correctness-grade kernels serve fp8, which is what the existing ladder already implies: the WMMA prefill kernels stage bf16 fragments, the vendored FA-2 launchers take bf16 pointers, and the vectorized decode-opt/GQA kernels read through `LoadRowN`/`LoadRow8`, 128-bit `uint4` loads specialized for bf16 and f32 only. Upstream draws the same line from the other side (`flash_attn.py:181-187,796-805`). WHAT ACTUALLY MADE THE ARM UNREACHABLE was neither kernel. Both W1 wrappers carried `VT_CHECK(q.device.type == DeviceType::kCPU, ... "a named later brick")` evaluated BEFORE provider lookup, so no CUDA kernel could ever have been reached however well it was registered. That is the RED this change was written against. The STORE now resolves through the provider table like every other op, because `kReshapeAndCacheFp8` is its own `OpId` that only CPU and CUDA register and an unimplemented backend refuses BY NAME inside `GetOp`. The READ cannot: it rides ADDITIVE fields on `PagedAttentionArgs` of an op `kMETAL` and `kROCM` already register for the FLOAT path, and nothing in the provider table separates the two arms, so an fp8 cache would reach a float kernel and return silent garbage. It keeps an explicit CPU-or-CUDA list whose message names the missing part. UNREACHED, DELIBERATELY, AND NOT NEW. Nothing calls the fp8 KV path from a production entry point on either backend: `vt::ReshapeAndCacheFp8` and `PagedAttentionArgs::kv_cache_dtype` have no caller outside their tests. W1 landed in that state and this does not change it. `KV-FP8` W3 owns the wiring — half-sized KV blocks in the runner, `--kv-cache-dtype` threaded from the CLI, and the checkpoint `k_scale`/`v_scale` path — and #1593 tracks it. Listed under `## Owed` in `.agents/specs/fp8-kv-cache.md`. THE DEVICE GATES ARE UNEXECUTED AND THE CUDA TRANSLATION UNITS ARE UNCOMPILED. The implementing session had no CUDA toolkit (`nvcc` is absent on `mudler-ubuntu-box`) and no device (the fleet was leased for #1574), so G2 (provider registration), G3 (store byte parity, f32 and bf16), G4 (paged-read parity, decode and prefill) and G5 (the e5m2 refusal) all skip with a MESSAGE naming what did not run. G1 and G1b — provider routing and the Metal/ROCm refusal — are the only cases that ran, and they run on the CPU leg. The first CUDA build or `rc` lease that touches this row must run `ctest -R test_cuda_fp8_kv_cache` before W2 counts as measured; the spec's `## Owed` says so. EVIDENCE. RED first: `test_cuda_fp8_kv_cache` 6 assertions / 6 failed, on `reshape_and_cache_fp8: only the CPU fp8-KV store is implemented in W1` at `ops.cpp:3478` and `paged_attention: only the CPU fp8-KV read is implemented in W1` at `ops.cpp:3811`. Green after: 6 cases / 10 assertions. Three negative mutations, each rebuilt, run and restored against a pre-taken sha256 — reinstate the store guard (3 failed), reinstate the read guard (7 failed), delete the Metal/ROCm refusal (4 failed). No sibling regressions: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Issue: #1593 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ld not fail for any CUDA defect (#1593, #1636) Repairs the six findings of the fresh review of #1593 W2. Four are records, two are gates, and the two merge-blocking ones land in the commit message, which `squash_merge_commit_message = PR_BODY` makes permanent. THE ANCHOR. Four new sites and the pull request body cited `scaled_vec_conversion<float, uint8_t>` at `quant_utils.cuh:302-308`. Verified against the pinned tree at `5559679229bc961848b121ccdeaa8fa5d79bec98`: lines 301-305 are the GENERIC primary template, whose body is `return x;`, and 307-314 are the `// fp8 -> half` `<uint16_t, uint8_t>` specialization. The `<float, uint8_t>` one is at `:419-429`. Same class, same change: `Fp8KVCacheDataType` was cited at `dtype_fp8.cuh:9-13`, the `#include <cuda_fp8.h>` guard; the enum is at `:15-19`. An anchor is how the next reader checks a port against the oracle, and one landing on an identity template invites the conclusion that the port is unfaithful. Three W1 copies of the same wrong anchor (`include/vt/fp8_kv.h:92`, `include/vt/ops.h:1129`, `src/vt/cpu/cpu_paged_attn.cpp:164`) are outside this change's authority and are filed as #1636, owned by `KV-FP8` and listed under the spec's `## Owed`. THE TRANSLATION UNITS COMPILE. The body said they do not. CI job `cuda-fat-build` built both changed files for `80;86;87;89;90a;100a;103a;110; 120a;121a` under `-Werror=all-warnings` and PASSED on `4d71e776e`, run 32495320287, job 96812232428. What stays true is that nothing has been EXECUTED on a device, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`. The spec now separates the two states. G5 COULD NOT FAIL FOR ANY CUDA DEFECT, and the fix needed three mutations to state correctly. The e5m2 refusal exists in the op wrapper, the CPU kernel and the CUDA kernel. Deleting the CUDA one on a CPU build gives `ninja: no work to do` and leaves the file 7/10 SUCCESS. Deleting the wrapper's leaves W1's `test_ops_fp8_kv_cache` GREEN at 8/511, because execution falls through to the CPU kernel's check and W1's case asserts `CHECK_THROWS_AS` on `std::runtime_error` rather than a message; only deleting both turns it red (7/8, 510/511). A layered refusal needs an assertion that names its layer, so G5 now resolves the registered provider with `GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA)`, calls it directly, and requires both `cuda reshape_and_cache_fp8` and `fp8_e5m2` in the message. TWO UNGATED INSTANTIATIONS. `LaunchPagedFp8Out<__nv_bfloat16, __nv_bfloat16>` is what a served bf16 model takes and what #1574's subject runs, and G4 exercised only `<float, float>`; the store's `DType::kF16 -> __half` arm was untouched while the wrapper's `IsFloat()` admits it. G4b and the f16 leg of G3 close both. Neither has a red-first mutation: they are device cases and this session has no `nvcc` and no device, which the first mutation above measures rather than assumes. They are listed under `## Owed`. TWO PROSE OVERSTATEMENTS. The store is elementwise-identical, not 1:1: upstream vectorizes the contiguous-heads arm through `vectorize_with_alignment<VEC_SIZE>` (`cache_kernels.cu:360-363`) and this is a scalar strided loop over the same elements in the same order, which is a bandwidth difference W4 owns. And the read is the CPU codec for every one of the 254 finite e4m3 codes, not "line for line": on `0x7F` and `0xFF` the CPU returns `quiet_NaN()` (`0x7FC00000`) and the device returns `CUDART_NAN_F` (`0x7FFFFFFF`), which no gate here can see because a NaN compares unequal to itself. Focused gate after: `test_cuda_fp8_kv_cache` 7 cases / 10 assertions SUCCESS, Release CPU build with `-Wall -Wextra -Werror`. Siblings unmoved: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Every mutation restored against a pre-taken sha256. Issue: #1593 Anchor debt: #1636 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…and a default scale that knows it was declared (#1593) W1 landed the CPU fp8-e4m3 KV store and read, W2 the CUDA arm, and both landed with nothing reaching either from a production entry point. W3 is the wiring: `--kv-cache-dtype` on the server flag, the checkpoint's own `kv_cache_quant_algo` honoured when no flag is typed, KV blocks sized at one byte per element, and the `k_scale`/`v_scale` path. The resolution chain mirrors vLLM at `555967922` step for step. `vllm::ResolveKvCacheDTypeString` (`include/vllm/config/cache.h`) is `resolve_kv_cache_dtype_string` + `get_kv_cache_quant_algo_string` (`utils/torch_utils.py:374-392,310-362,64-67`), called ONCE from `LoadedEngine::FromModelDir` where `EngineArgs.create_engine_config` calls it (`arg_utils.py:1915-1929`); an explicit value is returned unchanged and the checkpoint is never consulted (`:380-381`). `vllm::v1::ApplyCacheDType` is the runner's `self.kv_cache_dtype` reaching every attention spec (`gpu_model_runner.py:484-486`), and the halving is upstream's own arithmetic: `AttentionSpec.real_page_size_bytes` is linear in `get_dtype_size(self.dtype)` (`kv_cache_interface.py:204-218`) and every fp8 CacheDType stores as `torch.uint8` (`torch_utils.py:38-40`). THE ORDER IS THE FEATURE. `ApplyResolvedCacheDType` runs on the PROBE config, before `ResolveNumBlocks` reads its geometry, because the probe's `KVBytesPerBlock` is the divisor knob 2 sizes the pool with. Applying it afterwards serves the same pool in half the bytes instead of twice the pool, and the gate that catches that is an equality (`fp8_blocks == 2 * bf16_blocks` at one `--kv-cache-memory`), not a ratio. THE TRAP THE #1574 GATE CHECKPOINT SETS. `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a2` declares `kv_cache_quant_algo: "FP8"` and ships ZERO `k_scale`/`v_scale` tensors -- measured from the public `model.safetensors.index.json`: 2001 tensors, no `k_scale`, no `v_scale`, no `kv_scale`. So it serves on the default scale 1.0, and so would a checkpoint that declared no KV quantization at all if the default were reached by falling off the end of a missing-tensor lookup. The two are indistinguishable at runtime until somebody ships the second kind, at which point the accidental path invents a scale for a cache nobody asked to quantize. Upstream keeps them apart structurally -- the scale block runs only under `is_quantized_kv_cache` (`kv_cache.py:100-102`), and the both-sentinels arm inside it is a separate branch (`:112-116`) -- and `KvScaleOrigin` mirrors that with four named arms. `ScalesForFp8Store` refuses `kNotQuantized` BY NAME rather than answering 1.0, and G2 gates the difference between two calls whose numbers are identical. Half-sized blocks are silent-corruption territory, so the store and the read read the SAME `PagedKvCache::fp8_kind` the runner copied off the spec, and `IsFp8KvCache` refuses a view whose storage dtype and interpretation disagree. Routed: the shared `dense_attn::AttnBlock` seam and `qwen3_5.cpp` (the Qwen3.5/3.8 family the campaign measures). Everything else is refused BY NAME at `vt::ReshapeAndCache`, which now names `vt::ReshapeAndCacheFp8` and says the architecture is not routed, instead of letting a float store index a half-sized page. Gate: `tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp`, 19 cases / 87 assertions GREEN on a CPU-only Release build. G4 and G5 enter through the LoadedEngine constructor rather than by building a spec by hand, so G5 is a real reachability gate: an fp8 engine generates tokens through `Qwen3_5DenseModel::Forward` over a one-byte cache. UNREACHED AND OWED, all under #1593 and listed in the spec's `## Owed`: the C ABI carries no `kv_cache_dtype` field (`include/vllm.h` and `src/capi/` were outside this dispatch's authority), so a C-ABI caller reaches fp8 KV only through a declaring checkpoint; 17 architectures with their own attention preambles refuse rather than route; no weight loader extracts `k_scale`/`v_scale`, so the two checkpoint-loaded arms of the resolver are unit-gated and unreached; and every CUDA case in this feature is still UNEXECUTED for the reason W2 recorded. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ry production weight arm (#1593) The fresh review of PR #1641 returned FAIL with a functional bug, three unreached surfaces, two unrecorded facts and a stale record. This commit merges `origin/main` -- W2 landed as squash `53f5b74be`, so the pull request's base is dead -- and repairs all of them. THE BUG. `FullAttnBlockPaged` cast K and V to the cache dtype only under `if (kv.dtype == DType::kBF16)`, and on the fp8 route `kv.dtype` is `kI8`. K still arrived as f32, because `attn_dt` is f32 for every fp8 cache: both FA2 eligibility tests carry `kv.dtype == DType::kBF16` as a term. V arrived as whatever the v_proj GEMM emitted, which is bf16 on the block-wise fp8 arm, bf16 on the NVFP4 arm under the default `VT_BF16_GEMM_OUT`, and bf16 on ordinary torch safetensors. `vt::ReshapeAndCacheFp8` quantizes from ONE source dtype, so `--kv-cache-dtype fp8` threw at the first forward on every weight arm a user has -- including the #1574 campaign checkpoint this wave exists to serve. The CPU gate could not see it: `MakeDenseWeights` builds weights with no `nk` flag, which take `MatmulF32D`, and f32 pairs with f32. The fp8 route now normalises both to bf16 exactly as the bf16-cache route already did, and bf16 rather than f32 because bf16 is what upstream quantizes from -- its model IS bf16 where `reshape_and_cache_flash` takes key/value. THE MERGE. `main`'s SPEC-DRAFTER-CHAIN W1 refusal and this row's resolution stanza both insert as the FIRST statement of `FromModelDir`, and both are load-bearing: that row's G5 requires the refusal before any weight I/O, and `ReadQuantConfigJson` opens a file inside `model_dir`. Taking either side of the conflict drops a guarantee, so both survive, chain first, and G10 gates the order -- `test_drafter_chain_reach.cpp` cannot, because it points at a nonexistent directory that `ReadQuantConfigJson` answers "" for without opening anything. `.agents/engine-matrix.md`, `.agents/quantization-matrix.md` and `.agents/issue-index.md` take the target-branch version with the scoped edit re-applied by hand; the union merge had duplicated two index rows. THREE UNREACHED SURFACES, now entered through production entry points. The loader's own resolution stanza and `ReadQuantConfigJson` (deleting the whole stanza left the gate 19/19, 89/89 SUCCESS) are gated by G10 through `LoadedEngine::FromModelDir`. The `--kv-cache-dtype` flag itself (deleting both its parse arm and its assignment left the same green, and no test in the tree named the flag) is gated by a new `test_serve_kv_cache_dtype.cpp` that drives the REAL `VllmServerMain`. The `per_layer_attn_specs` retype branch is gated by G11 on the arithmetic it claims. TWO MIRRORS AND TWO RECORDS. `ReadQuantConfigJson` read `hf_quant_config.json` first; upstream reads `config.json:quantization_config` first and consults the legacy file only when that is absent, which matters for a checkpoint re-quantized in place. And enabling fp8 KV takes the model off FA-2 prefill and decode, the WMMA ladder and the vectorized decode kernels, all bf16-native, so the 11.06 vs 9.71 tok/s lead this wave exists to make matched may not survive turning it on. Nothing here measured that, and the spec, `docs/USAGE.md`, `docs/FEATURES.md` and both matrix rows now say so instead of implying the memory win is free. Gates: `test_kv_cache_fp8_wiring` 26 cases / 120 assertions and `test_serve_kv_cache_dtype` 3 cases / 26 assertions, both GREEN, CPU-only Release. Eight negative mutations, each applied with its sha256 and diff printed, built, run red and restored byte-identically: the fp8 normalisation, the flag's parse arm, the flag's assignment, the loader's checkpoint read, the file precedence, the per-layer retype, and the chain refusal both deleted and MOVED below the stanza. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
localai-bot
marked this pull request as ready for review
August 22, 2026 07:38
…ist, and USAGE did not say a refusal arrives after the sizing (#1593) The `FromModelDir` comment explaining why the drafter-chain refusal must precede the KV-FP8 resolution stanza cited `ChainRefusalPrecedesKvCacheDTypeResolution`, which is not the name of anything. The case is "kv-fp8 W3 G10: the drafter-chain refusal runs BEFORE the KV resolution" in `tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp`, and the comment now says so together with the reason `test_drafter_chain_reach.cpp` cannot hold the same order: its model path does not exist, so `ReadQuantConfigJson` answers "" for it without opening anything and the refusal arrives either way. An anchor a reader cannot follow is how a correct guarantee gets deleted by the next change. `docs/USAGE.md` gains the operator-visible consequence of that same ordering. `ApplyCacheDType` retypes every attention spec, group and heterogeneous per-layer alike, before an unrouted architecture's store refuses -- so an operator who passes `--kv-cache-dtype fp8` to Gemma-4 sees a doubled block count in the startup line and THEN a named refusal, rather than a served run. The sizing is the half a wrong answer would corrupt in silence, so it is made consistent first; the page said nothing about that sequence and a doubled count followed by a failure reads as a bug otherwise. Comment and document only. `test_kv_cache_fp8_wiring` 26/120 and `test_serve_kv_cache_dtype` 3/26 rebuilt and rerun GREEN. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
`origin/main` moved while this branch's gate was running, so the trailer and style gates skipped rather than reporting on this tree. Merged with no conflict: main's side is MUSIC3-DIT-ARM-REACH (#1131) plus its records, which touch no file this row changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
Second catch-up merge in this session: `origin/main` moved twice while the branch's gate was running, and the trailer and style gates skip rather than report on a branch that is behind. No conflict. Main's side is SERVE-REQUEST-LENGTH-GUARD and AB-ARMS-CONTROL plus their records and tests, which touch no file this row changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
Merge up so `scripts/agent-preflight.sh` stops SKIPPING `commit-trailers` and `commit-style`: both are diff-scoped against `origin/main`, and a base that is no longer an ancestor makes them report nothing at all about this tree rather than reporting a pass. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ritten to serve (#1593) `dense_attn::AttnBlock` is what AGENTS.md `## Shared seams` names as the decode seam, and W3 listed it as routed for fp8 KV. It was not. The routing was there -- `fp8_kv`, `WriteKvCache`, `ApplyKvCacheQuant` -- underneath a preamble guard that still admitted only `kBF16` and `kF32`, while `IsFp8KvCache` is true only for `kI8`. `fp8_kv` was therefore provably false at every call and neither fp8 arm could execute. That is a contradiction between two lines of one function, not an inference. The gate could not see it. A fresh review reverted the ENTIRE routing from the seam and the focused suite stayed at 26 cases / 120 assertions SUCCESS, because every case entered through `Qwen3_5DenseModel::Forward` and none entered here. The seam is the production forward for Qwen3 dense (`qwen3.cpp:185`), Qwen3-MoE (`qwen3_moe.cpp:84`), Voxtral (`voxtral.cpp:102`) and the Llama, Mistral and InternLM2 registries that share `Qwen3DenseModel`, so `--kv-cache-dtype fp8` on any of them sized the pool at half and then threw a message naming neither fp8, nor the flag, nor the row. The guard is widened exactly as `qwen3_5.cpp:5313` already was. G12 is the gate: it enters an fp8 cache through `Qwen3DenseModel::Forward`, the function `ForwardQwen3ForCausalLM` calls under `ModelRegistry::Forward` (`qwen3_dense.cpp:113`), and asserts that the half-width pages carry bytes after the forward, that two fresh stacks agree bit for bit, and that the fp8 logits DIFFER from the same forward over a bf16 cache -- the last one because an fp8 cache that read back like a float one would mean the dequant never happened. The second review also reported `GetKvCacheQuantAlgoString`'s acceptance of `producer.name` as a divergence, because `get_kv_cache_quant_algo_string` (`torch_utils.py:319`) gates on a top-level `quant_method` and nothing else. That reads one function where the answer needs two. `ModelArchConfigConvertorBase._normalize_quantization_config` (`transformers_utils/model_arch_config_convertor.py:208-247`) runs first, at `ModelConfig.__post_init__` (`config/model.py:577`), and INJECTS `quant_cfg["quant_method"] = "modelopt"` when `producer["name"] == "modelopt"`, mutating the same dict `hf_config.quantization_config` names. Both functions were extracted with `ast` from the files at pin `555967922` and RUN over the live documents rather than retyped: `nvidia/Llama-3.3-70B-Instruct-FP8`'s producer-only `hf_quant_config.json` answers `None` before normalization and `'fp8_e4m3'` after it. Accepting the producer name is the mirror; the comment that justified it on other grounds is replaced by the chain and the measurement. Reading the injector surfaced three real differences, so the marker set is now exactly upstream's. `quantization.quant_method` is no longer accepted, because upstream writes and reads the marker at the top level and never looks inside -- a document whose only marker was nested resolved fp8 KV here and nowhere else. A nested `modelopt_quant_config` key is now accepted (`:218-220`), which upstream treats as the legacy modelopt marker and this port ignored. And the producer name is compared raw against the literal, because `:222` is a bare `==`, while `quant_method` keeps the prefix match and the case fold that `:238-246` and `:319` between them perform. G1's marker case pins all six arms. One difference is recorded rather than copied. The injector RAISES `ValueError: Unknown ModelOpt quant algo: <algo>` (`:235`) when the producer is modelopt and the nested `quant_algo` is neither FP8-family nor NVFP4, and this resolver answers `fp8_e4m3` instead. That refusal is a weight-half validation living in a config convertor this port does not have, and moving it into the KV resolver would refuse a `MIXED_PRECISION` checkpoint whose weights `modelopt_mixed_precision.h` loads. Porting the convertor's validation is its own row; the spec's `## Owed` names it. Four record corrections ride along. The #1574 subject does NOT take the declared-fp8 path: verified from the live artifact @ `36f717a2`, its `config.json:quantization_config` carries `quant_method: "modelopt"`, `quant_algo: "MIXED_PRECISION"` and no `kv_cache_*` key at all, and only the legacy `hf_quant_config.json` declares `kv_cache_quant_algo: "FP8"`. The inline document wins, so neither engine auto-selects fp8 KV for it and the campaign has to type the flag on both sides -- which its competitors' `serve.sh` already does. Five places said the opposite and a new G10 case pins the answer at `auto` with both real documents in one directory. The refusal at the other architectures was described as arriving at `vt::ReshapeAndCache`, and for 13 of 16 it does not: `granite`, `minicpm`, `phi3`, `gemma3`, `opt`, `stablelm`, `glm4`, `commandr`, `gemma`, `gemma2`, `phi`, `muse_glimmer` and `olmo2` each refuse at their own `"<arch>: KV cache must be bf16 or f32"` guard first, which names neither fp8 nor the flag. The count was 16 architectures at 17 call sites rather than 17 architectures. And `docs/FEATURES.md` promised "halves the block, doubles the pool" unqualified, where `ResolveNumBlocks` returns the fixed 256-block fallback unless `--num-blocks` or `--kv-cache-memory` is given -- so on the default path fp8 KV halves the pool bytes instead. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…8x scale walk through 30/30 green (#1593) G12's counter-case asserted `max_abs < 1.0` on a delta its own MESSAGE reported as 3.4e-4, so it measured this toy model's insensitivity to its KV cache rather than the cache. The third review walked two mutations of `kv_cache_route.h:63` straight through the suite: storing V into `k_cache` and K into `v_cache` (delta 0.0247, 73x) and storing with `k_scale * 8` / `v_scale * 8` against an unscaled read (delta 0.0064, 19x). Neither is visible on that axis, and no number written on it would have been safe. The new G12 case compares the CACHE BYTES of layer 0 against the bf16 run's page inside the envelope the FORMAT defines: e4m3fn carries three explicit mantissa bits and rounds to nearest even, so a normal is within 2^-4 of its value and a subnormal within 2^-10 of the scale. Layer 0 is the whole population because its K and V are functions of the embedding and the input layernorm alone, so both runs hand the store bit-identical floats. Scales are 0.125 and 0.25, non-unit and unequal, so a dropped, swapped or one-sided scale leaves the envelope. Measured: 0/320 elements outside, 281/320 normals over 10/10 pages. Both mutations are red against it (316/320 outside at ratio 416.9; 319/320 at ratio 13.9), and both left G5 and G9 green, which is why an engine-level token comparison is not the repair. The `## Owed` refusal accounting was wrong for two of the three architectures it named. `gemma4:306-315` never reaches the store: with a kI8 page it casts into a kI8 destination and dies in `vt::CastF32` (`ops.cpp:4087`), naming neither fp8, nor the flag, nor gemma4. `nemotron_h_device:1589-1593` has its own `VT_CHECK` that names the fp8 KV scheme and fires first. Only `qwen3_vl:198-200` reaches the store guard, so the split is 14 / 1 / 1 rather than 13 / 3. Safety is unaffected; message quality is what moves. G7 hand-builds its tensors, so it gates the store guard's message and never the claim that any architecture reaches it. The marker mirror over-accepted three shapes. Upstream injects `quant_method` only `if quant_algo is not None` (`:224`), reading that key out of `quant_cfg.get("quantization", {})` -- an empty-object fallback, unlike the reader's `quant_cfg.get("quantization", quant_cfg)` at `torch_utils.py:321`. A modelopt producer with no `quant_algo`, a `quantization.modelopt_quant_config` with no `quant_algo`, and a top-level `modelopt_quant_config` with no `quantization` key each answered `fp8_e4m3` here while upstream answered `None`. All three are now refused and G1 arm (g) pins them. The MIXED_PRECISION raise stays owed: importing it would refuse a checkpoint this tree loads. This PR gave `hf_quant_config.json` its first production reader (`cache.cpp:207` from `model_loader.cpp:1988`) and left four statements elsewhere false. `modelopt_mixed_precision.h`, both `qwen38-27b-quant-arms.md` bullets and the `test_qwen38_27b_modelopt_mtp_arm` case name are narrowed to what is now true: the file is read, but only as the legacy fallback behind `config.json`'s inline `quantization_config`, and both artifacts ship that inline document. No assertion changed. `KvAlgoFromObject`'s `std::optional<int>` ternary is a statement now. GCC 13 at -O2 could not see through the inlined form and failed `-Werror=maybe-uninitialized`, which broke every Release and RelWithDebInfo build of this branch. CI configures without a build type and never saw it. CUDA remains unexecuted: no device in this session. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
Merge up so the diff-scoped `commit-trailers` and `commit-style` gates run against an ancestor rather than SKIPPING, and so #1649's fix `2f2a70925` -- which was NOT reachable from this branch's base -- is reachable here. The two `windows-msvc-*` reds this pull request has been carrying are that issue, and this is what settles whether they were the repo-wide failure or something of ours. No conflicts, and `.agents/issue-index.md` is unchanged by this branch. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
`origin/main` moved again while the round-3 gate was running, and the diff-scoped `commit-trailers` and `commit-style` checks report nothing at all about a branch their base is not an ancestor of. Merge up so both execute rather than SKIP. No conflicts. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…at carried 13 / 3 (#1593) The recount landed in the spec and stopped there, and four surfaces this same pull request adds still told a reader that Gemma-4, Qwen3-VL and Nemotron-H are the three that name the flag. None of them names the flag. Only Nemotron-H names fp8, only Qwen3-VL reaches the store guard, and Gemma-4 dies inside `vt::CastF32` with a message that names no architecture at all. `docs/USAGE.md` had the strongest version of the claim and now describes what an operator actually sees on each of the sixteen. `docs/FEATURES.md`'s cell says one architecture names fp8 back rather than implying all sixteen refuse by their own name. The engine-matrix and quantization-matrix rows carry the same correction, and both take the new gate count (31 cases, G1-G12) with one line on what G12's third case bounds -- the quantization-matrix row was still saying 26 cases and G1-G11, and its residual list still said 17 unrouted architectures. Records only. No product code and no assertion moves. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
The fourth review found `scripts/agent-integration.py --base origin/main` exiting 1 with commit-trailers and commit-style both SKIPPED: origin/main had moved to c98ffd4 while this branch's last merge took ff8f728, so neither trailer gate reported anything about this tree. Merging main up puts origin/main back in this head's ancestry and both gates execute again. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…d V with K's scale walked 31/31 green (#1593) The fourth review found the hole the third round's repair left. G12's envelope case decodes the cache bytes with its own `vt::LoadKvFp8E4M3` and never enters the production dequant at `cpu_paged_attn.cpp:167`, and `fp8_logits` -- the one value in the case that IS downstream of that dequant -- carried no assertion at all. Every other case in the file that asserts a number downstream of the read runs at `k_scale == v_scale == 1`, where a k/v scale swap on the read is arithmetically inert. Mutating the production read to `const float v_scale = args.k_scale;` therefore left `test_kv_cache_fp8_wiring` at 31/31 and `test_ops_fp8_kv_cache` at 8/8, both SUCCESS, while every V the softmax saw was halved. The spec and the pull request body claimed a dropped, swapped or one-sided scale leaves the envelope; that was true of the store and false of the read, and both now say so. The read is closed by SCALE INVARIANCE, and it is exact rather than a tolerance. e4m3fn's normal grid is relative -- for |y| in [2^e, 2^(e+1)) the representable points are m*2^(e-3) -- and dividing by a power of two is exact in binary floating point and shifts e without touching the mantissa. So for two power-of- two scales that both leave a value normal and unsaturated, `s * Dequant(Quantize(x/s))` is the same float bit for bit. Two fp8 runs at different power-of-two scales must produce BIT-IDENTICAL logits: their cache bytes differ in every element's exponent field and the floats the attention kernel is handed do not. The case runs the seam twice more, at (2^-7, 2^-13) and (2^-11, 2^-9), and requires memcmp-level agreement. Both pairs move BOTH sides, because a pair that moved only one would let a read-side defect that depends on the other scale reproduce itself identically in both runs and cancel out. Two anti-vacuity REQUIREs hold it up: all 320/320 elements are normal and unsaturated at all four scales, since one subnormal would round on the absolute 2^-9 grid which is not scale invariant; and the two caches really do hold different bytes, so the comparison is about the read and not about two identical buffers. MEASURED on the repaired tree: 0/320 logits differ, max |delta| exactly 0. Three read-side mutations are RED, and the store envelope reads 0/320 under all three -- 320/320 at max delta 0.0673 for `v_scale = args.k_scale`, 256/320 at 1.08e-4 for `k_scale = args.v_scale`, and 320/320 at 2.47e-3 for `v_scale = 1.0F`. The K-side one is why this is stated as exact equality: a tolerance sized at 1e-4 would have let it through. The existing envelope assertions and the case's structure are untouched; this adds to them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
Merged so that `commit-trailers` and `commit-style` execute against this tree rather than SKIP on a base that has moved. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…tation, and the exactness holds 320 of the 640 elements (#1593) The fifth review passed the code and rejected two sentences written about it. No production code changes here, and no assertion moves; the pull request body becomes the landed commit message under `squash_merge_commit_message = PR_BODY`, so both claims are repaired before the merge rather than after it. Every number below was re-measured on this tree. "A tolerance sized at `1e-4` would have let it through" is FALSE on either reading. The K-side mutation's measured delta is `1.08408e-4`, which is 8.4% ABOVE an absolute `1e-4`, so that bound catches it; and a relative `1e-4` taken against the invariance runs' largest logit -- measured `0.0627671` on both -- is `6.3e-6` absolute and catches it by 17x. The argument that survives is the SPAN: the three read signals run `1.08e-4` to `0.0673`, a factor of 621, so a bound sized against the largest keeps nothing for the smallest. Even a tight 12.4x margin against `0.0673` puts the constant at `5.4e-3`, which admits `1.08e-4` by 50x. That is why the gate is exact equality and not a number. The exactness claim was also broader than what is asserted. The comment said the gate is exact "as a property of the format rather than a lucky measurement" and that "the `scale_exact` REQUIRE below holds the population to it". `scale_exact` decodes `bf16.buf[0]`; `MakeSeamConfig` sets `num_hidden_layers = 2`, so it holds 320 of the 640 elements each run stores, while `inv_differing` compares LOGITS that are a function of both layers' caches. Extending the assertion over every layer was the preferred repair and this fixture cannot carry it: decoding each run's own `buf[1]` at its own scales measures `3/320` elements disagreeing by up to `7.62939e-06`, all K-side and an order of magnitude below layer 0's `1.76e-4` minimum, so at `kInvAKScale = 2^-7` they sit in e4m3's SUBNORMAL region where the grid is the absolute `2^-9` step and the power-of-two covariance does not hold. That layer carries 2 saturated elements besides. So the prose is narrowed instead: exact-by-format for layer 0, EMPIRICAL for layer 1, where a 7.6e-6 cache perturbation is absorbed in f32 accumulation before it reaches a logit. It is not knife-edge -- a third legal pair `(2^-9, 2^-11)` against pair A also measures `0/320` logits differing at max `|delta|` `0` -- and the fragility is recorded under `## Owed`, because a change to `MakeSeamWeights`, `kSeamTokens`, `num_hidden_layers`, the thread count or the accumulation order could push a layer-1 discrepancy into a logit and redden a CORRECT tree. A second `## Owed` line records what the gate still admits. An fp8-versus-fp8 comparison cannot see a read defect that is a function of BYTES and INDICES rather than of scales, because both runs commit it identically and it cancels. Measured on production code, both of these PASS the whole file at 31/31, 487/487 with `0/320 logits differ`: V served out of the K page with K's scale (`cpu_paged_attn.cpp:174` `v_base = k_cache.data` with `:167` `v_scale = args.k_scale`), the read-side twin of the `N1_KVSWAP` store mutation the envelope DOES catch; and the V read dropping the in-page token offset (`:270`, `off & 0`), which is pure indexing and scale-free. Closing that class needs a comparison against a REFERENCE and never a second fp8 run, so the design paragraph now claims the SCALE half of the read alone. `test_kv_cache_fp8_wiring` stays 31/31, 487/487 SUCCESS on this tree. Both mutation sites were restored byte-for-byte against a pre-taken md5. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
W1 landed the CPU fp8-e4m3 KV store and read, W2 the CUDA arm, and both landed
with nothing reaching either from a production entry point. W3 is the wiring:
--kv-cache-dtypeon the server flag, the checkpoint's ownkv_cache_quant_algohonoured when no flag is typed, KV blocks sized at one byte per element, and the
k_scale/v_scalepath.Now based on
main. W2 merged as squash53f5b74be, so this branch'sprevious base
row/KV-FP8-W2is dead and GitHub was computing mergeabilityagainst it.
origin/mainis merged in and the conflict set was re-derived here,not inherited.
Issue: #1593 (the issue that carries W2
and W3). Spec:
.agents/specs/fp8-kv-cache.md## W3 — the runner integration.What the FIRST fresh review found, and what changed
The review of
68d2e0d5areturned FAIL — not safe to merge, with onefunctional bug. Every finding is repaired here, and the second review re-proved
every one of them by mutation.
The bug:
--kv-cache-dtype fp8threw at the first forward on every real checkpointFullAttnBlockPagedcast K and V to the cache dtype only underif (kv.dtype == DType::kBF16), and on the fp8 routekv.dtypeiskI8. So Kstill arrived as f32 —
attn_dtis f32 for every fp8 cache, becausekv.dtype == DType::kBF16is a term of both FA2 eligibility tests — while Varrived as whatever the
v_projGEMM emitted. That is bf16 on the block-wisefp8 arm (
MatmulFp8BlockScaledD), bf16 on the NVFP4 arm under the defaultVT_BF16_GEMM_OUT, and bf16 on ordinary torch safetensors (MatmulBf16D).vt::ReshapeAndCacheFp8quantizes from ONE source dtype, so the flag threwreshape_and_cache_fp8: k/v must share one float dtypeat the first forward onevery weight arm a user has — including the #1574 campaign checkpoint this
wave exists to serve, so W3 as written did not deliver the matched comparison
it is on the critical path for.
The gate could not see it.
MakeDenseWeightsbuilds projection weights with nonkflag,ProjectFullAttnQkvserves those throughMatmulF32D, and f32 pairswith f32. Only the per-tensor fp8 arm and that synthetic arm ever paired.
The fp8 route now normalises both to bf16, exactly as the bf16-cache route
already did, reusing the two
DBufs that were already allocated. bf16 ratherthan f32 because bf16 is the dtype upstream quantizes from: vLLM's model IS bf16
where
reshape_and_cache_flashtakes key/value(
cache_kernels.cu:314-401). G9 is the new case, and it enters throughLoadedEnginewith the raw torch Linear layout ([N=out, K=in],nk) a realsafetensors checkpoint ships.
The merge conflict in
model_loader.cppwas SEMANTIC, and both sides survivemain's SPEC-DRAFTER-CHAIN W1 refusal (#1522) and this row's resolution stanzaboth insert as the FIRST statement of
FromModelDir, and both are load-bearing:that row's G5 requires the chain refusal before any weight I/O, and
ReadQuantConfigJsonopens a file insidemodel_dir. Taking either side drops aguarantee. Both survive, chain refusal first, and the order is gated.
tests/vllm/entrypoints/test_drafter_chain_reach.cppcannot catch an inversion:it points at a NONEXISTENT directory, and
ReadQuantConfigJsonanswers""forone of those without opening anything, so the chain refusal arrives either way.
G10's ordering case points at a directory that EXISTS and declares fp8, so an
inverted order announces the declaration first — and that line is the evidence.
Mutation M19b MOVES the refusal below the stanza (rather than deleting it)
and reddens exactly that assertion.
The keyed records —
.agents/engine-matrix.md,.agents/quantization-matrix.md,.agents/issue-index.md— take the target-branch version with the scoped editre-applied by hand. The union merge had duplicated two
issue-indexrows(#1593 and #1636 appear once on
mainand twice after the automatic merge);that is the known union-driver trap and it is undone here.
Three surfaces that no gate reached
FromModelDir's resolution stanza +ReadQuantConfigJson, whose only caller it isLoadedEngine::FromModelDiron a real directory--kv-cache-dtypeparse arm andengine_params.kv_cache_dtype = args.kv_cache_dtypetests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp, driving the REALVllmServerMainApplyCacheDType'sper_layer_attn_specsloopKVBytesPerBlockexactly halvedThe serve gate works on a polarity rather than on a printed value: the loader
announces the checkpoint's declaration only when nothing explicit was typed, so
CASE 1 (no flag) requires the line and CASE 2 (
--kv-cache-dtype bfloat16)requires its absence. Break either half of the flag's chain and the default
autoarrives, the checkpoint wins, and the line comes back.The file precedence was inverted against upstream
ReadQuantConfigJsonreadhf_quant_config.jsonfirst and fell back toconfig.json:quantization_config.vllm/transformers_utils/config.py:751-761does the reverse, under upstream's own comments: ModelOpt writes the inline
document from 0.31.0 on, and the standalone file is what 0.29.0 and before wrote,
consulted only when the inline one is
None. Mirrored, and gated at both theresolver and the loader by G10 — a checkpoint re-quantized in place, inline
document added and stale legacy file left beside it, otherwise resolves to the
OLD declaration and quantizes a KV cache nobody asked to quantize, at half the
page, silently.
What turning it on costs, recorded rather than claimed
An fp8 KV cache takes the model off every fast attention kernel this engine
has.
qwen3_5.cppmakeskv.dtype == DType::kBF16a term of bothfa2_prefillandfa2_decode;src/vt/cuda/cuda_paged_attn.cu:2941-2960routesan fp8 read only through tiled prefill and block decode, because the WMMA ladder,
the vendored FA-2 launchers and the vectorized decode-opt/GQA kernels are
bf16-native by construction.
W3's rationale is that fp8 KV makes campaign #1574 a matched comparison, where we
currently lead 11.06 vs 9.71 tok/s on the bf16 path. Turning it on plausibly
LOSES that lead, and this session had no GPU and measured nothing. So the spec's
## W3,docs/USAGE.md,docs/FEATURES.mdand both matrix rows now say theexchange rate is unmeasured, and the fast-kernel routing is named under
## Owed. No direction is claimed.What the SECOND fresh review found, and what changed
The review of
947c5b65freturned FAIL, with one blocking finding. Everyfinding is repaired here, and the first round's repairs are untouched.
BLOCKING: the shared seam's fp8 routing was DEAD CODE
dense_attn::AttnBlockis named in AGENTS.md## Shared seamsas the decodeseam, and W3 listed it as routed. It was not. The routing was written —
fp8_kv,WriteKvCache,ApplyKvCacheQuant— underneath a preamble guard that stilladmitted only
kBF16andkF32, andIsFp8KvCacheis true only forkI8. Sofp8_kvwas provably false at every call and neither fp8 arm could execute. Notan inference: a contradiction between two lines of the same function.
The gate was blind to it. The reviewer reverted the ENTIRE routing from the
seam and got 26 cases / 120 assertions SUCCESS, because every case in the file
entered through
Qwen3_5DenseModel::Forwardand none entered here.The seam is the production forward for Qwen3 dense (
qwen3.cpp:185), Qwen3-MoE(
qwen3_moe.cpp:84), Voxtral (voxtral.cpp:102) and the Llama, Mistral andInternLM2 registries that share
Qwen3DenseModel.--kv-cache-dtype fp8on anyof them sized the pool at half and then threw a message naming neither fp8, nor
the flag, nor the row.
The guard is widened exactly as
qwen3_5.cpp:5313already was, and G12 isthe new gate. It enters through
Qwen3DenseModel::Forward, which is whatForwardQwen3ForCausalLMcalls underModelRegistry::Forward(
qwen3_dense.cpp:113), and it asserts three things a "it ran" case would not:the half-width pages carry bytes after the forward, two fresh stacks agree bit
for bit, and the fp8 logits DIFFER from the same forward over a bf16 cache
(320/320, max |delta| 3.39e-4) — which is what says the read really dequantized
rather than the cache behaving like a float one. Mutation MSEAM reverts the
seam's routing the way the reviewer did and reddens both cases.
The #1574 checkpoint does NOT take the declared-fp8 path, and five places said it did
Verified from the live artifact @
36f717a2:config.json:quantization_configcarries
quant_method: "modelopt",quant_algo: "MIXED_PRECISION"and nokv_cache_*key at all; onlyhf_quant_config.jsondeclareskv_cache_quant_algo: "FP8". The first round's own file-precedence repair makesthe inline document win, so the declaration is never read for this checkpoint —
on this engine and on vLLM alike.
Nothing is functionally wrong; the mirror is correct on both halves. What was
wrong is what the spec,
kv_cache.h,cache.h,model_loader.cppand this bodysaid about it. The campaign consequence, stated plainly: for #1574 the flag has
to be typed explicitly on both sides. The competitors' own
serve.shalreadypasses
--kv-cache-dtype fp8, so this is consistent rather than a blocker. A newG10 case writes BOTH real documents into one directory and pins the answer at
auto.The modelopt marker: NOT a divergence, and the reason took running upstream
The review reported
src/vllm/config/cache.cppacceptingproducer.nameas adivergence, on the grounds that
get_kv_cache_quant_algo_string(
torch_utils.py:319) gates on a top-levelquant_methodand nothing else. Thattranscription is of one function; the answer needs two.
ModelArchConfigConvertorBase._normalize_quantization_config(
transformers_utils/model_arch_config_convertor.py:208-247) runs FIRST, atModelConfig.__post_init__(config/model.py:577), and INJECTSquant_cfg["quant_method"] = "modelopt"whenproducer["name"] == "modelopt",mutating the very dict
hf_config.quantization_confignames.MEASURED by extracting both functions with
astfrom the files at555967922and running them over the live documents rather than retyping them:
nvidia/Llama-3.3-70B-Instruct-FP8's producer-onlyhf_quant_config.jsonanswers
Nonebefore normalization and'fp8_e4m3'after it. So accepting theproducer name is the MIRROR. The comment that claimed it on the wrong grounds
("
modelopt_mixed_precision.halready reads both shapes") is replaced by thechain and the measurement.
Reading the injector did surface three real differences, and the marker set is
now exactly upstream's:
quantization.quant_methodis no longer accepted. Upstream writes and readsthe marker at the TOP level and never looks inside, so a document whose only
marker is nested resolved fp8 KV here and nowhere else.
modelopt_quant_configkey is now accepted (:218-220), whichupstream treats as the legacy modelopt marker and we ignored.
:222is==), notprefix-matched and case-folded.
quant_methodkeeps both, because:238-246lower-cases it before
:319prefix-tests it. Same file, two tests, mirroredseparately.
G1's new marker case pins all six arms; restoring the nested-marker widening
reddens two of them.
One difference remains and is recorded rather than copied. The injector
RAISES
ValueError: Unknown ModelOpt quant algo: <algo>(:235) when theproducer is modelopt and the nested
quant_algois neither FP8-family norNVFP4 — which is what the #1574 legacy file (
MIXED_PRECISION) gets, and weanswer
fp8_e4m3for it instead. That refusal is a WEIGHT-half validationliving in a config convertor this port does not have, and moving it into the KV
resolver would refuse a checkpoint whose weights
modelopt_mixed_precision.hloads. It is unreachable for that checkpoint anyway, because its inline document
wins. Spec
## Owednames it.The refusal mechanism was wrong for 13 of the 16 unrouted architectures
"Every other architecture is refused BY NAME at
vt::ReshapeAndCache" holds forthree.
granite:95,minicpm:96,phi3:78,gemma3:121,opt:125,stablelm:86,glm4:93,commandr:93,gemma:53,gemma2:135,phi:98,muse_glimmer:144andolmo2:94each refuse at their own"<arch>: KV cache must be bf16 or f32"guard first, which names neither fp8nor the flag. Only
gemma4,qwen3_vlandnemotron_h_devicereach the storeguard. The safety property holds — nothing writes floats into a half-sized page
— but the described mechanism did not, and the count was 16 architectures at 17
call sites rather than 17 architectures. Round 3 traced each of the three and
found only
qwen3_vlreaches the store guard; the split is 14 / 1 / 1. Seebelow. Corrected in the spec, in
docs/USAGE.md, indocs/FEATURES.mdand in the engine-matrix row, and the13-guard gap is now its own
## Owedbullet.docs/FEATURES.mdpromised a doubled pool on a path that halves itResolveNumBlocks(model_loader.cpp:1446-1451) returns the fixed 256-blockfallback when neither
--num-blocksnor--kv-cache-memoryis given, so on thedefault path fp8 KV halves the pool BYTES rather than doubling the pool. The
one-liner said "halves the block, doubles the pool" unqualified;
docs/USAGE.md's example passes--kv-cache-memoryand was correct, and nowsays so explicitly.
Preflight was reporting nothing about two gates
origin/mainhad moved to08c81a89and was no longer an ancestor, so thediff-scoped
commit-trailersandcommit-stylechecks SKIPPED.origin/mainismerged up and both execute.
What the THIRD fresh review found, and what changed
The review of
d0bcf80ddreturned FAIL, with one blocking finding. Everyfinding is repaired here, and rounds 1 and 2 are untouched. The round-2
producer.nameMEDIUM was FALSIFIED by that review, whichast-extracted bothpinned upstream functions and ran them:
_normalize_quantization_configmutateshf_config.quantization_configin place duringModelConfig.__post_init__,before
resolve_kv_cache_dtype_stringreads it, so acceptingproducer.nameisthe mirror. It stays.
BLOCKING: the gate could not distinguish a working fp8 KV cache from a corrupted one
G12's counter-case asserted
differing > 0andmax_abs < 1.0on a delta itsown
MESSAGEreported as0.000339303. The bound sat about 3000x above thesignal, so it measured this toy model's insensitivity to its KV cache and not
the cache. Two mutations of
kv_cache_route.h:63walked through the whole 30/30suite:
(q, v, k, k_cache, v_cache, ...), so the fp8 arm stores V intok_cacheand K intov_cache0.0247495(73x)kv.k_scale * 8.0F/kv.v_scale * 8.0Fwhile the read uses the unscaled values0.0063844(19x)The second one is exactly the defect the
k_scale/v_scaleplumbing exists toprevent. Nothing above G12 closed the hole either: G5 and G9, the two
production-entry cases, assert token counts and determinism across two fresh
stacks, and neither compares an fp8 run against a bf16-cache run.
The repair compares the CACHE BYTES, inside an envelope read off the FORMAT.
A new G12 case runs the same prompt twice — once with a bf16 page, once with an
fp8 page — and checks every layer-0 element of both K and V:
e4m3fn carries three explicit mantissa bits and rounds to nearest even
(
include/vt/fp8_kv.h), so a normal magnitude is within a half ulp — at most2^-4of its own value — and anything below the smallest normal2^-6sits ona uniform
2^-9grid whose half step is2^-10, in absolute terms scaled by theper-tensor scale. Not one constant is fitted to a measurement, which is the point:
a threshold sized to today's
3.4e-4is the same hole one decimal place tighter.Layer 0 is the whole population, and that is the design rather than a
limitation. Its K and V are functions of the embedding and the input layernorm
alone, so the bf16 run and the fp8 run hand the store bit-identical floats and
the float run's page IS the reference the fp8 page has to round. From layer 1 on
the fp8 run's inputs already carry the previous layer's dequantization and no
per-element envelope holds.
The scales are
0.125for K and0.25for V — non-unit and unequal, so adropped, swapped or one-sided scale IN THE STORE leaves the envelope rather than
staying inside it. Both are below one because the store divides by the scale and this
model's layer-0 K and V land around
1e-2; a scale above one pushes mostelements under
2^-6, where only the weaker absolute arm applies. Two anti-vacuityREQUIREs hold the population to that: every one of the ten pages (five tokens,K and V) must carry at least one NORMAL, and no reference may reach e4m3's finite
maximum, so the bound is the pure rounding envelope and never a clamp.
MEASURED on the repaired tree, and both mutations are RED against it:
A bf16-versus-fp8 comparison in G5/G9 through
LoadedEnginewas considered andis NOT what closes this. Both STORE mutations left every other case green, G5 and
G9 included, which is the direct measurement that an engine-level token or
determinism comparison cannot see a dequant defect this model absorbs. The old
two assertions are kept beside the new case, not replaced.
Round 4: that envelope gated the STORE, and the READ was still open
The fourth review found the hole the third one's repair left. The envelope case
decodes the cache bytes with its own
vt::LoadKvFp8E4M3and never enters theproduction dequant (
cpu_paged_attn.cpp:167); the one value in the case that ISdownstream of that dequant carried no assertion at all. Every other case in the
file that asserts a number downstream of the read runs at
k_scale == v_scale == 1, where a k/v scale SWAP on the read is arithmeticallyinert. Mutating the production read to
const float v_scale = args.k_scale;therefore left
test_kv_cache_fp8_wiringat 31/31 andtest_ops_fp8_kv_cacheat8/8, both SUCCESS, while every V the softmax saw was halved.
The read's SCALE is closed by INVARIANCE, and EXACTLY rather than by a
tolerance — its ROUTING is not, and that half is under
## Owed.e4m3fn's normal grid is relative — for
|y|in[2^e, 2^(e+1))therepresentable points are
m * 2^(e-3)— and dividing by a power of two is exactin binary floating point, shifting
ewithout touching the mantissa. So for anytwo power-of-two scales that both leave a value normal and unsaturated,
bit for bit: the two runs' cache BYTES differ in every element's exponent field,
and the floats the attention kernel is handed do not. The case runs the seam
twice more, at
(2^-7, 2^-13)and(2^-11, 2^-9), and requiresmemcmp-levelagreement on the logits. There is no constant to fit and none to widen later,
which is the same discipline the envelope above is written to. Both pairs move
BOTH sides on purpose: a pair that moved only one would let a read-side defect
that depends on the other scale reproduce itself identically in both runs and
cancel out.
That exactness is ASSERTED for LAYER 0, and the logit equality is an EMPIRICAL
result for this fixture rather than a theorem. The
scale_exactREQUIREdecodes
bf16.buf[0]and holds320/320LAYER-0 elements normal and unsaturatedat all four scales — the measured magnitudes are
1.76e-4to1.32e-1for K and5.41e-5to4.22e-2for V, against all-normal windows of(2.94e-4, 1.13e-2]and(9.43e-5, 3.46e-3]— because one subnormal would roundon the ABSOLUTE
2^-9grid, which is not scale invariant.MakeSeamConfigsetsnum_hidden_layers = 2, so that is 320 of the 640 elements each run stores,while the logits the case compares are a function of BOTH layers' caches. Layer 1
does NOT satisfy the precondition: decoding each run's own
buf[1]at its ownscales,
3/320of its elements disagree by up to7.62939e-06, all K-side andan order of magnitude below layer 0's
1.76e-4minimum, so atkInvAKScale = 2^-7they land in e4m3's SUBNORMAL region and that layer carries2 SATURATED elements besides.
CHECK(inv_differing == 0)therefore holds by theformat property for layer 0 and by ABSORPTION for layer 1 — a
7.6e-6cacheperturbation vanishing in f32 accumulation before it reaches a logit. Extending
scale_exactover every layer was the preferred repair and this fixture cannotsatisfy it, so the claim is narrowed here and the fragility is recorded under
## Owed.The second anti-vacuity
REQUIREis that the two caches really do hold differentbytes, so the comparison is a statement about the read rather than about two
identical buffers.
MEASURED on the repaired tree:
0/320logits differ, max|delta|exactly0.Three read-side mutations of
cpu_paged_attn.cpp:167are RED, and the storeenvelope above reads
0/320under all three:v_scale = args.k_scale0.0673k_scale = args.v_scale1.08e-4v_scale = 1.0F2.47e-3The K-side one is why this is stated as exact equality rather than a bound. The
three signals span 621x, so a bound sized against the largest keeps nothing
for the smallest: even a tight 12.4x margin against
0.0673puts the constant at5.4e-3, which admits1.08e-4by 50x. An absolute1e-4would in facthave CAUGHT the K-side one, by 8.4%, and a relative
1e-4against the fixture'slargest logit (
0.0627671) is6.3e-6absolute and catches it by 17x — theargument is the span, and the refusal to fit a constant to whichever defect
happened to be measured first, never any one number.
What this comparison still cannot see is a read-side defect that is a
function of BYTES and INDICES rather than of scales, because both fp8 runs commit
it identically and it cancels. Measured on production code, both of these PASS
the whole file at
31/31,487/487, with0/320 logits differ: serving V outof the K page with K's scale (
cpu_paged_attn.cpp:174v_base = k_cache.datawith
:167v_scale = args.k_scale), the read-side twin of theN1_KVSWAPstore mutation the envelope DOES catch; and dropping the in-page token offset
from the V read (
:270,off & 0), which is pure indexing and scale-free.Closing that class needs a comparison against a REFERENCE — the bf16 run's page,
or the case's own decode — and never a second fp8 run. It is under
## Owed.The same review found
scripts/agent-integration.py --base origin/mainexiting 1with
commit-trailersandcommit-styleboth SKIPPED:origin/mainhad moved toc98ffd4d0while this branch's last merge tookff8f72807, so neither gatereported anything about this tree.
mainis merged up and both execute again.Round 5: the code was correct, and two sentences about it were not
The fifth review passed the change and rejected two claims written around it. It
did not ask for a production edit and none was made; the two paragraphs above are
where the repair lands, and both numbers in them were re-measured here rather
than carried over.
origin/mainhad moved again, toea9b7e30e, so it ismerged up once more and
commit-trailersandcommit-styleexecute againstthis tree rather than SKIP on a base that has moved.
The first was the tolerance sentence. "A tolerance sized at
1e-4would have letit through" is false as stated and false on either reading:
1.08408e-4is 8.4%ABOVE an absolute
1e-4, and a relative1e-4against the fixture's largestlogit
0.0627671is6.3e-6and catches it by 17x. The argument that survivesis the 621x SPAN between the three read signals, not that constant.
The second was the exactness scope. The comment said the invariance gate is exact
"as a property of the format rather than a lucky measurement" and that
"the
scale_exactREQUIRE below holds the population to it".scale_exactdecodes
bf16.buf[0], which is 320 of the 640 elements each run stores, and thelogits it compares depend on both layers. Extending the assertion over every
layer was the preferred repair and the fixture cannot carry it: layer 1 measures
3/320elements disagreeing by up to7.62939e-06, subnormal at2^-7. So theprose is narrowed to what is asserted, the empirical part is named as empirical,
and the fragility and the unclosed routing half are both under
## Owed.The
## Owedrefusal accounting was wrong for 2 of the 3 architectures it namedRound 2 said
gemma4,qwen3_vlandnemotron_h_devicereach the store guard.Traced with a
kI8page against a bf16 model dtype, onlyqwen3_vldoes. Thetrue split is 14 / 1 / 1:
gemma4.cpp:306-315takeskv.dtype != adt, allocatesDBuf kcast(d, kv.dtype /* kI8 */, ...)and callsvt::CastF32, which refusesat
src/vt/ops.cpp:4087with"cast_f32: out must be f32". It never reachesthe store, and that message names neither fp8, nor the flag, nor
gemma4— soit is the WORST of the sixteen, not one of the three best.
nemotron_h_device.cpp:1589-1593carries an explicitelse { VT_CHECK(false, "NemotronH paged forward: ... The fp8 KV scheme ... is not selected here"); }on the same cast, which fires first. It is the ONErefusal in the sixteen that names fp8 to the operator.
qwen3_vl.cpp:198-200carries no guard and no cast, sovt::ReshapeAndCacheis what refuses it.
Safety is unaffected — nothing writes floats into a half-sized page on any of the
sixteen — and what moves is message quality: the store guard's better message is
reached by 1 of 16, not 3. Corrected in the spec's narrative and in its
## Owedbullet, which now also records that G7 hand-builds its K/V tensors and calls
vt::ReshapeAndCachedirectly, so it gates the store guard's MESSAGE and neverthe claim that any particular architecture reaches it.
The marker mirror accepted three shapes upstream answers
Nonefor_normalize_quantization_configinjectsquant_methodonlyif quant_algo is not None(:224), and it reads that key out ofquant_cfg.get("quantization", {})— an EMPTY-object fallback, unlike thereader's
quant_cfg.get("quantization", quant_cfg)attorch_utils.py:321.cache.cppaccepted the producer name and the legacy key unconditionally, andread the legacy key out of the reader's
inner. Three documents thereforeresolved to
fp8_e4m3here while upstream resolvedNone:producer.name == "modelopt",quantizationwith noquant_algoNonefp8_e4m3autoquantization.modelopt_quant_configwith noquant_algoNonefp8_e4m3automodelopt_quant_config, noquantizationkey at allNonefp8_e4m3autoEach of the three would halve the KV page on a checkpoint vLLM runs at the model
dtype. Arm (g) of G1's marker case pins all three, alongside the same
documents made acceptable by adding the
quant_algoupstream requires, so whatrefuses them is the guard and not the shape. No shipped fixture moved: every
real document in the suite carries a
quant_algo.The MIXED_PRECISION raise stays under
## Owed, and that is why the mirroredtest is exactly
quant_algo is not Nonerather than the FP8/NVFP4 family set:the only two upstream outcomes for a
quant_algothat IS present are "inject"and "raise", and taking the raise's arm collapses them into one. Importing the
raise would refuse a checkpoint whose weights this tree loads.
This PR falsified four statements elsewhere and now carries the repair
Before this branch nothing in production read
hf_quant_config.json.src/vllm/config/cache.cpp:207now does, fromLoadedEngine::FromModelDir(
model_loader.cpp:1988). Four statements written against the old world readfalse at this head, and a record edit rides in the pull request whose change made
it stale:
modelopt_mixed_precision.h:1006config.json's inline document, and this checkpoint ships that document.agents/specs/qwen38-27b-quant-arms.md:1629.agents/specs/qwen38-27b-quant-arms.md:1841hf_quant_config.json"test_qwen38_27b_modelopt_mtp_arm.cpp:482(case NAME) and:487-488config.jsonoutranks", with the precedence spelled outNo assertion changed. They stay valid because
config.json's inlinequantization_configwins on the one path that reads either.The 14 / 1 / 1 split had to reach four more records
The recount above landed in the spec and stopped there.
docs/USAGE.md,docs/FEATURES.md, theKV-FP8engine-matrix row and theQUANT-KV-FP8quantization-matrix row — all four written by THIS pull request — still told a
reader that Gemma-4, Qwen3-VL and Nemotron-H are the three that name the flag.
None of them names the flag. All four now describe what an operator actually
sees, and the two matrix rows take the new gate count (31 cases, G1-G12) with one
line on what G12's third case bounds; the quantization-matrix row was still
saying 26 cases, G1-G11, and 17 unrouted architectures.
A GCC 13 -O2 build of this branch did not compile
KvAlgoFromObject'sconst auto num_bits = cond ? std::optional<int>(...) : std::nullopt;fails-Werror=maybe-uninitializedat-O2on GCC 13.3, which isevery
ReleaseandRelWithDebInfobuild of this tree. CI configures with noCMAKE_BUILD_TYPEand never saw it;build-newest-gccusesReleaseon anewer compiler that does not emit it. Written as a statement instead, and the
translation unit now compiles clean under
-O2 -g -DNDEBUG -Wall -Wextra -Werror.Gates
tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp— 31 cases / 481assertions GREEN, and
tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp— 3 cases / 26assertions GREEN, plus
tests/vllm/models/test_qwen38_27b_modelopt_mtp_arm.cpp— 22 cases / 1687 assertions GREEN, whose case name this round narrowed.
CPU-only,
-Wall -Wextra -Werror, zero warnings over a full build. Thesrc/vllm/config/cache.cpptranslation unit is compiled a second time at-O2 -g -DNDEBUGfor the GCC 13-Werror=maybe-uninitializedrepair above,because CI configures with no
CMAKE_BUILD_TYPEand cannot see it.ctest584/584 PASSED, 0 failed on the pushed head, on an idle box, over thewhole tree built from scratch with
-Wall -Wextra -Werrorand zero warnings.test_runnerPASSES here in 36.9 s, so the CI red(#1602) is not reproducible on
this host; it is still counted as the pre-existing red rather than as fixed,
because nothing in this branch touched it. Four suites are Skipped for absent
checkpoints or an absent config:
test_modelopt_mixed_precision_checkpoint,test_voxtral_e2e,test_cli_offload_config,test_qwen35_paged_engine. The count is 584 ratherthan CI's 593 because this configure sets
-DVLLM_CPP_BUILD_EXAMPLES=OFF.The round-2 measurement, kept for the record:
ctest591/593 after the second round, on a CONTENDED box (1-minute loadaverage 62 to 100 through the run, other sessions building in their own
worktrees).
test_runneris the pre-existing red(#1602) and fails on an
attention-backend selection message under a non-multiple-of-16 block size, which
this change does not touch; it fails the same way on a serial re-run
(20 cases, 543/544 assertions), so it is a code red that predates this branch and
not a load artefact.
test_engine_core_procfailed underctest -j 4andPASSES on a serial re-run (14 cases / 118 assertions,
Status: SUCCESS!), soit is counted as green rather than as a regression. Three suites are Skipped for
absent checkpoints:
test_modelopt_mixed_precision_checkpoint,test_voxtral_e2e,test_qwen35_paged_engine.scripts/agent-preflight.sh --stagedis green excepttest_cpu_x86_llamacpp_floor(#618), which fails on one run
and passes on the next depending on the box's load: its harness leg is DISCARDED
and RETRIED under contention and the case asserts
returncode == 2against the4 the retry returns. Nothing here was repaired.
CI READ on this head (
b7921ceb0, run32627552373,
CONCLUDED): 17 pass, 6 skipping, 3 fail, and all three are pre-existing and
filed.
build-test-cpuPASSES here, 593/593.test_runner(#1602) is green in that lane
and in both sanitize lanes on this head, and green locally in 36.9 s. On the
previous head it was the ONLY failing test in all three of
build-test-cpu,sanitize-cpu (address,undefined)andsanitize-cpu (thread)— one flakereported as three job failures, not three defects.
cuda-fat-build,build-newest-gcc,build-test-cpu-arm64,build-test-vulkan, bothverifylanes,
commit-protocol-tag,documentation-checkpoint,pr-size,device-leakage,last-gated-commit,cuda-arch-features,vulkan-spirv-freshness,plananddrift-checkall pass.The three reds:
windows-msvc-cputest_openai_api_server.exe exited with status -1073740791(0xC0000409,STATUS_STACK_BUFFER_OVERRUN)windows-msvc-vulkanagent-recordFileNotFoundError: [Errno 2] No such file or directory: 'hugo'attests/scripts/test_check_site.py:130The
windows-msvc-*signature MOVED, and that settles what the catch-up mergebought. Both jobs were failing at step 3, "Prove PowerShell, static CRT, and
unsupported-tier contracts" —
#1649, whose fix
2f2a70925was NOT reachable from this branch's previous base. On this head step 3 PASSES in
both, and what fails is step 4, on the pre-existing #584 crash. #1649 is closed
for this branch; #584 is not this branch's and is not repaired here.
agent-record's case was added by #1714 (1db7e59cf, onmain) and shells outto a binary the job never installs (
grep hugo .github/workflows/ci.ymlfindsnothing). It passes here 7/7, because this host has
hugo. Red onmainand onevery pull request; not repaired here.
commit-trailersandcommit-stylenow RUN.origin/mainhad moved to08c81a89and was no longer an ancestor, so both were SKIPPED and reportednothing at all about this tree. After the catch-up merge preflight prints
Commit trailers vs origin/main 08c81a89218906cac08209a63c6301f03fdc8ec7:withboth green, alongside
now-current rangeandissue-index append-only.Twelve mutations from the first round remain valid, eight more were run for the
second, two for the third, and three more for this round, each applied with its sha256 and applied
diff printed, built with its exit status printed, run red, and restored
byte-identically:
--kv-cache-dtypeparse arm is droppedEngineParamscast_dtback tokv.dtype,WriteKvCache->vt::ReshapeAndCache,ApplyKvCacheQuantdeleted)quantization.quant_methodwidening is restored(q, v, k, ...), so V lands ink_cacheand K inv_cachek_scale * 8/v_scale * 8against an unscaled readcpu_paged_attn.cpp:167)test_ops_fp8_kv_cache8/8 before this roundv_scale = 1.0F)inner)The driver prints the applied diff and the build exit status for every one,
because
git diff --statis EMPTY for a file git does not track yet, and amutation that never applied — or that failed to build — reads as a passing test.
Unreached and owed
All under #1593 and listed in the spec's
## Owed:kv_cache_dtypefield. A C-ABI caller reaches an fp8KV cache only through a checkpoint that declares one, which the loader does
honour on that path.
above). Owned by
KV-FP8under KV-FP8 W2 and W3: the CUDA fp8 KV store, its paged-attention read, and the runner integration — the only thing now gating every cell of the #1574 three-way #1593.arithmetic; the only architecture that populates
per_layer_attn_specsisGemma-4, which is one of the 16 that refuse the store, so a Gemma-4 run gets a
correctly halved pool and then a named refusal — the intended order.
glm4,minicpm,opt,gemma,gemma2,gemma3,gemma4(two sites),commandr,phi,phi3,muse_glimmer,stablelm,qwen3_vl,olmo2,graniteandnemotron_h_device.Thirteen have their own
"<arch>: KV cache must be bf16 or f32"guard;gemma4dies earlier still, insidevt::CastF32, with a message that namesno architecture at all. Only
qwen3_vlreaches the store guard, and onlynemotron_h_devicenames fp8 to the operator. Widening those 15 is the sameedit that routes them, so it is owed with the bullet above.
quant_algoIS present and is neither FP8-family nor NVFP4, and upstream'sconfig convertor raises instead of injecting. The
quant_algo is not Noneguard itself is now mirrored (round 3 above), so this is the one remaining
arm. Recorded, argued and unreachable for the Qwen3.8-27B three-way at each engine's best: one checkpoint ours/vLLM/SGLang all serve, and three of the four recorded NVFP4 blockers are properties of the unsloth artifact rather than the format #1574 subject; porting the
convertor's weight-half validation is its own row.
the flag on both sides and the checkpoint-declaration path this row builds is
gated by G1 and G10 rather than exercised by the benchmark it was built for.
k_scale/v_scale, so the twocheckpoint-loaded arms of the resolver are unit-gated and unreached, and the
scale is per-engine rather than per-layer.
--calculate-kv-scalesis refused, not implemented (upstream deprecates itfor v0.19).
recorded: no
nvccand no device in this session, and a CPU build compileszero
.cutranslation units.scale_exactdecodes
bf16.buf[0], 320 of the 640 elements each run stores), and layer 1carries
3/320elements that dequantize7.62939e-06apart between the twoinvariance runs, in e4m3's subnormal region where the covariance argument does
not hold. The logit equality is empirical there, absorbed below f32, and a
change to the fixture, the layer count, the thread count or the accumulation
order could redden a CORRECT tree with a defect-shaped message.
half. Two fp8 runs cannot see a defect that is a function of bytes and
indices; a V-from-the-K-page read and a V read that drops the in-page offset
both PASS the whole file. Closing that class needs a REFERENCE comparison.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]