Skip to content

feat(KV-FP8): W3 -- half-sized KV blocks the runner actually serves, and a default scale that knows it was declared (#1593) - #1641

Merged
localai-bot merged 17 commits into
mainfrom
row/KV-FP8-W3
Aug 23, 2026
Merged

feat(KV-FP8): W3 -- half-sized KV blocks the runner actually serves, and a default scale that knows it was declared (#1593)#1641
localai-bot merged 17 commits into
mainfrom
row/KV-FP8-W3

Conversation

@localai-bot

@localai-bot localai-bot commented Aug 21, 2026

Copy link
Copy Markdown
Collaborator

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.

Now based on main. W2 merged as squash 53f5b74be, so this branch's
previous base row/KV-FP8-W2 is dead and GitHub was computing mergeability
against it. origin/main is 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 68d2e0d5a returned FAIL — not safe to merge, with one
functional bug. Every finding is repaired here, and the second review re-proved
every one of them by mutation.

The bug: --kv-cache-dtype fp8 threw at the first forward on every real checkpoint

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. So K
still arrived as f32 — attn_dt is f32 for every fp8 cache, because
kv.dtype == DType::kBF16 is a term of both FA2 eligibility tests — while V
arrived as whatever the v_proj GEMM emitted. That is bf16 on the block-wise
fp8 arm (MatmulFp8BlockScaledD), bf16 on the NVFP4 arm under the default
VT_BF16_GEMM_OUT, and bf16 on ordinary torch safetensors (MatmulBf16D).
vt::ReshapeAndCacheFp8 quantizes from ONE source dtype, so the flag threw
reshape_and_cache_fp8: k/v must share one float dtype at the first forward on
every 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. MakeDenseWeights builds projection weights with no
nk flag, ProjectFullAttnQkv serves those through MatmulF32D, and f32 pairs
with 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 rather
than f32 because bf16 is the dtype upstream quantizes from: vLLM's model IS bf16
where reshape_and_cache_flash takes key/value
(cache_kernels.cu:314-401). G9 is the new case, and it enters through
LoadedEngine with the raw torch Linear layout ([N=out, K=in], nk) a real
safetensors checkpoint ships.

The merge conflict in model_loader.cpp was SEMANTIC, and both sides survive

main's SPEC-DRAFTER-CHAIN W1 refusal (#1522) 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 chain refusal before any weight I/O, and
ReadQuantConfigJson opens a file inside model_dir. Taking either side drops a
guarantee. Both survive, chain refusal first, and the order is gated.

tests/vllm/entrypoints/test_drafter_chain_reach.cpp cannot catch an inversion:
it points at a NONEXISTENT directory, and ReadQuantConfigJson answers "" for
one 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 edit
re-applied by hand. The union merge had duplicated two issue-index rows
(#1593 and #1636 appear once on main and twice after the automatic merge);
that is the known union-driver trap and it is undone here.

Three surfaces that no gate reached

Surface The review's mutation Now gated by
FromModelDir's resolution stanza + ReadQuantConfigJson, whose only caller it is deleting the whole stanza left 19/19, 89/89 SUCCESS G10, through LoadedEngine::FromModelDir on a real directory
the --kv-cache-dtype parse arm and engine_params.kv_cache_dtype = args.kv_cache_dtype deleting both left the same green, and no test in the tree named the flag new tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp, driving the REAL VllmServerMain
ApplyCacheDType's per_layer_attn_specs loop deleting it left the same green G11, on the arithmetic it claims: every non-null spec retyped, and KVBytesPerBlock exactly halved

The 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
auto arrives, the checkpoint wins, and the line comes back.

The file precedence was inverted against upstream

ReadQuantConfigJson read hf_quant_config.json first and fell back to
config.json:quantization_config. vllm/transformers_utils/config.py:751-761
does 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 the
resolver 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.cpp makes kv.dtype == DType::kBF16 a term of both
fa2_prefill and fa2_decode; src/vt/cuda/cuda_paged_attn.cu:2941-2960 routes
an 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.md and both matrix rows now say the
exchange 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 947c5b65f returned FAIL, with one blocking finding. Every
finding is repaired here, and the first round's repairs are untouched.

BLOCKING: the shared seam's fp8 routing was DEAD CODE

dense_attn::AttnBlock is named in AGENTS.md ## Shared seams as the decode
seam, and W3 listed it as routed. It was not. The routing was written — fp8_kv,
WriteKvCache, ApplyKvCacheQuant — underneath a preamble guard that still
admitted only kBF16 and kF32, and IsFp8KvCache is true only for kI8. So
fp8_kv was provably false at every call and neither fp8 arm could execute. Not
an 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::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. --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, and G12 is
the new gate. It enters through Qwen3DenseModel::Forward, which is what
ForwardQwen3ForCausalLM calls under ModelRegistry::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_config
carries quant_method: "modelopt", quant_algo: "MIXED_PRECISION" and no
kv_cache_* key at all
; only hf_quant_config.json declares
kv_cache_quant_algo: "FP8". The first round's own file-precedence repair makes
the 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.cpp and this body
said about it. The campaign consequence, stated plainly: for #1574 the flag has
to be typed explicitly on both sides.
The competitors' own serve.sh already
passes --kv-cache-dtype fp8, so this is consistent rather than a blocker. A new
G10 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.cpp accepting producer.name as a
divergence, on the grounds that get_kv_cache_quant_algo_string
(torch_utils.py:319) gates on a top-level quant_method and nothing else. That
transcription is of one function; 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 very dict hf_config.quantization_config names.

MEASURED by extracting both functions with ast from the files at 555967922
and running them over the live documents rather than retyping them:
nvidia/Llama-3.3-70B-Instruct-FP8's producer-only hf_quant_config.json
answers None before normalization and 'fp8_e4m3' after it. So accepting the
producer name is the MIRROR. The comment that claimed it on the wrong grounds
("modelopt_mixed_precision.h already reads both shapes") is replaced by the
chain and the measurement.

Reading the injector did surface three real differences, and the marker set is
now exactly upstream's:

  • quantization.quant_method is no longer accepted. Upstream writes and reads
    the marker at the TOP level and never looks inside, so a document whose only
    marker is 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 we ignored.
  • The producer name is compared raw against the literal (:222 is ==), not
    prefix-matched and case-folded. quant_method keeps both, because :238-246
    lower-cases it before :319 prefix-tests it. Same file, two tests, mirrored
    separately.

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 the
producer is modelopt and the nested quant_algo is neither FP8-family nor
NVFP4 — which is what the #1574 legacy file (MIXED_PRECISION) gets, and we
answer fp8_e4m3 for it 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 checkpoint whose weights modelopt_mixed_precision.h
loads. It is unreachable for that checkpoint anyway, because its inline document
wins. Spec ## Owed names it.

The refusal mechanism was wrong for 13 of the 16 unrouted architectures

"Every other architecture is refused BY NAME at vt::ReshapeAndCache" holds for
three. granite:95, minicpm:96, phi3:78, gemma3:121, opt:125,
stablelm:86, glm4:93, commandr:93, gemma:53, gemma2:135, phi:98,
muse_glimmer:144 and olmo2:94 each refuse at their own
"<arch>: KV cache must be bf16 or f32" guard first, which names neither fp8
nor the flag. Only gemma4, qwen3_vl and nemotron_h_device reach the store
guard. 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_vl reaches the store guard; the split is 14 / 1 / 1. See
below.
Corrected in the spec, in
docs/USAGE.md, in docs/FEATURES.md and in the engine-matrix row, and the
13-guard gap is now its own ## Owed bullet.

docs/FEATURES.md promised a doubled pool on a path that halves it

ResolveNumBlocks (model_loader.cpp:1446-1451) returns the fixed 256-block
fallback when neither --num-blocks nor --kv-cache-memory is given, so on the
default 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-memory and was correct, and now
says so explicitly.

Preflight was reporting nothing about two gates

origin/main had moved to 08c81a89 and was no longer an ancestor, so the
diff-scoped commit-trailers and commit-style checks SKIPPED. origin/main is
merged up and both execute.

What the THIRD fresh review found, and what changed

The review of d0bcf80dd returned FAIL, with one blocking finding. Every
finding is repaired here, and rounds 1 and 2 are untouched. The round-2
producer.name MEDIUM was FALSIFIED by that review, which ast-extracted both
pinned upstream functions and ran them: _normalize_quantization_config mutates
hf_config.quantization_config in place during ModelConfig.__post_init__,
before resolve_kv_cache_dtype_string reads it, so accepting producer.name is
the mirror. It stays.

BLOCKING: the gate could not distinguish a working fp8 KV cache from a corrupted one

G12's counter-case asserted differing > 0 and max_abs < 1.0 on a delta its
own MESSAGE reported as 0.000339303. The bound sat about 3000x above the
signal
, so it measured this toy model's insensitivity to its KV cache and not
the cache. Two mutations of kv_cache_route.h:63 walked through the whole 30/30
suite:

Mutation Delta on the old axis Old suite
N1_KVSWAP — pass (q, v, k, k_cache, v_cache, ...), so the fp8 arm stores V into k_cache and K into v_cache 0.0247495 (73x) 30/30 SUCCESS
N4_STORE_SCALE8 — store with kv.k_scale * 8.0F / kv.v_scale * 8.0F while the read uses the unscaled values 0.0063844 (19x) 30/30 SUCCESS

The second one is exactly the defect the k_scale/v_scale plumbing exists to
prevent. 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:

|dequant(store(ref)) - ref|  <=  2^-4 * |ref|  +  2^-10 * scale

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 most
2^-4 of its own value — and anything below the smallest normal 2^-6 sits on
a uniform 2^-9 grid whose half step is 2^-10, in absolute terms scaled by the
per-tensor scale. Not one constant is fitted to a measurement, which is the point:
a threshold sized to today's 3.4e-4 is 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.125 for K and 0.25 for V — non-unit and unequal, so a
dropped, 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 most
elements under 2^-6, where only the weaker absolute arm applies. Two anti-vacuity
REQUIREs 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:

Tree Elements outside the envelope Worst ratio to the bound
repaired 0 / 320 (281/320 normals over 10/10 pages) 0
N1_KVSWAP 316 / 320 416.9
N4_STORE_SCALE8 319 / 320 13.9

A bf16-versus-fp8 comparison in G5/G9 through LoadedEngine was considered and
is 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::LoadKvFp8E4M3 and never enters the
production dequant (cpu_paged_attn.cpp:167); 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 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)) the
representable points are m * 2^(e-3) — and dividing by a power of two is exact
in binary floating point, shifting e without touching the mantissa. So for any
two power-of-two scales that both leave a value normal and unsaturated,

s * Dequant(Quantize(x / s))  ==  s' * Dequant(Quantize(x / s'))

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 requires memcmp-level
agreement 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_exact REQUIRE
decodes bf16.buf[0] and holds 320/320 LAYER-0 elements normal and unsaturated
at all four scales — the measured magnitudes are 1.76e-4 to 1.32e-1 for K and
5.41e-5 to 4.22e-2 for V, against all-normal windows of
(2.94e-4, 1.13e-2] and (9.43e-5, 3.46e-3] — because one subnormal would round
on the ABSOLUTE 2^-9 grid, which is not scale invariant. MakeSeamConfig sets
num_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 own
scales, 3/320 of its elements disagree 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 land in e4m3's SUBNORMAL region and that layer carries
2 SATURATED elements besides. CHECK(inv_differing == 0) therefore holds by the
format property for layer 0 and by ABSORPTION for layer 1 — a 7.6e-6 cache
perturbation vanishing in f32 accumulation before it reaches a logit. Extending
scale_exact over every layer was the preferred repair and this fixture cannot
satisfy it, so the claim is narrowed here and the fragility is recorded under
## Owed.

The second anti-vacuity REQUIRE is that the two caches really do hold different
bytes, so the comparison is a statement about the read rather than about two
identical buffers.

MEASURED on the repaired tree: 0/320 logits differ, max |delta| exactly 0.
Three read-side mutations of cpu_paged_attn.cpp:167 are RED, and the store
envelope above reads 0/320 under all three:

Read mutation Logits differing max delta
v_scale = args.k_scale 320 / 320 0.0673
k_scale = args.v_scale 256 / 320 1.08e-4
v_scale = 1.0F 320 / 320 2.47e-3

The 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.0673 puts the constant at
5.4e-3, which admits 1.08e-4 by 50x. An absolute 1e-4 would in fact
have CAUGHT the K-side one, by 8.4%, and a relative 1e-4 against the fixture's
largest logit (0.0627671) is 6.3e-6 absolute and catches it by 17x — the
argument 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, with 0/320 logits differ: serving V 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 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/main exiting 1
with commit-trailers and commit-style both SKIPPED: origin/main had moved to
c98ffd4d0 while this branch's last merge took ff8f72807, so neither gate
reported anything about this tree. main is 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/main had moved again, to ea9b7e30e, so it is
merged up once more and commit-trailers and commit-style execute against
this tree rather than SKIP on a base that has moved.

The first was the tolerance sentence. "A tolerance sized at 1e-4 would have let
it through" is false as stated and false on either reading: 1.08408e-4 is 8.4%
ABOVE an absolute 1e-4, and a relative 1e-4 against the fixture's largest
logit 0.0627671 is 6.3e-6 and catches it by 17x. The argument that survives
is 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_exact REQUIRE below holds the population to it". scale_exact
decodes bf16.buf[0], which is 320 of the 640 elements each run stores, and the
logits 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/320 elements disagreeing by up to 7.62939e-06, subnormal at 2^-7. So the
prose 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 ## Owed refusal accounting was wrong for 2 of the 3 architectures it named

Round 2 said gemma4, qwen3_vl and nemotron_h_device reach the store guard.
Traced with a kI8 page against a bf16 model dtype, only qwen3_vl does. The
true split is 14 / 1 / 1:

  • gemma4.cpp:306-315 takes kv.dtype != adt, allocates
    DBuf kcast(d, kv.dtype /* kI8 */, ...) and calls vt::CastF32, which refuses
    at src/vt/ops.cpp:4087 with "cast_f32: out must be f32". It never reaches
    the store, and that message names neither fp8, nor the flag, nor gemma4 — so
    it is the WORST of the sixteen, not one of the three best.
  • nemotron_h_device.cpp:1589-1593 carries an explicit
    else { VT_CHECK(false, "NemotronH paged forward: ... The fp8 KV scheme ... is not selected here"); } on the same cast, which fires first. It is the ONE
    refusal in the sixteen that names fp8 to the operator.
  • qwen3_vl.cpp:198-200 carries no guard and no cast, so vt::ReshapeAndCache
    is 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 ## Owed
bullet, which now also records that G7 hand-builds its K/V tensors and calls
vt::ReshapeAndCache directly, so it gates the store guard's MESSAGE and never
the claim that any particular architecture reaches it.

The marker mirror accepted three shapes upstream answers None for

_normalize_quantization_config injects quant_method only
if quant_algo is not None (:224), and it reads 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.
cache.cpp accepted the producer name and the legacy key unconditionally, and
read the legacy key out of the reader's inner. Three documents therefore
resolved to fp8_e4m3 here while upstream resolved None:

Document Upstream Before Now
producer.name == "modelopt", quantization with no quant_algo None fp8_e4m3 auto
quantization.modelopt_quant_config with no quant_algo None fp8_e4m3 auto
top-level modelopt_quant_config, no quantization key at all None fp8_e4m3 auto

Each 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_algo upstream requires, so what
refuses 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 mirrored
test is exactly quant_algo is not None rather than the FP8/NVFP4 family set:
the only two upstream outcomes for a quant_algo that 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:207 now does, from LoadedEngine::FromModelDir
(model_loader.cpp:1988). Four statements written against the old world read
false at this head, and a record edit rides in the pull request whose change made
it stale:

Site Was Now
modelopt_mixed_precision.h:1006 "a file no production path in this tree reads at all" reads it, but only as the LEGACY FALLBACK behind config.json's inline document, and this checkpoint ships that document
.agents/specs/qwen38-27b-quant-arms.md:1629 "the declaration is invisible to the loader rather than ignored by it" the same conclusion, reached through the precedence rather than through absence
.agents/specs/qwen38-27b-quant-arms.md:1841 "No production path in this tree reads hf_quant_config.json" names the reader, the fallback order, and that BOTH artifacts carry the inline document
test_qwen38_27b_modelopt_mtp_arm.cpp:482 (case NAME) and :487-488 "where NO production path reads it" / "the ONLY document any production path in this tree reads" "in the file this artifact's config.json outranks", with the precedence spelled out

No assertion changed. They stay valid because config.json's inline
quantization_config wins 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, the KV-FP8 engine-matrix row and the QUANT-KV-FP8
quantization-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's const auto num_bits = cond ? std::optional<int>(...) : std::nullopt; fails -Werror=maybe-uninitialized at -O2 on GCC 13.3, which is
every Release and RelWithDebInfo build of this tree. CI configures with no
CMAKE_BUILD_TYPE and never saw it; build-newest-gcc uses Release on a
newer 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.cpp31 cases / 481
assertions GREEN
, and
tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp3 cases / 26
assertions 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. The
src/vllm/config/cache.cpp translation unit is compiled a second time at
-O2 -g -DNDEBUG for the GCC 13 -Werror=maybe-uninitialized repair above,
because CI configures with no CMAKE_BUILD_TYPE and cannot see it.

ctest 584/584 PASSED, 0 failed on the pushed head, on an idle box, over the
whole tree built from scratch with -Wall -Wextra -Werror and zero warnings.
test_runner PASSES 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 rather
than CI's 593 because this configure sets -DVLLM_CPP_BUILD_EXAMPLES=OFF.

The round-2 measurement, kept for the record:
ctest 591/593 after the second round, on a CONTENDED box (1-minute load
average 62 to 100 through the run, other sessions building in their own
worktrees). test_runner is 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_proc failed under ctest -j 4 and
PASSES on a serial re-run (14 cases / 118 assertions, Status: SUCCESS!), so
it 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 --staged is green except
test_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 == 2 against the
4 the retry returns. Nothing here was repaired.

CI READ on this head (b7921ceb0, run
32627552373,
CONCLUDED): 17 pass, 6 skipping, 3 fail, and all three are pre-existing and
filed.

build-test-cpu PASSES 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) and sanitize-cpu (thread) — one flake
reported as three job failures, not three defects. cuda-fat-build,
build-newest-gcc, build-test-cpu-arm64, build-test-vulkan, both verify
lanes, commit-protocol-tag, documentation-checkpoint, pr-size,
device-leakage, last-gated-commit, cuda-arch-features,
vulkan-spirv-freshness, plan and drift-check all pass.

The three reds:

Job What it reports Whose
windows-msvc-cpu test_openai_api_server.exe exited with status -1073740791 (0xC0000409, STATUS_STACK_BUFFER_OVERRUN) #584, open, reproduced unchanged on PR #1744
windows-msvc-vulkan the same executable, the same status #584
agent-record FileNotFoundError: [Errno 2] No such file or directory: 'hugo' at tests/scripts/test_check_site.py:130 #1722, #1754, #1764; PR #1726 is the fix in flight

The windows-msvc-* signature MOVED, and that settles what the catch-up merge
bought.
Both jobs were failing at step 3, "Prove PowerShell, static CRT, and
unsupported-tier contracts"

#1649, whose fix 2f2a70925
was 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, on main) and shells out
to a binary the job never installs (grep hugo .github/workflows/ci.yml finds
nothing). It passes here 7/7, because this host has hugo. Red on main and on
every pull request; not repaired here.

commit-trailers and commit-style now RUN. origin/main had moved to
08c81a89 and was no longer an ancestor, so both were SKIPPED and reported
nothing at all about this tree. After the catch-up merge preflight prints
Commit trailers vs origin/main 08c81a89218906cac08209a63c6301f03fdc8ec7: with
both green, alongside now-current range and issue-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:

Mutation Gate that went red
M13 the fp8 route stops normalising K/V to one dtype G9 (the production arm)
M14 the --kv-cache-dtype parse arm is dropped serve CASE 2 and CASE 3
M15 the parsed flag never reaches EngineParams serve CASE 2 and CASE 3
M16 the loader never reads the checkpoint's declaration G10
M17 the legacy file outranks the current one again G10 (precedence, both halves)
M18 the per-layer specs are never retyped G11 (9 assertions)
M19 the chain refusal is deleted G10 (ordering)
M19b the chain refusal MOVES below the stanza G10 (ordering) — the mutation M19 cannot make
MSEAM the shared seam's fp8 routing is reverted whole (cast_dt back to kv.dtype, WriteKvCache -> vt::ReshapeAndCache, ApplyKvCacheQuant deleted) G12, both cases — the mutation that stayed GREEN before this round
MMARKER the nested quantization.quant_method widening is restored G1's marker case, arms (c) and (d)
N1_KVSWAP the fp8 store is handed (q, v, k, ...), so V lands in k_cache and K in v_cache G12's envelope case: 316/320 elements outside, worst ratio 416.9 — GREEN 30/30 before this round
N4_STORE_SCALE8 the fp8 store uses k_scale * 8 / v_scale * 8 against an unscaled read G12's envelope case: 319/320 outside, worst ratio 13.9 — GREEN 30/30 before this round
N5_READ_VSWAP the production read dequantizes V with K's scale (cpu_paged_attn.cpp:167) G12's envelope case, the new invariance half: 320/320 logits differ at max delta 0.0673 — GREEN 31/31 AND test_ops_fp8_kv_cache 8/8 before this round
N6_READ_KSWAP the production read dequantizes K with V's scale same: 256/320 logits differ at max delta 1.08e-4
N7_READ_VDROP the production read drops the V scale (v_scale = 1.0F) same: 320/320 logits differ at max delta 2.47e-3
N_LOW1 the marker predicate is restored to its pre-repair form (unguarded producer/legacy key, legacy key read out of inner) G1's marker case, arm (g): 7 assertions red across all three shapes

The driver prints the applied diff and the build exit status for every one,
because git diff --stat is EMPTY for a file git does not track yet, and a
mutation 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:

  • The C ABI carries no kv_cache_dtype field. A C-ABI caller reaches an fp8
    KV cache only through a checkpoint that declares one, which the loader does
    honour on that path.
  • An fp8 KV cache is served by the SLOW attention kernels only (the section
    above). Owned by KV-FP8 under 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.
  • The heterogeneous per-layer fp8 arm is SIZING-only. G11 gates the
    arithmetic; the only architecture that populates per_layer_attn_specs is
    Gemma-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.
  • 16 architectures refuse rather than route, at 17 call sites: glm4,
    minicpm, opt, gemma, gemma2, gemma3, gemma4 (two sites),
    commandr, phi, phi3, muse_glimmer, stablelm, qwen3_vl, olmo2,
    granite and nemotron_h_device.
  • 15 of those 16 refuse with a message naming neither fp8 nor the flag.
    Thirteen have their own "<arch>: KV cache must be bf16 or f32" guard;
    gemma4 dies earlier still, inside vt::CastF32, with a message that names
    no architecture at all. Only qwen3_vl reaches the store guard, and only
    nemotron_h_device names fp8 to the operator. Widening those 15 is the same
    edit that routes them, so it is owed with the bullet above.
  • The resolver accepts one document upstream REFUSES, where the nested
    quant_algo IS present and is neither FP8-family nor NVFP4, and upstream's
    config convertor raises instead of injecting. The quant_algo is not None
    guard 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.
  • No engine auto-selects fp8 KV 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, so the campaign types
    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.
  • No weight loader extracts k_scale/v_scale, so the two
    checkpoint-loaded arms of the resolver are unit-gated and unreached, and the
    scale is per-engine rather than per-layer.
  • --calculate-kv-scales is refused, not implemented (upstream deprecates it
    for v0.19).
  • Every CUDA case in this feature is still UNEXECUTED, for the reason W2
    recorded: no nvcc and no device in this session, and a CPU build compiles
    zero .cu translation units.
  • G12's read-side exactness is ASSERTED for LAYER 0 only (scale_exact
    decodes bf16.buf[0], 320 of the 640 elements each run stores), and layer 1
    carries 3/320 elements that dequantize 7.62939e-06 apart between the two
    invariance 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.
  • The read comparison closes the SCALE half of the read and not the ROUTING
    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]

mudler added 4 commits August 21, 2026 14:49
… 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
localai-bot changed the base branch from row/KV-FP8-W2 to main August 22, 2026 07:37
@localai-bot
localai-bot marked this pull request as ready for review August 22, 2026 07:38
mudler added 5 commits August 22, 2026 07:50
…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]
mudler added 8 commits August 23, 2026 07:44
…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]
@localai-bot
localai-bot merged commit 27d8bfa into main Aug 23, 2026
2 of 4 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants