From 16a8cdd156b9b367403ffd3dd561d7805e2ea9bd Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 01:33:35 +0200 Subject: [PATCH 1/2] feat(BACKEND-TENSTORRENT-KEEPQUANT): dot Q4_K weights on device via the W1 decode (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit W2 of the row: the keep-quant DOT. MatmulBTQuantKernel registers kMatmulBTQuant on the P150 and is reached through vt::MatmulBT's block-weight dispatch (ops.cpp:163) — the entry every model matmul helper already uses, so the test drives the public seam and not a hand-cast op pointer. tt-metal has no packed-weight GEMM primitive, so the composition decodes the blocks through the W1 bit-exact chain (DecodeQ4KBlocksF32, factored out of the decode kernel verbatim — one decode, two consumers, one numerics authority), rounds the decoded weight once to bf16 (RNE, the device round-once convention), and runs the same bf16 tile matmul as kMatmulBT. The M=1 GEMV — the production decode shape — is in the sweep. The second half is the kTENSTORRENT arm of DeviceKeepQuantSupported, admitting exactly {Q4_K}. The P150 is discrete with no CPU fallback tier, so an arm wider than the registered kernel set throws at first forward with the model resident — the exact failure this predicate exists to prevent, and the reason Q5_K/Q6_K/Q8_0 stay on the expand-bf16 residency until their W4 kernels land. The routing test pins the set host-side; a reviewer who widens the arm without widening the kernel reds it. Red-first on both halves, captured before the implementation: the device test's REQUIRE(OpRegistered(kMatmulBTQuant, kTENSTORRENT)) fatalled, and all six unimplemented encodings routed kKeepQuant on TT under the default arm — the too-wide answer this change replaces. Green after: the dot sits inside the analytic bf16 operand-rounding envelope of the decode-based reference (worst bound-ratio 0.53 across M {1,5} x K {256,512}; the bound is computed from the data, 1.05 * 2^-8 * (sum |a_k w_k| + |ref|), not a picked tolerance), full suites 54/54 and 53/53 on the device. Reachability note: the op-level dot is reachable, the LOADER path is not until the W3 e2e battery runs a Q4_K GGUF end to end on the P150; that scope is named here, owned by the row, and listed in the spec's Owed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- .../model_loader/gguf_keep_quant.cpp | 10 ++ src/vt/tenstorrent/tenstorrent_ops.cpp | 104 +++++++++++++--- tests/vllm/test_gguf_keep_quant.cpp | 34 ++++++ tests/vt/test_tenstorrent_backend.cpp | 111 ++++++++++++++++++ 4 files changed, 241 insertions(+), 18 deletions(-) diff --git a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp index 65cc5db456..4e670018f8 100644 --- a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp +++ b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp @@ -141,6 +141,16 @@ bool DeviceKeepQuantSupported(vt::DType dt, vt::DeviceType dev) { // owed (recorded in .agents/specs/rocm-gg-keep-quant.md). return dt == vt::DType::kQ8_0 || dt == vt::DType::kQ4_K || dt == vt::DType::kQ5_K || dt == vt::DType::kQ6_K; + case vt::DeviceType::kTENSTORRENT: + // KEEPQUANT W2: the P150 is discrete with no CPU fallback tier, so this + // arm admits exactly what src/vt/tenstorrent/tenstorrent_ops.cpp has a + // registered decode for — MatmulBTQuantKernel implements Q4_K (through + // the W1 bit-exact DecodeQ4KBlocksF32 chain). Q5_K/Q6_K/Q8_0 are the + // row's W4; admitting them here before their kernels land throws at + // first forward with the model resident, which is the exact failure + // this predicate exists to prevent. tests/vllm/test_gguf_keep_quant.cpp + // pins the set; widening the arm without widening the kernel reds it. + return dt == vt::DType::kQ4_K; default: // CUDA falls back to the CPU kernel for anything it lacks // (cuda_quant_dot.cu:1841-1846); the CPU list IS the CPU capability. diff --git a/src/vt/tenstorrent/tenstorrent_ops.cpp b/src/vt/tenstorrent/tenstorrent_ops.cpp index fb770f9b6f..2454d8afea 100644 --- a/src/vt/tenstorrent/tenstorrent_ops.cpp +++ b/src/vt/tenstorrent/tenstorrent_ops.cpp @@ -1743,21 +1743,14 @@ void EmbeddingKernel(Queue&, Tensor& out, const Tensor& table, const Tensor& ids // host value and the device buffer (multiply and the i32<->f32 bitcast both // canonicalize -0 to +0). where() consumes it as the true branch of the // signed-zero repair. -void KeepQuantDecodeKernel(Queue&, Tensor& out, const Tensor& packed) { - TT_OP_TRACE("KeepQuantDecode"); - VT_CHECK(packed.rank == 2 && out.rank == 2, - "tenstorrent kKeepQuantDecode: packed rank-2 [rows, nb], out " - "rank-2 [rows, nb*256]"); - VT_CHECK(packed.dtype == DType::kQ4_K, - "tenstorrent kKeepQuantDecode: packed dtype must be kQ4_K"); - VT_CHECK(out.dtype == DType::kF32, - "tenstorrent kKeepQuantDecode: out must be f32"); - VT_CHECK(packed.IsContiguous() && out.IsContiguous(), - "tenstorrent kKeepQuantDecode: contiguous required"); - const int64_t rows = packed.shape[0]; - const int64_t nb = packed.shape[1]; - VT_CHECK(out.shape[0] == rows && out.shape[1] == nb * 256, - "tenstorrent kKeepQuantDecode: out shape mismatch"); +// +// THE SHARED DECODE (KEEPQUANT W2): the packed-Q4_K -> f32 chain verbatim, +// returning the repaired f32 {rows, nb*256} in ROW_MAJOR. KeepQuantDecode +// commits it to the host-visible output; the quant matmul (MatmulBTQuantKernel +// below) consumes it as the weight operand. One decode, two consumers, one +// numerics authority. +ttnn::Tensor DecodeQ4KBlocksF32(const Tensor& packed, int64_t rows, int64_t nb, + MeshDevice& device) { const int64_t b64 = rows * nb; const uint32_t B = static_cast(b64); @@ -1777,7 +1770,6 @@ void KeepQuantDecodeKernel(Queue&, Tensor& out, const Tensor& packed) { (static_cast(p[3]) << 24))); } } - MeshDevice& device = SharedMeshDevice(); ttnn::Tensor w = ttnn::Tensor::from_vector( std::move(words), SpecOf(tt::tt_metal::Shape({B, 36u}), ttnn::DataType::INT32, @@ -1968,12 +1960,86 @@ void KeepQuantDecodeKernel(Queue&, Tensor& out, const Tensor& packed) { m1 = ttnn::to_layout( ttnn::where(pred8, neg0T8, m1T), ttnn::Layout::ROW_MAJOR); ttnn::Tensor y = ttnn::subtract(prod, m1); - y = ttnn::reshape(ttnn::to_layout(std::move(y), ttnn::Layout::ROW_MAJOR), - ttnn::Shape({B, 256u})); + return ttnn::reshape(ttnn::to_layout(std::move(y), ttnn::Layout::ROW_MAJOR), + ttnn::Shape({static_cast(rows), + static_cast(nb) * 256u})); +} + +void KeepQuantDecodeKernel(Queue&, Tensor& out, const Tensor& packed) { + TT_OP_TRACE("KeepQuantDecode"); + VT_CHECK(packed.rank == 2 && out.rank == 2, + "tenstorrent kKeepQuantDecode: packed rank-2 [rows, nb], out " + "rank-2 [rows, nb*256]"); + VT_CHECK(packed.dtype == DType::kQ4_K, + "tenstorrent kKeepQuantDecode: packed dtype must be kQ4_K"); + VT_CHECK(out.dtype == DType::kF32, + "tenstorrent kKeepQuantDecode: out must be f32"); + VT_CHECK(packed.IsContiguous() && out.IsContiguous(), + "tenstorrent kKeepQuantDecode: contiguous required"); + const int64_t rows = packed.shape[0]; + const int64_t nb = packed.shape[1]; + VT_CHECK(out.shape[0] == rows && out.shape[1] == nb * 256, + "tenstorrent kKeepQuantDecode: out shape mismatch"); + ttnn::Tensor y = DecodeQ4KBlocksF32(packed, rows, nb, SharedMeshDevice()); CommitDeviceLogical2D(out, std::move(y), static_cast(rows), static_cast(nb * 256)); } +// kMatmulBTQuant (KEEPQUANT W2): `a` is [M,K] float, `b` is [N,K] packed +// Q4_K blocks — out = a @ b^T, reached through vt::MatmulBT's block-weight +// dispatch (ops.cpp:163), the entry every model matmul helper already uses. +// tt-metal has no packed-weight GEMM primitive (row survey), so the +// composition decodes the blocks through the W1 bit-exact f32 chain +// (DecodeQ4KBlocksF32 above), rounds the decoded weight ONCE to bf16 (RNE, +// the device's round-once convention), and runs the same bf16 tile matmul as +// kMatmulBT. The activation takes the device's ordinary bf16 tile path (an +// f32 master rounds once on device, the same widen-on-load rule as bf16 +// weights). Numerics sit inside the analytic bf16 operand-rounding envelope +// of the decode-based reference — the W2 test pins it, not a picked +// tolerance. Eager staging only today: the decode reads the host-mapped +// blocks per call; capture-safe residency is W3 (#2959). +void MatmulBTQuantKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& b) { + TT_OP_TRACE("MatmulBTQuant"); + VT_CHECK(a.rank == 2 && b.rank == 2 && out.rank == 2, + "tenstorrent kMatmulBTQuant: rank-2 a/b/out required"); + // Exactly the encodings DeviceKeepQuantSupported admits on kTENSTORRENT + // (gguf_keep_quant.cpp). Q5_K/Q6_K/Q8_0 decodes are owed by the row's W4; + // refusing here BY NAME keeps an admitted-but-unimplemented encoding from + // reaching the device. + VT_CHECK(b.dtype == DType::kQ4_K, + "tenstorrent kMatmulBTQuant: only kQ4_K decodes on TENSTORRENT " + "today; kQ5_K/kQ6_K/kQ8_0 are owed by BACKEND-TENSTORRENT-" + "KEEPQUANT W4"); + VT_CHECK(b.shape[1] % 256 == 0, + "tenstorrent kMatmulBTQuant: K must be a whole number of Q4_K " + "blocks (256 elems)"); + VT_CHECK(IsFloatDType(a.dtype) && + (out.dtype == DType::kF32 || out.dtype == DType::kBF16), + "tenstorrent kMatmulBTQuant: float activation, f32/bf16 out"); + const uint32_t M = static_cast(a.shape[0]); + const uint32_t K = static_cast(a.shape[1]); + const uint32_t N = static_cast(b.shape[0]); + VT_CHECK(b.shape[1] == K, "tenstorrent kMatmulBTQuant: a/b inner dim mismatch"); + VT_CHECK(out.shape[0] == M && out.shape[1] == N, + "tenstorrent kMatmulBTQuant: out shape mismatch"); + VT_CHECK(a.IsContiguous() && b.IsContiguous() && out.IsContiguous(), + "tenstorrent kMatmulBTQuant: strided tensors are not supported in W2"); + + MeshDevice& device = SharedMeshDevice(); + ttnn::Tensor w_f32 = DecodeQ4KBlocksF32(b, N, K / 256, device); + ttnn::Tensor w_bf16 = ttnn::to_layout( + ttnn::typecast(std::move(w_f32), ttnn::DataType::BFLOAT16), + ttnn::Layout::TILE); + ttnn::Tensor dev_a = EnsureDevice2D(a, device); + if (a.dtype == DType::kF32) + dev_a = ttnn::to_layout( + ttnn::typecast(std::move(dev_a), ttnn::DataType::BFLOAT16), + ttnn::Layout::TILE); + ttnn::Tensor dev_c = ttnn::operations::matmul::matmul( + dev_a, w_bf16, /*transpose_a=*/false, /*transpose_b=*/true); + CommitDevice2D(out, std::move(dev_c)); +} + // Upload a rank-1 affine vector as TILE BFLOAT16 [1, d], caching on the weight's // host buffer slot so RmsNorm/LayerNorm do not re-upload every layer call. // ttnn's TILE-gamma path requires padded height == tile_height (32); from_vector @@ -6658,6 +6724,8 @@ struct Registrar { reinterpret_cast(static_cast(&EmbeddingKernel))); RegisterOp(OpId::kKeepQuantDecode, DeviceType::kTENSTORRENT, reinterpret_cast(static_cast(&KeepQuantDecodeKernel))); + RegisterOp(OpId::kMatmulBTQuant, DeviceType::kTENSTORRENT, + reinterpret_cast(static_cast(&MatmulBTQuantKernel))); RegisterOp(OpId::kLayerNorm, DeviceType::kTENSTORRENT, reinterpret_cast(static_cast(&LayerNormKernel))); RegisterOp(OpId::kRmsNorm, DeviceType::kTENSTORRENT, diff --git a/tests/vllm/test_gguf_keep_quant.cpp b/tests/vllm/test_gguf_keep_quant.cpp index c7f1e74487..39e4b89eaa 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -274,6 +274,40 @@ TEST_CASE("keep-quant routing respects the RUNNING DEVICE's format set (review # CHECK(!pol.keep_f16); } +TEST_CASE("keep-quant routing on TENSTORRENT admits exactly the registered decodes (KEEPQUANT W2)") { + // The P150 is discrete with no CPU fallback tier, so the same #523 shape as + // ROCm applies: the TT arm may admit only what tenstorrent_ops.cpp has a + // registered decode for. W2 registers Q4_K (MatmulBTQuantKernel, decoding + // through the W1 bit-exact chain); Q5_K/Q6_K/Q8_0 are owed by the row's W4 + // and MUST keep the pre-existing expand-bf16 residency until their kernels + // land — admitting one early throws at first forward with the model + // resident. This is the mutation red made structural: a reviewer who + // widens the TT arm past the registered set reds this case. + const std::vector shape = {4, 256}; // [out, in]: whole blocks + const auto route = [&](uint32_t ty) { + return RouteGgufTensor(/*keep_quant=*/true, /*keep_f16=*/true, + /*nvfp4_fp4=*/false, /*cpu_ref=*/false, + GgufTensorRole::kMatmulWeight, ty, shape, + vt::DeviceType::kTENSTORRENT); + }; + CHECK(route(kQ4_K) == GgufResidency::kKeepQuant); + CHECK(route(kQ8_0) == GgufResidency::kExpandBf16); // owed W4 + CHECK(route(kQ5_K) == GgufResidency::kExpandBf16); // owed W4 + CHECK(route(kQ6_K) == GgufResidency::kExpandBf16); // owed W4 + CHECK(route(kQ4_0) == GgufResidency::kExpandBf16); // no TT arm at all + CHECK(route(kQ2_K) == GgufResidency::kExpandBf16); + CHECK(route(kQ3_K) == GgufResidency::kExpandBf16); + // The loader boolean flips only when the op is registered, so a host with a + // P150 resolves keep-quant on by default; without the card the default arm + // stays false and the load is unchanged. + if (vt::OpRegistered(vt::OpId::kMatmulBTQuant, vt::DeviceType::kTENSTORRENT)) { + const GgufLoadPolicy tt = GgufLoadPolicy::FromEnv(vt::DeviceType::kTENSTORRENT); + CHECK(tt.keep_quant); + } else { + MESSAGE("kMatmulBTQuant not registered on this host; loader boolean unchecked"); + } +} + TEST_CASE("quant_repack is decided WITH the resolved device (#2406)") { // THE DEFECT. `quant_repack` selects the CIQ G7 load-time permutation of a // Q8_0 weight into the ARM i8mm `block_q8_0x4` interleave. Only the CPU diff --git a/tests/vt/test_tenstorrent_backend.cpp b/tests/vt/test_tenstorrent_backend.cpp index b33e8543f2..87b6f59968 100644 --- a/tests/vt/test_tenstorrent_backend.cpp +++ b/tests/vt/test_tenstorrent_backend.cpp @@ -5207,3 +5207,114 @@ TEST_CASE("kTENSTORRENT kKeepQuantDecode matches vt::cpu::BlockToFloat bit-exact } } } + +// W2: the keep-quant DOT. Enters through vt::MatmulBT's public dispatch (the +// entry point every model matmul helper already uses — ops.cpp:163 routes a +// block-typed [N,K] weight to kMatmulBTQuant), not through a hand-cast op +// pointer, so the test proves the routing a GGUF load actually takes. The +// oracle is the DECODE-based bf16 reference computed here: the weight decoded +// by vt::cpu::BlockToFloat (bit-exact per W1) and BOTH operands rounded to +// bf16 once (RNE), accumulated in f32 in ascending k — the same convention +// the device arm runs (decode f32 → one bf16 RNE → bf16 tile matmul). The +// envelope is the analytic bf16 operand-rounding bound, not a picked +// tolerance: rounding a and w to bf16 perturbs each product by at most +// 2^-8 relative to |a||w|, and the bf16 matmul output adds one more output +// rounding, so |device - ref| <= 1.05 * 2^-8 * (sum_k |a_k w_k| + |ref|) +// must hold elementwise; a bigger gap is a real defect, not noise. +TEST_CASE("kTENSTORRENT kMatmulBTQuant Q4_K via vt::MatmulBT matches the decode-based bf16 oracle") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kMatmulBTQuant, vt::DeviceType::kTENSTORRENT)); + + const int64_t kBlockBytes = vt::BlockBytes(vt::DType::kQ4_K); + constexpr int64_t N = 16; + std::mt19937 rng(20260906u); + + Backend& backend = vt::GetBackend(vt::DeviceType::kTENSTORRENT); + Queue q = backend.CreateQueue(); + + auto widen = [](uint16_t u) { + uint32_t bits = static_cast(u) << 16; + float f; + std::memcpy(&f, &bits, 4); + return f; + }; + + for (int64_t M : {int64_t{1}, int64_t{5}}) { // M=1 is the decode GEMV + for (int64_t nb : {int64_t{1}, int64_t{2}}) { // K spans 1 and 2 blocks + const int64_t K = nb * 256; + std::vector packed(N * nb * kBlockBytes); + for (int64_t b = 0; b < N * nb; ++b) { + uint8_t* blk = packed.data() + b * kBlockBytes; + const float d = (0.05f + 0.35f * static_cast(rng() % 64) / 64.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + const float dmin = (0.005f + 0.02f * static_cast(rng() % 32) / 32.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + const uint16_t d_bits = vt::F32ToF16(d); + const uint16_t dmin_bits = vt::F32ToF16(dmin); + std::memcpy(blk + 0, &d_bits, sizeof(d_bits)); + std::memcpy(blk + 2, &dmin_bits, sizeof(dmin_bits)); + for (int i = 0; i < 12; ++i) blk[4 + i] = static_cast(rng() & 0xFF); + for (int i = 0; i < 128; ++i) blk[16 + i] = static_cast(rng() & 0xFF); + } + std::vector a_f32(M * K); + for (auto& v : a_f32) v = (static_cast(rng() % 401) - 200.0f) / 100.0f; + + // Oracle: decode (bit-exact W1 authority) -> round ONCE to bf16 -> + // f32 accumulate in ascending k; plus the analytic envelope. + std::vector w_f32(N * K); + vt::cpu::BlockToFloat(vt::DType::kQ4_K)(packed.data(), w_f32.data(), N * K); + std::vector a_bf(M * K), w_bf(N * K); + for (size_t i = 0; i < a_f32.size(); ++i) a_bf[i] = vt::F32ToBF16(a_f32[i]); + for (size_t i = 0; i < w_f32.size(); ++i) w_bf[i] = vt::F32ToBF16(w_f32[i]); + std::vector ref(M * N), bound(M * N); + for (int64_t m = 0; m < M; ++m) + for (int64_t n = 0; n < N; ++n) { + float acc = 0.0f, mag = 0.0f; + for (int64_t k = 0; k < K; ++k) { + const float p = widen(a_bf[static_cast(m) * K + k]) * + widen(w_bf[static_cast(n) * K + k]); + acc += p; + mag += std::fabs(p); + } + ref[static_cast(m) * N + n] = acc; + bound[static_cast(m) * N + n] = + 1.05f * std::ldexp(1.0f, -8) * (mag + std::fabs(acc)); + } + + void* mem_a = backend.Alloc(M * K * sizeof(uint16_t)); + void* mem_b = backend.Alloc(packed.size()); + void* mem_o = backend.Alloc(M * N * sizeof(float)); + backend.Copy(q, mem_a, a_bf.data(), a_bf.size() * sizeof(uint16_t)); + backend.Copy(q, mem_b, packed.data(), packed.size()); + Tensor a_t = Tensor::Contiguous(mem_a, vt::DType::kBF16, + Device{vt::DeviceType::kTENSTORRENT, 0}, {M, K}); + Tensor b_t = Tensor::Contiguous(mem_b, vt::DType::kQ4_K, + Device{vt::DeviceType::kTENSTORRENT, 0}, {N, K}); + Tensor o_t = Tensor::Contiguous(mem_o, vt::DType::kF32, + Device{vt::DeviceType::kTENSTORRENT, 0}, {M, N}); + vt::MatmulBT(q, o_t, a_t, b_t); // the PUBLIC dispatch: block weight -> kMatmulBTQuant + std::vector out(M * N, 0.0f); + backend.Copy(q, out.data(), mem_o, out.size() * sizeof(float)); + backend.Free(mem_a); + backend.Free(mem_b); + backend.Free(mem_o); + + float worst = 0.0f, worst_ratio = 0.0f; + for (int64_t i = 0; i < M * N; ++i) { + const float diff = std::fabs(out[static_cast(i)] - + ref[static_cast(i)]); + worst = std::max(worst, diff); + worst_ratio = std::max(worst_ratio, diff / bound[static_cast(i)]); + CHECK(std::isfinite(out[static_cast(i)])); + CHECK_MESSAGE(diff <= bound[static_cast(i)], + "M=" << M << " K=" << K << " i=" << i << " out=" << out[i] + << " ref=" << ref[i] << " bound=" << bound[i]); + } + MESSAGE("kMatmulBTQuant M=", M, " K=", K, + ": worst_abs=", worst, " worst bound-ratio=", worst_ratio); + } + } +} From 40c69f454ac41d80ab9877d3d5a1fdc6c01d271b Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 01:33:35 +0200 Subject: [PATCH 2/2] record(BACKEND-TENSTORRENT-KEEPQUANT): W2 complete; repair the GDN anchor again (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The spec's Now moves to W2 complete and the backend matrix's KEEPQUANT row states the new position: W1 landed in #2989, W2 (dot + predicate) is on the branch, W3 owns capture-safe residency plus the e2e vehicle battery, W4 owns Q5_K/Q6_K/Q8_0, the int8 lever and the 27B arm. The vehicle GGUF was fetched under the developer's 2026-09-05 session grant and recorded in developer-preferences.md; its pin lands in docs/USAGE.md when the arm first runs (W3), per the trigger the spec body sets. The GDN row's registration citation shifted again — this row's W2 kernel insert moved tenstorrent_ops.cpp's registration block from 6687-6702 to 6755-6770 and the record gate flagged it; the citation now points at the block's new home. Anchor repaired, not the baseline — the same repair W1's record commit made, which is the cost of a keyed record citing line ranges in a file this row keeps extending. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- .agents/backend-matrix.md | 4 ++-- .agents/specs/tenstorrent-keepquant.md | 22 ++++++++++++---------- 2 files changed, 14 insertions(+), 12 deletions(-) diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 570096cd15..529e60eb7a 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -246,9 +246,9 @@ this repository. State remains `ACTIVE`; no lifecycle transition is claimed. | `BACKEND-TENSTORRENT` | Tenstorrent Blackhole (Tensix multicore, discrete PCIe, no unified memory) — thin `vt::` adapter over ttnn's existing C++ op library rather than hand-written kernels, mirroring the Metal/MLX decision (E1); vLLM has no Tenstorrent platform anywhere | vllm.cpp extension through upstream seam `platforms/interface.py:134-229` (same pattern as Metal/Vulkan) | **ACTIVE 2026-08-10.** `vt::tenstorrent::Backend` + registrar [tenstorrent_backend.cpp](../src/vt/tenstorrent/tenstorrent_backend.cpp); shared mesh-device lifecycle [tenstorrent_device.cpp](../src/vt/tenstorrent/tenstorrent_device.cpp); 17 registered ops cover OPT-125m and the Qwen3-0.6B forward (`kMatmul`, `kMatmulBT`, `kAdd`, `kRelu`, `kEmbedding`, `kLayerNorm`, `kRmsNorm`, `kSiluAndMul`, bf16/f32 casts, three RoPE forms, `kQkvSplit`, `kReshapeAndCache`, host-oracle `kPagedAttention`, `kGreedyArgmax`) [tenstorrent_ops.cpp](../src/vt/tenstorrent/tenstorrent_ops.cpp); platform allow-list selects OPT and Qwen3 [platforms/tenstorrent.cpp](../src/vllm/platforms/tenstorrent.cpp). `DeviceType::kTENSTORRENT` [device.h](../include/vt/device.h) | [test_tenstorrent_backend.cpp](../tests/vt/test_tenstorrent_backend.cpp) carries real-Blackhole op gates; [test_qwen3_paged_engine.cpp](../tests/parity/test_qwen3_paged_engine.cpp) selects Tenstorrent device-specific anchor and teacher-forced near-tie goldens. OPT-125m STRICT 6/6 passed. Qwen3 short warm smoke ran 4 tokens at about 0.28 tok/s; full 16x16 gate remains pending behind host paged attention | [tenstorrent-backend.md](specs/tenstorrent-backend.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-SPIKE` | | `BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | Child of `BACKEND-TENSTORRENT` — the owed op-level numerics evidence at the residual-RMS device boundary (`kDeviceResidualMinRows == 32`): device path does `ttnn::add`+`ttnn::rms_norm` in bf16; host/CPU path accumulates in f32. Bot-flagged on #289; never measured at the boundary. | vllm.cpp CPU oracle `RmsNormKernel` mirrors vLLM `fused_add_rms_norm` (add in model dtype, variance in f32); `src/vt/cpu/cpu_ops.cpp:371-398` | `src/vt/tenstorrent/tenstorrent_ops.cpp:1067-1117` (host/device split, `kDeviceResidualMinRows=32`) | [test_tenstorrent_backend.cpp](../tests/vt/test_tenstorrent_backend.cpp) `kRmsNorm residual: device vs CPU f32 oracle across the rows=32 boundary`: 22/22 cases on real Blackhole P150. **Measured 2026-08-11:** host path `rows<32` bit-identical to CPU (`max_abs=0`); device bf16 path `rows>=32` diverges by constant **0.0459 abs** (1.9–2.6× rel on near-zero outputs) — bf16 rounding signature, not accumulation. Decision pending the e2e golden tie-break | [tenstorrent-residual-golden.md](specs/tenstorrent-residual-golden.md) | `SPIKE` | `CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | | `BACKEND-TENSTORRENT-MISTRAL` | Child of `BACKEND-TENSTORRENT` — allowlist `MistralForCausalLM` (Mistral-7B-v0.3: GQA 32/8, head_dim 128, plain rope theta 1e6, untied lm_head, full attention) on the TT platform + device-aware SACRED gate. Mistral reuses the Qwen3-dense forward verbatim (qk-norm skipped); every op already registered. No new kernel. | vLLM `mistral.py::MistralForCausalLM(LlamaForCausalLM)` (already ported to the shared dense machinery); gate pattern mirrored from `test_qwen3_paged_engine.cpp:221-296` | `src/vllm/platforms/tenstorrent.cpp:52-54` (allowlist) + `tests/parity/test_mistral_paged_engine.cpp` (device-aware wiring + Backend Proof) | **Gate PASSED on Blackhole P150 (2026-08-12):** [test_mistral_paged_engine.cpp](../tests/parity/test_mistral_paged_engine.cpp) 16/16 prompts PASS (12/16 strict-exact, 4/16 near-tie, 0 forward-divergent), max gap **0.062 nats**, BACKEND PROOF 0 declines (kMatmul selections=256 = untied lm_head on device, kPagedAttention=8192). Goldens `our_ids_tenstorrent.npy` + `neartie_gap_mnats_tenstorrent.npy` (transformers alternative-oracle; POL-ORACLE deviation recorded, same as Qwen3-0.6B TT precedent). Exit SIGSEGV 139 is the known MeshDevice teardown crash, not a gate failure | [tenstorrent-mistral.md](specs/tenstorrent-mistral.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-MISTRAL` | -| `BACKEND-TENSTORRENT-GDN` | Child of `BACKEND-TENSTORRENT` — the GDN linear-attention op chain as native TT kernels, the hard prerequisite for the Qwen3.5/3.8 family (#1715): `kGdnPrefill`, `kGdnDecode`, `kL2Norm`, `kRmsNormGated`, `kCausalConv1dFwd`/`kCausalConv1dUpdate`, `kGdnStateGather`/`kGdnStateScatter`. The P150 is discrete, so an op miss refuses by name — the ops must land before any `Qwen3_5*` arch registration. Correctness oracle is our own CPU f32 arm (residual-golden precedent); no vLLM mirror exists for TT | Substrate: pinned tt-metal `ttnn::transformer::chunk_gated_delta_rule` (FLA chunked GDN forward, on-core recurrent state, `initial_state`/`final_state`) behind a varlen+state-permute adapter for `kGdnPrefill`; decode = rank-1 update composed from ttnn matmul+eltwise with a device shadow keyed by host pointer (`PagedKvShadow` pattern); contracts at `src/vt/ops.cpp:1823-2500`, CPU reference `src/vt/cpu/cpu_ops.cpp:1537-1740` | GDN kernels [GdnPrefillKernel :4220](../src/vt/tenstorrent/tenstorrent_ops.cpp#L4220) + [GdnDecodeKernel :5049](../src/vt/tenstorrent/tenstorrent_ops.cpp#L5049), registered [tenstorrent_ops.cpp:6687-6702](../src/vt/tenstorrent/tenstorrent_ops.cpp#L6687-L6702) (`kL2Norm`..`kGdnStateScatter`); op-level cases vs the CPU f32 oracle (T-sweep, indexed-`state_idx` forms, prefill↔decode state round-trip) | [test_tenstorrent_backend.cpp:1749-3340](../tests/vt/test_tenstorrent_backend.cpp#L1749-L3340): L2Norm, RmsNormGated, CausalConv1dFwd/Update, GdnPrefill, GdnDecode, prefill↔decode round-trip, StateGather/Scatter, edge shapes — every op family vs the CPU f32 oracle. W1 `34fde3502` (prefill set) + W2 `c85af0aaf` (decode+state-I/O set) landed, both fresh-review PASS; production-reached via the `Qwen3_5*` wiring row (BACKEND-TENSTORRENT-QWEN35: allow-list, e2e sacred pair 16/16 STRICT both legs); capture compatibility unmeasured behind #1625; closure [parity-ledger.md#L946](parity-ledger.md#L946) | [tenstorrent-gdn.md](specs/tenstorrent-gdn.md) | `DONE` | `c85af0aaf` | +| `BACKEND-TENSTORRENT-GDN` | Child of `BACKEND-TENSTORRENT` — the GDN linear-attention op chain as native TT kernels, the hard prerequisite for the Qwen3.5/3.8 family (#1715): `kGdnPrefill`, `kGdnDecode`, `kL2Norm`, `kRmsNormGated`, `kCausalConv1dFwd`/`kCausalConv1dUpdate`, `kGdnStateGather`/`kGdnStateScatter`. The P150 is discrete, so an op miss refuses by name — the ops must land before any `Qwen3_5*` arch registration. Correctness oracle is our own CPU f32 arm (residual-golden precedent); no vLLM mirror exists for TT | Substrate: pinned tt-metal `ttnn::transformer::chunk_gated_delta_rule` (FLA chunked GDN forward, on-core recurrent state, `initial_state`/`final_state`) behind a varlen+state-permute adapter for `kGdnPrefill`; decode = rank-1 update composed from ttnn matmul+eltwise with a device shadow keyed by host pointer (`PagedKvShadow` pattern); contracts at `src/vt/ops.cpp:1823-2500`, CPU reference `src/vt/cpu/cpu_ops.cpp:1537-1740` | GDN kernels [GdnPrefillKernel :4220](../src/vt/tenstorrent/tenstorrent_ops.cpp#L4220) + [GdnDecodeKernel :5049](../src/vt/tenstorrent/tenstorrent_ops.cpp#L5049), registered [tenstorrent_ops.cpp:6755-6770](../src/vt/tenstorrent/tenstorrent_ops.cpp#L6755-L6770) (`kL2Norm`..`kGdnStateScatter`); op-level cases vs the CPU f32 oracle (T-sweep, indexed-`state_idx` forms, prefill↔decode state round-trip) | [test_tenstorrent_backend.cpp:1749-3340](../tests/vt/test_tenstorrent_backend.cpp#L1749-L3340): L2Norm, RmsNormGated, CausalConv1dFwd/Update, GdnPrefill, GdnDecode, prefill↔decode round-trip, StateGather/Scatter, edge shapes — every op family vs the CPU f32 oracle. W1 `34fde3502` (prefill set) + W2 `c85af0aaf` (decode+state-I/O set) landed, both fresh-review PASS; production-reached via the `Qwen3_5*` wiring row (BACKEND-TENSTORRENT-QWEN35: allow-list, e2e sacred pair 16/16 STRICT both legs); capture compatibility unmeasured behind #1625; closure [parity-ledger.md#L946](parity-ledger.md#L946) | [tenstorrent-gdn.md](specs/tenstorrent-gdn.md) | `DONE` | `c85af0aaf` | | `BACKEND-TENSTORRENT-GDN-DEVICE-PURE` | Child of `BACKEND-TENSTORRENT-GDN` — make the decode-side GDN ops (`CausalConv1dUpdateKernel`, `GdnDecodeKernel`) device-resident so a tt-metal trace capture admits them, unblocking the Qwen3.5-0.8B captured arm, the last capture-blocked family with a committed eager pair (#2907, owed from #2812) | Move the per-call host orchestration on-device: resident inputs from the graph's producer ops, device-side indexed state update (baked slot addressing re-primed by the recapture cadence, or `kGdnStateGather`/`kGdnStateScatter` indirection — decided on recapture-cost measurement), conv two-views coherent on device, token readback outside the captured span | [CausalConv1dUpdateKernel :5013](../src/vt/tenstorrent/tenstorrent_ops.cpp#L5013), [GdnDecodeKernel :5225](../src/vt/tenstorrent/tenstorrent_ops.cpp#L5225) (EnsureHost ×5, ReadIdxHost, UploadTensor ×5+), q35 harness pair selection [test_qwen35_paged_engine.cpp:233](../tests/parity/test_qwen35_paged_engine.cpp#L233), fatal repro `fd_mesh_command_queue.cpp:760` 2/2 | Red-first: the #2812 loud-skip opt-in cells run captured; captured dump ×2 byte-identity with reset between; teacher-forced pair vs the transformers oracle inside the eager band; Qwen3.5 joins `DecodeCaptureDefaultArch`, ambient adjudicates CAPTURED, env=0 eager; tamper + arch-deletion mutations red; gate arm selection [test_qwen35_paged_engine.cpp:301-330](../tests/parity/test_qwen35_paged_engine.cpp#L301-L330); closure [parity-ledger.md:947](parity-ledger.md#L947) | [tenstorrent-gdn-device-pure.md](specs/tenstorrent-gdn-device-pure.md) | `DONE` | `bcade48d6` | -| `BACKEND-TENSTORRENT-KEEPQUANT` | Child of `BACKEND-TENSTORRENT` — keep-quant dense dot on TT-Metal ([#2959](https://github.com/mudler/vllm.cpp/issues/2959)): GGUF k-quant arms refuse on the P150 today — `DeviceKeepQuantSupported` has no `kTENSTORRENT` arm (`gguf_keep_quant.cpp:136-148`) and no TT quant kernel exists — while the smallest Qwen3.8 artifact that fits the card is the dense 27B Q4_K_M (17.1 GB; bf16 is 53.8 GB) and a quantized Qwen3.5 (0.8B, ~0.5 GB) is the test vehicle | ggml k-quant block formats via our reader, pinned bit-exact vs llama.cpp `b10451` (#2240/#2260 lineage); no vLLM mirror for TT (secondary substrate, deviation by design); substrate surveyed 2026-09-05: tt-metal has NO packed-weight matmul (`ttnn` quantization is per-tensor activation dquant only), so the kernel is ours — resident blocks, on-core decode to bf16 tiles through f32 (decode bit-exact vs `BlockToFloat`), the dot reusing the `kMatmulBT` tile path, device-bf16 band at the residual-golden boundary, the int8 dot as the named next lever | OWED W1: the Q4_K block-decode device path + provider registration + the predicate arm admitting exactly the registered encodings (never wider — the GLM-5.3 W10 lesson) | OWED W1: decode bit-exact vs `vt::cpu::BlockToFloat` across a shape sweep (red-first, op-level suite pattern); capture dump ×2 byte-identity (#2907 discipline); e2e 16-prompt battery on the P150 vs the bf16 arm's committed pair (STRICT or inside the ≤500-mnat near-tie band) | [tenstorrent-keepquant.md](specs/tenstorrent-keepquant.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-KEEPQUANT` — implementer on the row branch; spec commit on `row/BACKEND-TENSTORRENT-KEEPQUANT`, claim in [.agents/claims/](claims/CLAIM-BACKEND-TENSTORRENT-KEEPQUANT.md) | +| `BACKEND-TENSTORRENT-KEEPQUANT` | Child of `BACKEND-TENSTORRENT` — keep-quant dense dot on TT-Metal ([#2959](https://github.com/mudler/vllm.cpp/issues/2959)): GGUF k-quant arms refuse on the P150 until this row lands — W1 (#2989, open) lands the Q4_K decode, W2 (branch) lands the dot and the `kTENSTORRENT` predicate arm (`gguf_keep_quant.cpp:136-148`) admitting exactly `{Q4_K}` — while the smallest Qwen3.8 artifact that fits the card is the dense 27B Q4_K_M (17.1 GB; bf16 is 53.8 GB) and a quantized Qwen3.5 (0.8B, ~0.5 GB) is the test vehicle | ggml k-quant block formats via our reader, pinned bit-exact vs llama.cpp `b10451` (#2240/#2260 lineage); no vLLM mirror for TT (secondary substrate, deviation by design); substrate surveyed 2026-09-05: tt-metal has NO packed-weight matmul (`ttnn` quantization is per-tensor activation dquant only), so the kernel is ours — resident blocks, on-core decode to bf16 tiles through f32 (decode bit-exact vs `BlockToFloat`), the dot reusing the `kMatmulBT` tile path, device-bf16 band at the residual-golden boundary, the int8 dot as the named next lever | W1 LANDED (#2989): the Q4_K block-decode device path. W2 ON BRANCH: the dot (`MatmulBTQuantKernel`, decode → one bf16 RNE → `kMatmulBT` tile matmul, reached through `vt::MatmulBT`'s block-weight dispatch) + the predicate arm admitting exactly the registered set (never wider — the GLM-5.3 W10 lesson). OWED W3: capture-safe residency + the e2e vehicle. OWED W4: Q5_K/Q6_K/Q8_0, the int8 lever, the 27B arm | W1 LANDED (#2989): decode bit-exact vs `vt::cpu::BlockToFloat` across a shape sweep (red-first, op-level suite pattern). W2 ON BRANCH: the dot vs the decode-based bf16 oracle inside the analytic operand-rounding envelope + the predicate routing set (both red-first). OWED W3: capture dump ×2 byte-identity (#2907 discipline); e2e 16-prompt battery on the P150 vs the bf16 arm's committed pair (STRICT or inside the ≤500-mnat near-tie band) | [tenstorrent-keepquant.md](specs/tenstorrent-keepquant.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-KEEPQUANT` — implementer on the row branch; spec commit on `row/BACKEND-TENSTORRENT-KEEPQUANT`, claim in [.agents/claims/](claims/CLAIM-BACKEND-TENSTORRENT-KEEPQUANT.md) | | `BACKEND-TENSTORRENT-QWEN35` | Child of `BACKEND-TENSTORRENT` — the wiring row: `Qwen3_5ForConditionalGeneration` (dense text GDN hybrid) on the TT allow-list, the op delta the family forward refuses by name (`kGdnPostConv`, `kSigmoidGateBf16`, `kAttnQkNormRopeGate`, `kAttnQkNormRope` — pinned empirically by a W0 refusal sweep), and the first e2e gate. Makes the GDN row's ops production-reached (#1715 stays open until the family runs; GDN row lifecycle moves in the same change this lands) | Substrate: the GDN row's 8 kernels + the 27-op TT registry; e2e mirrors the ratified Mistral TT golden treatment (`VT_DUMP_IDS` bootstrap → `scripts/qwen3-neartie-gap-transformers.py` teacher-forced near-tie gaps → committed device-golden pair, POL-ORACLE deviation); op-level oracle is the CPU f32 arm. Capacity: `Qwen/Qwen3.5-0.8B` bf16 ≈1.6 GB fits (proven envelope Mistral-7B ≈14.5 GB); 27B bf16 ≈53.8 GB and GGUF k-quant arms (no TT kernels) refused by name, owed | `src/vllm/platforms/tenstorrent.cpp` (allow-list, lands last) + `src/vt/tenstorrent/tenstorrent_ops.cpp` (op delta) + `tests/parity/test_qwen35_paged_engine.cpp` (TT arm) | Spec committed 2026-08-23; W0 sweep, W1 op delta, W2 e2e, W3 GDN-reviewer leftovers (d2h counter completeness, `conv_transposed` fast-path check) owed | [tenstorrent-qwen35.md](specs/tenstorrent-qwen35.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-QWEN35` | | `BACKEND-TENSTORRENT-TRACE-RUNNER` | Child of `BACKEND-TENSTORRENT` — wire the landed graph-capture foundation (#354 / `59568772`) into a capturable forward region. Handoff §8/§9 "not done". **Decision (2026-08-13): NO-GO for pure T=1 decode capture — MEASURED, not assumed.** | CUDA is the ONLY backend with `SupportsGraphCapture()==true` (`cuda_backend.cu:184-240`); Metal/Vulkan `false`. Shared decode-graph framework `Qwen3DenseDecodeGraph` (qwen3.cpp:489, used by Qwen3/Mistral/Llama/InternLM2) gated on `support_static_graph_mode()` which TT does NOT override (base `false`) | `src/vt/tenstorrent/tenstorrent_backend.cpp:70-76` (capture surface, landed) + `tenstorrent_ops.cpp` `Trace*` (landed); NOT wired into a dense forward | **Measured on Blackhole P150 (2026-08-13):** Q1 no host-free region at T=1; Q2 all-device-at-T=1 costs 12.5→10.7 tok/s; Q2b capture attempt with both overrides flipped → ttnn **`TT_FATAL: Reads are not supported during trace capture`** (backtrace through `to_vector`), `0 replays`. The T=1 forward does device→host readbacks that ttnn trace prohibits; flipping the two thresholds is insufficient. Requires a host-free `ForwardLayers` redesign, not a threshold tweak. Next: prefill capture (separate row) must first audit its `to_vector` readbacks | [tenstorrent-trace-runner.md](specs/tenstorrent-trace-runner.md) | `SPIKE` (decision record complete) | `CLAIM-BACKEND-TENSTORRENT-TRACE-RUNNER` | | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | Child of `BACKEND-TENSTORRENT` — make the per-decode-layer forward region host-free (zero `to_vector` readbacks) so mesh-trace capture can run. Decomposes into R1 RmsNorm+RoPE all-device, R2 QkvSplit+ReshapeAndCache device, R3 PA decode device-resident metadata, R4 capture wire+measure. Prerequisite revealed by the trace-runner spike (capture aborts on `to_vector`). | CUDA decode-graph capture contract (`cuda_backend.cu:184-197`): captured region is async, no host sync, no malloc, fixed ptrs. TT must match: no `to_vector` between Begin/EndCapture | `src/vt/tenstorrent/tenstorrent_ops.cpp` (RmsNorm:1067, PreferDeviceRope:1344, QkvSplit:1460, ReshapeAndCache:1527, PagedAttention:2009) | **R1-R3b MEASURED on P150** (env-gated `VT_TT_HOST_FREE_DECODE`, inert by default; 23/23 TT tests incl. a dedicated default-path inertness guard, M1-mutation-proved): R1 threshold flip landed; R2 device->device copy (`CopyDeviceDeviceIfCapture`, ttnn::copy+empty) landed; R3 program-cache warm (`enable_program_cache` + eager-warm) landed; R3b device zero-fill (`MemsetDeviceIfCapture`) landed. Capture enters the forward and reaches layer ops (CastBf16/RmsNorm fire). Remaining item-5 blocker = per-op enqueue_write; answer = persistent device tensors + before-replay populate (tt-metal vLLM plugin design). Full blocker map + architecture in [tenstorrent-host-free-r1.md](specs/tenstorrent-host-free-r1.md). **R5 (2026-08-21, #1604): host-free decode is the DEFAULT** — `HostFreeDecodeEnabled()` centralizes the polarity (exact `0` = pre-flip opt-out), both device golden pairs re-dumped and re-adjudicated under the new default (Qwen3 max 375 mnats, Mistral max 250 mnats, 0 outside top-K), both paged-engine gates 16/16 green (125/125 + 128/128 assertions), default leg 10.94-11.06 tok/s vs 5.34 opt-out (2.1x, Qwen3-0.6B b1). Capture declined by default: multi-request captured hangs (#1625); TT async scheduling stays off, no `SupportsAsyncSampledTokenReadback` override (#1627) | [tenstorrent-host-free-forward.md](specs/tenstorrent-host-free-forward.md) | `ACTIVE` (R1-R3b + R5 default flip gated on `row/BACKEND-TENSTORRENT-HOST-FREE-1604`; capture declined per #1625, TT async readback owed per #1627) | `CLAIM-BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | diff --git a/.agents/specs/tenstorrent-keepquant.md b/.agents/specs/tenstorrent-keepquant.md index 2ed1f40198..6cd4c5480d 100644 --- a/.agents/specs/tenstorrent-keepquant.md +++ b/.agents/specs/tenstorrent-keepquant.md @@ -200,13 +200,15 @@ to make a failure pass. ## Now -`ACTIVE`, 2026-09-05. W1 complete: the Q4_K block-decode device kernel -(`KeepQuantDecodeKernel`, `tenstorrent_ops.cpp`) is bit-exact against -`vt::cpu::BlockToFloat` on the full sweep (rows {1,3,17} x nb {1,2,16}, -memcmp-level, 53/53 cases green), with the signed-zero repair all-TILE: -`ttnn::where` demands a TILE predicate and silently writes only its first 16 -elements when a TILE predicate is mixed with ROW_MAJOR branches; the f32 -{0,1}-mask multiply pred replaces `logical_and` (rank promotion of a broadcast -predicate); m1 is repaired at its final {B,8,1} shape. Next: W2, the dot -provider + `kTENSTORRENT` predicate arm + registration, mutation-red. The -vehicle pin stays owed until the arm first runs (W3). +`ACTIVE`, 2026-09-06. W1 complete (#2989, open). W2 complete on the row +branch: the dot (`MatmulBTQuantKernel` — `DecodeQ4KBlocksF32` factored out of +the W1 kernel, one bf16 RNE, the `kMatmulBT` tile matmul) is reached through +`vt::MatmulBT`'s block-weight dispatch and sits inside the analytic +operand-rounding envelope (bound ratios 0.28-0.53, M=1 GEMV included); the +`kTENSTORRENT` predicate arm admits exactly `{Q4_K}` and the routing test +reds any widening past the registered set. Both red-first: the registration +REQUIRE and the six wrongly-admitted encodings reded before the +implementation. Next: W3, the capture leg (dump ×2 byte-identity, capture- +safe staging for the per-call decode upload) + the e2e vehicle battery on +the P150 (vehicle fetched and hashed). W4 owed: Q5_K/Q6_K/Q8_0, the int8 +lever, the 27B arm.