feat(VT-MATMUL-FP8-BLOCK-REF): add vt::MatmulFp8BlockScaled, the block-scaled fp8 GEMM whose scales cannot live in an epilogue

[localai-bot]

Milestone M2 of #1189: vt::MatmulFp8BlockScaled, the CPU reference arm of the
128x128 block-scaled FP8 GEMM. It is the numerical oracle every later block-fp8
kernel is measured against. It consumes exactly what vt::QuantFp8Group emits —
M1, landed as ad5f175e7 — and stops there.

Spec: .agents/specs/vt-matmul-fp8-block-ref.md,
committed before the implementation (de5acfd8c, then 4e9e5d651). Pinned oracle:
vLLM 5559679229bc961848b121ccdeaa8fa5d79bec98, asserted as the local checkout's
HEAD before any anchor below was read.

The constraint that decides correctness

The scales apply in the GEMM mainloop, once per K-block, into an f32
accumulator — not in the epilogue:

acc = 0
for kt in [0, cdiv(K, block_k)):
    part = 0                                  a SEPARATE register
    for k in the k-tile:  part += f8(a[m,k]) * f8(b[n,k])
    acc += part * ( a_scale[m,kt] * b_scale[n / block_n, kt] )

vt::MatmulFp8Cutlass folds one scalar alpha after the whole K reduction. An
epilogue has exactly one degree of freedom per output element; this scheme has
cdiv(K, block_k) of them. An epilogue-only application therefore cannot express
a per-K-block scale at all. That is a correctness constraint rather than an
optimisation choice, and it is why this is a new op rather than a parameter of the
existing one.

G4 is the instrument, and it is built so no single-alpha implementation can pass.
Two K-blocks whose partials are 128 and 256 and whose scales are 0.25 and 0.5 give
160; swapping the two K-block scales leaves every per-tensor summary of the
scale tensor identical — same set, sum, product and max — and gives 128. Every value
is exact in f32 and bf16, so these are equalities rather than tolerances. Mutating
the kernel to the epilogue form compiles, runs, and returns the same number for
both.

Which implementation actually executes upstream

M1 found that upstream's readable Triton source is not what runs and that the two
arms disagree in polarity, so the question was asked again here rather than
inherited. The dispatch decision and the kernel, each by file:line:

Step	Where
CUDA block-FP8 kernel priority: FlashInfer, DeepGEMM, CUTLASS, Marlin, Triton, Humming	vllm/model_executor/kernels/linear/__init__.py:355-377
DeepGEMM auto-disabled for qwen3_5_text on device-capability family 120	vllm/utils/deep_gemm.py:27-46
the selected kernel's apply: ops.cutlass_scaled_mm(A, B.T, scale_a=As, scale_b=Bs.T)	vllm/model_executor/kernels/linear/scaled_mm/cutlass.py:312-326
every sm >= 120 device routes to the sm120 entry	csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu:220-226
the sm120 entry hands cutlass_scaled_mm_blockwise_sm120_fp8 to the shared dispatcher	.../cutlass/scaled_mm_c3x_sm120.cu:13-20
the blockwise branch: both scales f32 and 2-D, shapes checked with ceil_div, bias refused	.../c3x/scaled_mm_helper.hpp:15-18,39-55
ElementAccumulator = float, and both scale pointers are mainloop arguments	.../c3x/scaled_mm_blockwise_sm120_fp8_dispatch.cuh:56-58,218-235
CUTLASS 4.5.0, the mainloop line: accum(i) += tmp_accum(i) * tCrScaleAViewAsC(i) * tCrScaleBViewAsC(i)	include/cutlass/gemm/collective/sm120_mma_tma_blockwise_scaling.hpp:714-717

So the executing kernel is CUTLASS, and unlike M1 it agrees with the
reference: same placement, same polarity, f32 accumulator, a multiply. The one
difference is association — CUTLASS multiplies left to right where
native_w8a8_block_matmul forms the scale product first
(tests/kernels/quant_utils.py:150-151) — worth at most one f32 ULP per K-block,
which is precisely what upstream's own gate between the two arms admits at
rel_diff < 0.001 (tests/kernels/quantization/test_block_fp8.py:194-200). We
mirror the reference's association, because this op is the reference port and
milestone M5's CUDA kernel will be measured against it.

The ragged edge is covered, and it is what the grid is for

cdiv, never floor, on both dimensions: upstream asserts
triton.cdiv(N, block_n) == Bs.shape[0] and triton.cdiv(K, block_k) == Bs.shape[1]
(fp8_utils.py:935-936), so a short final block is legal and must work.
n / block_n indexes the b_scale row by output column (fp8_utils.py:823),
which agrees with a tile counter for round N and does not for ragged N.

G3 runs upstream's own non-round shapes — N=576 (4*128 + 64) and K=3884
(30*128 + 44) — separately and together, plus upstream's dedicated DSV3
kv_a_proj_with_mqa case M=32, N=576, K=7168 (test_block_fp8.py:156-200).

This is measured rather than asserted. Two separate floor-versus-ceil mutations —
k_tiles = k / block_k in the kernel, and a b_scale row index of
min(col / block_n, n / block_n - 1) — each compiled, each ran, and each left G2
entirely green, because every N and every K in the round-shape block is a multiple
of 128. Only G3 caught either. A grid of round shapes passes while being wrong,
which is why the checkpoint's own all-multiple-of-128 shapes are not the gate.

Tests and evidence

tests/vt/test_ops_matmul_fp8_block_cpu.cpp, registered in tests/CMakeLists.txt.

RED, re-taken on the merged tree with the implementation reversed out of the
four source files (git diff --stat printing 4 files changed, 200 deletions(-),
so the revert is proven to have landed): the focused build fails compile_rc=1 with
29 errors — 23 'MatmulFp8BlockScaled' is not a member of 'vt' and 6
'kMatmulFp8BlockScaled' is not a member of 'vt::OpId'.

GREEN: 6 cases, 80 assertions, 0 failed, 3.4 s. Per block, filtered through a
-tc prefix that contains no comma, because doctest splits that flag on commas and
a name carrying one reports 0 cases ran under a SUCCESS! banner:

Block	Cases	Assertions
G1 registration and name	1	2
G2 the ported grid	1	22
G3 the ragged edges	1	13
G4 the mainloop constraint	1	14
G5 the refusals	1	25
G6 the M1 seam	1	4
sum	6	80

The buckets sum to the whole-run count, so no block is silently empty and no filter
selected nothing. test_op_provider, test_ops_quant_fp8_group_cpu and
test_ops_fp8_cpu all pass unchanged. scripts/agent-preflight.sh --fail-on-skip
reports All gates green with 80 ok, no FAIL and no SKIP before the verdict
line; test_cpu_x86_llamacpp_floor (#618) passed rather than reporting
NO_QUIET_WINDOW, so no pristine-baseline reproduction was owed.

Ten mutations, each printing compile_rc and git diff --stat because a
mutation that fails to build and a mutation that never applied both read as a
passing test. One of the ten did exactly that — the b_scale index forced to a
constant, which orphaned block_n and died on -Werror=unused-parameter — and it is
reported as proving nothing and re-run in a form that keeps the parameter live. The
full table is in the spec's ## Evidence, along with a false red that cost a cycle:
a direct run reproduced the zeroed-store mutation's exact signature against clean
sources, because the binary was the one the mutation harness had last linked.

The upstream grid is adapted and the adaptation is written into the test. Upstream
runs 125 combinations whose largest is 4096x13824x16384, about 9.3e11
multiply-accumulates; that is a GPU grid, and a naive CPU reference nest run twice
cannot execute it. Every axis value appears at least once and the ragged values
appear separately and together; only the pairing is dropped, and the parameters,
dtypes, tolerance and failure criterion are preserved exactly. The reference is
independently written: upstream's loop nest in double, with an fp8 decode derived
from the e4m3fn field layout rather than from any codec in src/.

Nothing reaches this op yet

vt::MatmulFp8BlockScaled is dispatched by no production entry point at this merge
commit: include/vllm.h does not expose it, no loader builds an Fp8BlockWeight,
and ModelRegistry::Forward has no block-FP8 linear method to call it from. The
wiring is owned by #1189 milestone M4 (layers::Fp8BlockLinearMethod and the
Qwen3.5 dense forward), which needs M3 first. This is the staged-slice exception of
.agents/reachability.md, named here, in the commit
body, and under ## Owed in the spec. G6 composes M1's quantizer with this GEMM end
to end on a CPU queue — a composition test, not a reachability claim.

No CUDA arm lands here: milestone M5 owns the mainloop-scaled CUTLASS kernel for
sm_121a, needs a GPU, and this row took no lease by design. No
.agents/issue-index.md row is appended either: M1 already added the #1189 row and
that file is append-only, so re-stating the issue would duplicate rather than merge.

Closes nothing; #1189 stays open for M3 through M6.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5-1m [claude-code]