From c4e67e8ad2b463207e7e301e1506a870b2c7b94b Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 02:09:53 +0200 Subject: [PATCH 1/5] record(BACKEND-TENSTORRENT-KEEPQUANT): the q4km vehicle is mixed-quant; W3 takes the decode set (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The fetched Q4_K_M vehicle is not a uniform Q4_K artifact: token_embd is Q6_K (tied head), attn_qkv/ssm_out are Q5_K, ssm_alpha/ssm_beta are Q8_0, and ffn_down/attn_v mix Q6_K into Q4_K layers — read from the GGUF tensor headers of the hashed bytes. Under the W2 predicate a first forward refuses by name, which is the predicate working. The spec's W3/W4 split assumed a pure-Q4_K vehicle, so the e2e battery it schedules could never run; this reconciles the record with the evidence instead of re-scoping in a commit body: the Q5_K/Q6_K/Q8_0 decodes and the predicate widening move from W4 into W3, and W4 keeps the int8-dot lever and the 27B arm. The amendment also commits the W3 staging design the survey produced: stage the i32 word form once per weight through the existing persistent-shadow machinery (36 i32 words are exactly 144 packed bytes — zero expansion, residency thesis intact), decode on-core from the resident words every call, CHECK on a mid-capture miss, and prove it with a capture dump x2 byte-identity leg that reds against W2's per-call host round trip. The Q6_K embedding gather-vs-staged-table decision and the teacher-forcing oracle (transformers on the dequantized artifact, never the bf16 checkpoint) are named as implementation-time decisions with their constraints. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- .agents/specs/tenstorrent-keepquant.md | 62 +++++++++++++++++++++++--- 1 file changed, 56 insertions(+), 6 deletions(-) diff --git a/.agents/specs/tenstorrent-keepquant.md b/.agents/specs/tenstorrent-keepquant.md index 6cd4c5480..bb2b490f1 100644 --- a/.agents/specs/tenstorrent-keepquant.md +++ b/.agents/specs/tenstorrent-keepquant.md @@ -127,6 +127,50 @@ the rows>=32 boundary). | — (local; the predicate is the GLM-5.3 W10 lesson made structural) | OWED W2: the `kTENSTORRENT` arm of `DeviceKeepQuantSupported` (`src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:136-148`) | Admits exactly the encodings whose device kernels are registered; never wider, because the `default:` arm's CPU fallback does not exist on a discrete P150. | | — (local; the #2812/#2907 capture discipline applied to a new weight path) | OWED W3: capture-safe residency of the decode path | No mid-capture host download; CHECK-before-download ordering; capture dump ×2 byte-identity as the gate. | +## The vehicle falsified the W3/W4 split (found 2026-09-06) + +The Q4_K_M artifact is mixed-quant, not uniform Q4_K. The tensor histogram +of the fetched vehicle (unsloth/Qwen3.5-0.8B-GGUF at `6ab46149`, sha256 +`bd258782...`): `token_embd.weight` **Q6_K** (the tied LM head reads the +same tensor), `attn_qkv` and `ssm_out` **Q5_K** (18 each), `ssm_alpha` and +`ssm_beta` **Q8_0** (18 each), `ffn_down` Q6_K x12 / Q4_K x12, `attn_v` +Q6_K x4 / Q4_K x2, everything else Q4_K or F32. A first forward on TT under +the W2 predicate refuses by name — correctly. Consequence: the e2e vehicle +battery needs the Q5_K/Q6_K/Q8_0 decodes first. The decode set moves from +W4 into W3; W4 keeps the int8-dot lever and the 27B arm. The predicate +widens `{Q4_K}` to exactly the registered four in the same change as the +kernels — never before. + +The Q6_K token embedding is a design point, not a footnote: kEmbedding on +TT must reach a Q6_K weight either by a device gather that decodes rows +on-core, or by staging the decoded table (248320 x 1024 bf16, ~0.5 GiB — +bounded here, unlike the 27B expansion the design thesis forbids). Decide +at implementation against the CPU arm's behavior. + +## W3 capture-safe staging (design, from the W2 survey) + +`DecodeQ4KBlocksF32` today EnsureHosts the packed tensor, repacks to i32 +words on host, and `from_vector`-uploads per call — a per-call host round +trip that inside a capture is the #2812 class. The fix keeps the residency +thesis (packed stays the resident master; a decoded-bf16 cache would triple +the 0.8B and blow the 27B budget): stage the **i32 word form once per +weight** through the existing persistent-shadow machinery +(`EnsureMatmulWeightDevice`'s keyed-slot pattern, built for the view-staging +fatality — 36 i32 words are exactly 144 packed bytes, zero expansion), and +run the on-core decode chain from the resident words every call. Captured +replay recomputes deterministically from fixed device bytes; a cache miss +during capture CHECKs ("warm the keep-quant arm eagerly first"), the +zero-cache precedent. Red-first leg: with W2 code the captured vehicle run +traces staging writes during capture and must fail the dump x2 byte-identity +leg; after, zero writes during capture and dumps byte-identical across two +reset runs. + +The gap golden for the q4km arm teacher-forces `transformers` on the +DEQUANTIZED artifact — `from_pretrained(gguf_file=...)` if the pinned +transformers 5.14.1 parses qwen35 GGUF, otherwise our own bit-exact decoder +(W1-proven) writing a safetensors dir first. Never teacher-force against +the bf16 safetensors checkpoint: those logits are a different model's. + ## Tests to port vLLM has no TT backend and no keep-quant test to port; the oracle chain is the @@ -155,10 +199,14 @@ a scratch copy), capture dump ×2 byte-identity with reset between, and the - **W2**: the dot provider + `kTENSTORRENT` predicate arm + registration; predicate mutation red (admit one encoding too many → first-forward throw reproduced in a scratch copy). -- **W3**: capture leg (dump ×2 byte-identity with reset between) + the e2e - vehicle battery on the P150 vs the bf16 pair. -- **W4** (owed): Q5_K / Q6_K / Q8_0; the int8-dot perf lever; the 27B Q4_K_M - arm as the first qwen3.8 artifact on TT. +- **W3**: Q5_K / Q6_K / Q8_0 decodes (bit-exact vs `BlockToFloat` sweeps, + the W1 template generalizes); the predicate widening to exactly the + registered four in the same change; capture-safe resident-word staging; + the capture leg (dump ×2 byte-identity with reset between); the e2e + vehicle battery on the P150 vs its own captured pair under the committed + near-tie conventions. +- **W4** (owed): the int8-dot perf lever; the 27B Q4_K_M arm as the first + qwen3.8 artifact on TT. ## Risks @@ -210,5 +258,7 @@ 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. +the P150 (vehicle fetched and hashed). AMENDED 2026-09-06: the vehicle is +mixed-quant, so W3 now carries the Q5_K/Q6_K/Q8_0 decodes and the predicate +widening before the capture leg and the e2e battery (see the falsification +section). W4 owed: the int8 lever, the 27B arm. From f2bf5fec628218280cdb344d74862a9f8883b72d Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 14:41:51 +0200 Subject: [PATCH 2/5] feat(BACKEND-TENSTORRENT-KEEPQUANT): Q5_K/Q6_K/Q8_0 decode chains, MoeSiluMul, host-side twin (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The q4km vehicle is mixed-quant (token_embd Q6_K, attn_qkv/ssm_out Q5_K, ssm_alpha/ssm_beta Q8_0), so the decode set the row's falsification moved into W3 lands here: bit-exact DecodeKeepQuantBlocksF32 chains for Q5_K, Q6_K, and Q8_0 beside the W1 Q4_K chain, each pinned against vt::cpu::BlockToFloat across the shape sweep. The weight twin is built host-side, not on device. W2's device-side twin decode OOM-fatalled the vehicle's first forward with a 4,068,474,880 B ttnn::where (addr2line: the signed-zero repair inside DecodeKeepQuantBlocksF32). The kernel now decodes blocks on the host (vt::cpu::BlockToFloat), pre-rounds each f32 through BF16ToF32(F32ToBF16(x)) so the uploaded values are exactly bf16, and uploads once via from_vector — ttnn converts to BFLOAT16 Tile during the upload, the EnsureEmbedTableDevice precedent — memoized per host buffer in DecodedWeightShadow and dropped in UnregisterHostBuffer. vt::MoeSiluMul lands because the GGUF dense MLP's split gate/up path dispatches it (qwen3_5.cpp:7610); this vehicle never dispatches kSiluAndMul, and the e2e ladder's backend-proof list now asserts kMoeSiluMul for the keep-quant arm instead. Capture safety stays structural: EnsureKeepQuantWords refuses a capture-time arrival by name, and the staging counter (KeepQuantCaptureStagingWrites) proves the captured e2e stages zero word uploads — the accessors live in tenstorrent_device.h with inline no-ops for device-agnostic TUs. The i32 word-shadow machinery and kKeepQuantDecode stay registered and tested, but no production path reads them now; the residency reconciliation is recorded owed in the spec. Red-first: per-encoding decode-equality pins, the MoeSiluMul golden case, and the stages-zero-words-during-capture case each reded before their implementation. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- src/vt/tenstorrent/tenstorrent_device.h | 17 + src/vt/tenstorrent/tenstorrent_ops.cpp | 969 ++++++++++++++++++------ tests/vt/test_tenstorrent_backend.cpp | 436 +++++++++++ 3 files changed, 1186 insertions(+), 236 deletions(-) diff --git a/src/vt/tenstorrent/tenstorrent_device.h b/src/vt/tenstorrent/tenstorrent_device.h index 28f323631..9cb69a39a 100644 --- a/src/vt/tenstorrent/tenstorrent_device.h +++ b/src/vt/tenstorrent/tenstorrent_device.h @@ -162,6 +162,23 @@ bool MemsetDeviceIfCapture(void* p, int value, size_t bytes); bool MemsetDeviceFill(void* p, int value, size_t bytes); bool CopyDeviceDeviceIfResident(void* dst, const void* src, size_t bytes); +// BACKEND-TENSTORRENT-KEEPQUANT W3 capture-safety probe: how many staging +// writes (host repack + from_vector upload) the keep-quant decode path has +// performed while a trace capture was active. The staged arm performs ZERO — +// the i32 word shadow is resident before capture and the per-call decode +// recomputes from it — so any positive count inside a captured run is the +// #2812 class (a captured graph reading a buffer its replay cannot refresh). +// The reset hook exists for the op-level red-first test only. Defined only +// when the backend is built (tenstorrent_ops.cpp); inline no-ops otherwise, +// the WarmPagedKvShadow pattern, so device-agnostic TUs may call them. +#ifdef VLLM_CPP_TENSTORRENT +int64_t KeepQuantCaptureStagingWrites(); +void ResetKeepQuantCaptureStagingWritesForTest(); +#else +inline int64_t KeepQuantCaptureStagingWrites() { return 0; } +inline void ResetKeepQuantCaptureStagingWritesForTest() {} +#endif + // ITEM 5 (rope): driver-side warm hook — populate the persistent device // cos/sin tensors for the step's positions BEFORE BeginCapture (the // SizeSlot::Refresh slot), so the captured rope cache-HITs. No-op unless diff --git a/src/vt/tenstorrent/tenstorrent_ops.cpp b/src/vt/tenstorrent/tenstorrent_ops.cpp index 2454d8afe..d20cf4b35 100644 --- a/src/vt/tenstorrent/tenstorrent_ops.cpp +++ b/src/vt/tenstorrent/tenstorrent_ops.cpp @@ -21,6 +21,7 @@ #include "vt/cpu/cpu_threadpool.h" #include "vt/dtype.h" #include "vt/ops.h" +#include "vt/quant.h" // This OBJECT library is not the `vllm` target, so it does not inherit the // PUBLIC VLLM_CPP_TENSTORRENT define. Force the real declarations; the // header's inline no-ops are only for CPU/Vulkan/Windows TUs. @@ -192,6 +193,18 @@ bool& tt_capture_active() { } } // namespace +// KEEPQUANT W3 capture-safety probe (tenstorrent_device.h): staging writes the +// keep-quant decode performed while a capture was active. The staged arm must +// hold this at zero across a captured run (the W3 red-first test reads it). +// The counter lives at internal linkage; the accessors below the file's +// anonymous namespace give the parity/test TUs the external surface. +namespace { +std::atomic& KeepQuantCaptureStagingWritesCounter() { + static std::atomic* c = new std::atomic(0); // never destroyed (#1486) + return *c; +} +} // namespace + // ITEM 5 (rope): persistent device cos/sin (expanded per head), built OUTSIDE // capture and ttnn::copy'd in-region — the UploadRows in RopeApplyDeviceNeox // was the enqueue_write that killed capture at mid-layer-0. The cache is @@ -270,9 +283,9 @@ std::string ZeroCacheKey(const ttnn::Shape& shape, ttnn::DataType dt, } } // namespace -ttnn::Tensor ZeroCacheGet(const ttnn::Tensor& like, MeshDevice& device) { - const std::string key = ZeroCacheKey(like.logical_shape(), like.dtype(), - like.layout()); +ttnn::Tensor ZeroCacheGet(const ttnn::Shape& shape, ttnn::DataType dt, + ttnn::Layout lt, MeshDevice& device) { + const std::string key = ZeroCacheKey(shape, dt, lt); std::lock_guard g(ZeroCacheMutex()); auto& c = ZeroCache(); auto it = c.find(key); @@ -280,13 +293,16 @@ ttnn::Tensor ZeroCacheGet(const ttnn::Tensor& like, MeshDevice& device) { VT_CHECK(!tt_capture_active(), "tenstorrent: zero-cache miss during capture — warm the " "host-free path eagerly (VT_TT_HOST_FREE_DECODE warmup) first"); - it = c.emplace(key, ttnn::zeros(like.logical_shape(), like.dtype(), - like.layout(), std::ref(device))) - .first; + it = c.emplace(key, ttnn::zeros(shape, dt, lt, std::ref(device))).first; } return it->second; } +ttnn::Tensor ZeroCacheGet(const ttnn::Tensor& like, MeshDevice& device) { + return ZeroCacheGet(like.logical_shape(), like.dtype(), like.layout(), + device); +} + void ZeroCachePrime(const ttnn::Shape& shape, ttnn::DataType dt, ttnn::Layout lt, MeshDevice& device) { const std::string key = ZeroCacheKey(shape, dt, lt); @@ -1713,17 +1729,14 @@ void EmbeddingKernel(Queue&, Tensor& out, const Tensor& table, const Tensor& ids if (!in_place) CommitDevice2D(out, std::move(dev_out)); } -// kKeepQuantDecode: Q4_K block decode as a device compute chain — the packed -// stream is decoded on the device, not staged through a host decode. Contract -// is DequantQ4_K (cpu_quant_dequant.cpp:167, GetScaleMinK4 just above it; -// llama.cpp dequantize_row_q4_K): per block_q4_K {f16 d; f16 dmin; u8 -// scales[12]; u8 qs[128]}, y[g*32+l] = d1*(nib) - m1 with d1 = d*sc, m1 = -// dmin*mm, sc/mm from GetScaleMinK4(is, scales), groups 0..7 over the 8 -// 32-nibble lanes (low nibbles first, then high nibbles, of the same 32 -// bytes). The eager arm mirrors kEmbedding's staging: Alloc hands back a -// host-mapped pointer, so the packed bytes are read from t.data, assembled to -// little-endian u32 words (the GGUF stream order ReadF16 already assumes), -// and uploaded as one INT32 tensor. Every unpack step is device work: +// kKeepQuantDecode: packed-block decode as a device compute chain — the packed +// stream is decoded on the device, not staged through a host decode. Contracts +// are DequantQ4_K/DequantQ5_K/DequantQ6_K/DequantQ8_0 (cpu_quant_dequant.cpp; +// llama.cpp dequantize_row_q4_K:1471, q5_K:1673, q6_K:1881, q8_0:495). Q4_K: +// per block_q4_K {f16 d; f16 dmin; u8 scales[12]; u8 qs[128]}, y[g*32+l] = +// d1*(nib) - m1 with d1 = d*sc, m1 = dmin*mm, sc/mm from GetScaleMinK4(is, +// scales), groups 0..7 over the 8 32-nibble lanes (low nibbles first, then +// high nibbles, of the same 32 bytes). Every unpack step is device work: // bitwise masks/shifts for the nibbles, scale bytes, and f16 halves; // reshape/permute only reorder lanes; d/dmin widen f16->f32 exactly via the // integer bit-construction chain of vt::F16ToF32 (loader @@ -1735,54 +1748,210 @@ void EmbeddingKernel(Queue&, Tensor& out, const Tensor& table, const Tensor& ids // every finite pattern is bit-exact. nibbles and scale factors typecast to // f32 exactly (values <= 255); d1/m1 and the final y = d1*nib - m1 are // separate f32 ttnn ops (no FMA), the same IEEE order as the host under -// -ffp-contract=off. Trace capture (BACKEND-TENSTORRENT-KEEPQUANT W3) owns a -// device-authoritative staging arm; this eager arm reads the host-mapped -// staging the same way kEmbedding does. -// A device tensor filled with -0.0f. from_vector is a host memcpy, so the -// negative-zero bit pattern survives staging intact — no SFPU op between the -// 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. +// -ffp-contract=off. // -// 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); +// THE SHARED DECODE (KEEPQUANT W2, generalized W3): the packed-block -> f32 +// chains, each returning the repaired f32 {rows, nb*elems} in ROW_MAJOR. +// KeepQuantDecode commits it to the host-visible output; the quant matmul +// (MatmulBTQuantKernel below) consumes it as the weight operand. One decode +// per encoding, two consumers, one numerics authority: +// Q4_K 144B/256 and Q5_K 176B/256: {f16 d; f16 dmin; u8 scales[12]; ...} +// — Q5_K inserts qh[32] between scales and qs (the 5th bit, +16) and its +// tail is Q4_K's verbatim; Q6_K 210B/256: {u8 ql[128]; u8 qh[64]; i8 +// scales[16]; f16 d} — SIGNED 8-bit scales, 16 sub-blocks of 16, no min +// term, q = nib6 - 32; Q8_0 34B/32: {f16 d; i8 qs[32]} — plain int8 times +// the f16 scale. 210 and 34 are not multiples of four, so those blocks pad +// to 53 / 9 words (2 pad bytes at the stream tail, zero-filled). +// +// STAGING (W3): the packed stream is staged as a resident i32 word shadow ONCE +// per weight (EnsureKeepQuantWords below, the EnsureMatmulWeightDevice +// persistent-shadow pattern), so the per-call decode runs entirely on-core +// from the resident words — no host repack, no from_vector upload. The -0.0f +// signed-zero repair constant is built by BIT ops from the zero-cache's +0 +// (multiply canonicalizes -0 to +0 on the SFPU, so the sign bit is set below +// the float domain), because a per-call from_vector of the constant would be +// exactly the captured-graph write this wave removes. + + +// ---- KEEPQUANT W3: the resident i32 word shadow ----------------------------- +// Words per staged packed block: 144B -> 36 and 176B -> 44 (multiples of +// four); 210B -> 53 and 34B -> 9 for Q6_K/Q8_0, whose 2-byte tails the +// staging zero-fills (the GGUF stream never carries them, and every byte +// position the decoders read is inside the true block bytes, so the pad +// never reaches an output). +int KeepQuantWordsPerBlock(DType enc) { + switch (enc) { + case DType::kQ4_K: return 36; + case DType::kQ5_K: return 44; + case DType::kQ6_K: return 53; + case DType::kQ8_0: return 9; + default: return 0; + } +} + +struct KeepQuantWordShadow { + ttnn::Tensor words; // {B, wpb} INT32 ROW_MAJOR — the packed stream, word-staged + int64_t rows = 0; + int64_t nb = 0; + int wpb = 0; +}; +std::mutex& KeepQuantWordMutex() { + static std::mutex m; + return m; +} +// Keyed by the packed tensor's host base pointer — the EnsureMatmulWeightDevice +// persistent-shadow pattern. An interior view carries its own pointer, so a +// differently-offset view stages its own shadow and never consumes another +// slice's bytes (the Qwen3.5 BA interior-view fatality). +std::map& KeepQuantWordShadows() { + static std::map* m = + new std::map(); // never destroyed (#1486) + return *m; +} + +// Stage the packed keep-quant stream as a resident i32 word tensor ONCE per +// weight — the eager pre-capture step or the first eager call pays it — and +// serve it forever after: the per-call decode runs entirely on-core from the +// resident words (no host repack, no from_vector upload). The packed master is +// immutable post-load, so a serve can never go stale (same assumption as the +// BufferSlot/WeightViewShadow resident shadows). +ttnn::Tensor EnsureKeepQuantWords(const Tensor& packed, DType enc, int64_t rows, + int64_t nb, MeshDevice& device) { + VT_CHECK(packed.rank == 2 && packed.IsContiguous(), + "tenstorrent EnsureKeepQuantWords: contiguous rank-2 packed"); + const int wpb = KeepQuantWordsPerBlock(enc); + VT_CHECK(wpb > 0, "tenstorrent EnsureKeepQuantWords: unsupported encoding"); + { + std::lock_guard g(KeepQuantWordMutex()); + auto it = KeepQuantWordShadows().find(packed.data); + if (it != KeepQuantWordShadows().end() && it->second.rows == rows && + it->second.nb == nb && it->second.wpb == wpb) { + return it->second.words; + } + } + // MISS = the one staging write this weight ever pays. A capture-time arrival + // refuses by name (the ServeActF32 / zero-cache precedent): a captured graph + // pins capture-time bytes its replay cannot refresh (the #2812 class), and + // the engine's pre-capture eager step ("run one EAGER step", qwen3_5.cpp) + // warms the shadow, so every later call — captured or not — hits it. + if (tt_capture_active()) { + KeepQuantCaptureStagingWritesCounter()++; + if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr) + std::fprintf(stderr, + "[TT-KQ] EnsureKeepQuantWords word staging during capture\n"); + } + VT_CHECK(!tt_capture_active(), + "tenstorrent: keep-quant word-shadow miss during trace capture (" + + std::string(Name(enc)) + + ") — warm the keep-quant arm eagerly first"); EnsureHost(packed); const uint8_t* bytes = packed.Ptr(); - // 36 u32 words per 144-byte block, little-endian (GGUF stream order; the - // f16 halves land in word 0: d in bits 0..15, dmin in bits 16..31). + const int64_t b64 = rows * nb; + const int64_t block_bytes = BlockBytes(enc); std::vector words; - words.reserve(static_cast(b64) * 36); + words.reserve(static_cast(b64) * static_cast(wpb)); + std::vector padded(static_cast(wpb) * 4u, 0u); for (int64_t b = 0; b < b64; ++b) { - const uint8_t* blk = bytes + b * 144; - for (int wj = 0; wj < 36; ++wj) { - const uint8_t* p = blk + wj * 4; + std::memcpy(padded.data(), bytes + b * block_bytes, + static_cast(block_bytes)); + for (int wj = 0; wj < wpb; ++wj) { + const uint8_t* p = padded.data() + wj * 4; words.push_back(static_cast( static_cast(p[0]) | (static_cast(p[1]) << 8) | (static_cast(p[2]) << 16) | (static_cast(p[3]) << 24))); } } - ttnn::Tensor w = ttnn::Tensor::from_vector( + ttnn::Tensor staged = ttnn::Tensor::from_vector( std::move(words), - SpecOf(tt::tt_metal::Shape({B, 36u}), ttnn::DataType::INT32, - ttnn::Layout::ROW_MAJOR), + SpecOf(tt::tt_metal::Shape({static_cast(b64), + static_cast(wpb)}), + ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR), &device); + std::lock_guard g(KeepQuantWordMutex()); + KeepQuantWordShadows()[packed.data] = + KeepQuantWordShadow{staged, rows, nb, wpb}; + return staged; +} + +// Stream bytes [first, last) of the word-staged block as u8 {B, last-first} +// (little-endian lanes). concat stacks the four lane tensors, so the +// (lane, word)->(word, lane) permute restores stream order — the same trick +// the Q4_K scale extraction uses, generalized to any byte range. +ttnn::Tensor KeepQuantByteRange(const ttnn::Tensor& w, uint32_t B, int first, + int last) { + const int w0 = first / 4; + const int nwords = (last + 3) / 4 - w0; + ttnn::Tensor sw = ttnn::slice( + w, ttsl::SmallVector{0u, static_cast(w0)}, + ttsl::SmallVector{B, + static_cast(w0 + nwords)}, + ttsl::SmallVector{1u, 1u}); + auto byte_lane = [](const ttnn::Tensor& t, int shift) { + return ttnn::bitwise_and(ttnn::bitwise_right_shift(t, shift), 0xFF); + }; + ttnn::Tensor stream = ttnn::reshape( + ttnn::permute( + ttnn::reshape( + ttnn::concat(std::vector{ + byte_lane(sw, 0), byte_lane(sw, 8), byte_lane(sw, 16), + byte_lane(sw, 24)}, + /*dim=*/1), + ttnn::Shape({B, 4u, static_cast(nwords)})), + ttsl::SmallVector{0, 2, 1}), + ttnn::Shape({B, static_cast(4 * nwords)})); + const int off = first - 4 * w0; + return ttnn::slice( + stream, ttsl::SmallVector{0u, static_cast(off)}, + ttsl::SmallVector{B, + static_cast(off + last - first)}, + ttsl::SmallVector{1u, 1u}); +} + +// The -0.0f signed-zero repair constant, one cached device tensor per shape. +// It cannot be BUILT on device: the i32->f32 bitcast maps -0 to +0 (the W1 +// comment was literal — bit-31-setting scalars also die in the scalar +// binding, and multiply/subtract canonicalize what survives), so every +// device-side construction arrived +0 and a forced-true ternary wrote +0. +// Host bytes survive the whole chain: -0.0f uploaded once per shape — a +// one-time staging fill, warmed by the same eager pre-pass as the word +// shadows and refused on a capture-time miss — then selected by where(), +// pure data movement on both ends. +ttnn::Tensor Neg0CacheGet(const ttnn::Shape& shape, MeshDevice& device) { + const std::string key = "neg0/" + ZeroCacheKey(shape, + ttnn::DataType::FLOAT32, + ttnn::Layout::TILE); + std::lock_guard g(ZeroCacheMutex()); + auto& c = ZeroCache(); + auto it = c.find(key); + if (it == c.end()) { + VT_CHECK(!tt_capture_active(), + "tenstorrent: -0 cache miss during capture — warm the keep-quant " + "decode eagerly first"); + uint64_t n = 1; + for (const auto d : shape.view()) n *= d; + std::vector z(static_cast(n), -0.0f); + it = c.emplace(key, + ttnn::to_layout( + ttnn::Tensor::from_vector( + std::move(z), + SpecOf(shape, ttnn::DataType::FLOAT32, + ttnn::Layout::ROW_MAJOR), + &device), + ttnn::Layout::TILE)) + .first; + } + return it->second; +} + +ttnn::Tensor DecodeKeepQuantBlocksF32(const Tensor& packed, DType enc, + int64_t rows, int64_t nb, + MeshDevice& device) { + const uint32_t B = static_cast(rows * nb); + const ttnn::Tensor w = EnsureKeepQuantWords(packed, enc, rows, nb, device); // f16 bit pattern (held in INT32) -> f32 value, the integer chain of - // vt::F16ToF32 (minimax_h3_vae_loader.cpp:47): normal/inf/NaN bits are - // sign | (exp+112)<<23 | mant<<13; subnormal mant*2^-24 is exact in f32 - // (integer typecast times a power of two); zero keeps its sign bit. where() - // demands TILE, so the selection runs there; finite patterns come out - // bit-exact, and a NaN payload canonicalizes to inf across the SFPU - // bitcast (hardware pin — see the kernel comment above). auto f16_bits_to_f32 = [](ttnn::Tensor t) { t = ttnn::to_layout(t, ttnn::Layout::TILE); const ttnn::Tensor sign_b = ttnn::bitwise_left_shift( @@ -1806,213 +1975,464 @@ ttnn::Tensor DecodeQ4KBlocksF32(const Tensor& packed, int64_t rows, int64_t nb, return ttnn::to_layout(ttnn::bitcast(out_bits, ttnn::DataType::FLOAT32), ttnn::Layout::ROW_MAJOR); }; - ttnn::Tensor w0 = ttnn::slice( - w, ttsl::SmallVector{0u, 0u}, - ttsl::SmallVector{B, 1u}, ttsl::SmallVector{1u, 1u}); - ttnn::Tensor d_bits = ttnn::bitwise_and(w0, 0xFFFF); - ttnn::Tensor dmin_bits = - ttnn::bitwise_and(ttnn::bitwise_right_shift(w0, 16), 0xFFFF); - ttnn::Tensor d = f16_bits_to_f32(d_bits); - ttnn::Tensor dmin = f16_bits_to_f32(dmin_bits); - - // Scale bytes: words 1..3 hold scales[0..11]. Four byte-lane shifts, then a - // (t,word)->(word,t) lane order fix: concat stacks the shift tensors, so - // col = 3*lane + word; reshape to (lane,word), permute to (word,lane), and - // the flat order is scales[k], k = 4*word + lane. - ttnn::Tensor sw = ttnn::slice( - w, ttsl::SmallVector{0u, 1u}, - ttsl::SmallVector{B, 4u}, ttsl::SmallVector{1u, 1u}); - auto byte_lane = [](const ttnn::Tensor& t, int shift) { - return ttnn::bitwise_and(ttnn::bitwise_right_shift(t, shift), 0xFF); + // u8 byte -> SIGNED i8 value as exact f32: v - 256*bit7 (the int8 sign + // extension; every operand and step is exact in the integer domain). + // Q6_K's scales and Q8_0's qs are the signed-byte consumers. + auto signed_byte_f32 = [](const ttnn::Tensor& v) { + const ttnn::Tensor u = ttnn::bitwise_and(v, 0xFF); + const ttnn::Tensor bit7 = ttnn::bitwise_right_shift(u, 7); + return ttnn::subtract( + ttnn::typecast(u, ttnn::DataType::FLOAT32), + ttnn::multiply(ttnn::typecast(bit7, ttnn::DataType::FLOAT32), + 256.0f)); }; - ttnn::Tensor sb = ttnn::reshape( - ttnn::permute( - ttnn::reshape( - ttnn::concat(std::vector{ - byte_lane(sw, 0), byte_lane(sw, 8), byte_lane(sw, 16), - byte_lane(sw, 24)}, - /*dim=*/1), - ttnn::Shape({B, 4u, 3u})), - ttsl::SmallVector{0, 2, 1}), - ttnn::Shape({B, 12u})); - ttnn::Tensor sa = ttnn::slice( - sb, ttsl::SmallVector{0u, 0u}, - ttsl::SmallVector{B, 4u}, ttsl::SmallVector{1u, 1u}); - ttnn::Tensor sm = ttnn::slice( - sb, ttsl::SmallVector{0u, 4u}, - ttsl::SmallVector{B, 8u}, ttsl::SmallVector{1u, 1u}); - ttnn::Tensor sc = ttnn::slice( - sb, ttsl::SmallVector{0u, 8u}, - ttsl::SmallVector{B, 12u}, ttsl::SmallVector{1u, 1u}); - // GetScaleMinK4(is, scales) for is = 0..7, group-major: is<4 low pair from - // scales[is]/scales[is+4]; is>=4 high pair from scales[is+4] low bits and - // scales[is-4]/scales[is] top bits. - auto top6l4 = [](const ttnn::Tensor& t) { - return ttnn::bitwise_left_shift(ttnn::bitwise_right_shift(t, 6), 4); - }; - ttnn::Tensor sc_f = ttnn::typecast( - ttnn::concat(std::vector{ - ttnn::bitwise_and(sa, 63), - ttnn::bitwise_or(ttnn::bitwise_and(sc, 0xF), top6l4(sa))}, - /*dim=*/1), - ttnn::DataType::FLOAT32); // {B, 8}, values <= 63: exact - ttnn::Tensor mm_f = ttnn::typecast( - ttnn::concat(std::vector{ - ttnn::bitwise_and(sm, 63), - ttnn::bitwise_or(ttnn::bitwise_right_shift(sc, 4), top6l4(sm))}, - /*dim=*/1), - ttnn::DataType::FLOAT32); // {B, 8}, values <= 63: exact - - // Nibbles: words 4..35 are qs[128]; 8 nibble-lane shifts -> {B,256} with - // idx = 8*lane + word. Flat idx = 16l + 8h + 8q + r over (byte lane l, - // nibble half h, quarter q, word-in-quarter r); permute (l,h,q,r) -> - // (q,h,r,l) so the flat order is the output order q*64 + h*32 + 4r + l — - // scale group 2q+h covers the same 32-value run. - ttnn::Tensor qw = ttnn::slice( - w, ttsl::SmallVector{0u, 4u}, - ttsl::SmallVector{B, 36u}, ttsl::SmallVector{1u, 1u}); - auto nib = [](const ttnn::Tensor& t, int shift) { - return ttnn::bitwise_and(ttnn::bitwise_right_shift(t, shift), 0xF); - }; - ttnn::Tensor nibf = ttnn::typecast( - ttnn::reshape( - ttnn::permute( - ttnn::reshape( - ttnn::concat(std::vector{ - nib(qw, 0), nib(qw, 4), nib(qw, 8), nib(qw, 12), - nib(qw, 16), nib(qw, 20), nib(qw, 24), nib(qw, 28)}, - /*dim=*/1), - ttnn::Shape({B, 4u, 2u, 4u, 8u})), - ttsl::SmallVector{0, 3, 2, 4, 1}), - ttnn::Shape({B, 8u, 32u})), - ttnn::DataType::FLOAT32); // {B, 8, 32}, values <= 15: exact - - // y = (d*sc)*nib - (dmin*mm): the host's exact f32 order, as separate ops. - // Signed-zero repair first: the device multiply canonicalizes (-x)*0 to +0 - // and the i32->f32 bitcast maps -0 to +0, while the host chain is - // IEEE-exact for zeros — y keeps d's sign through a zero product and - // dmin's through a zero m1 (mm is unsigned, so the IEEE sign of both - // intermediates is the f16 sign bit). The sign is read from the f16 BIT - // patterns — never from a device float, whose zero sign is already - // canonicalized — and where() demands TILE, so the repairs run there. The - // combine subtracts back in ROW_MAJOR: a TILE {B,8,1} operand physically - // pads to 32 columns and a TILE broadcast reads those padding columns as - // data, which corrupted every non-zero logical column (238/256 at B=1). - auto neg01 = [](const ttnn::Tensor& bits) { - return ttnn::to_layout( - ttnn::bitwise_and(ttnn::bitwise_right_shift(bits, 15), 1), - ttnn::Layout::TILE); + // Signed-zero repair, third mechanism, after two device-falsified drafts. + // Draft one selected a -0 constant through where() over an arithmetic-f32 + // mask: the ternary never took its true branch (every sweep mismatch was + // dev +0 where the oracle keeps -0, none the reverse). Draft two OR-ed the + // sign bit into the product's own bits — bitwise work over a BITCAST of an + // F32 TILE tensor, which mangled every bit pattern through a + // reduced-precision path and landed the whole sweep on the f16 grid. This + // one keeps to primitives the working f16 decode above proves: + // - the zero test is a FLOAT compare (IEEE eq() is true for both zeros) + // — never a bitcast of the product; + // - the -0 constant is born-int: cached INT32 zeros OR the sign bit + // (scalar), then the chain's own bitcast(int -> float); + // - the where() predicate is a COMPARISON output (gt(mask, 0)) — the + // one predicate form the f16 decode exercises (its two where() calls); + // - the {0,1} masks are broadcast f32 arithmetic over sign bits + // typecast OUT of born-int patterns (the proven direction). + auto neg0_full = [&device](const ttnn::Shape& shape) { + return Neg0CacheGet(shape, device); }; - auto is_zero = [](ttnn::Tensor v) { - v = ttnn::to_layout(std::move(v), ttnn::Layout::TILE); - return ttnn::eq( - ttnn::bitwise_and( - ttnn::bitcast(std::move(v), ttnn::DataType::INT32), 0x7FFFFFFF), - 0); + // {0,1} f32 masks, built in ROW_MAJOR elementwise ops end to end: the + // sign bit travels as an integer (shift/and over the BIT pattern — never + // a device float, whose zero sign is already canonicalized), and the + // zero test is a float compare. Callers combine masks with broadcast f32 + // arithmetic and reshape ROW_MAJOR only. + auto sign_bit_f32 = [](const ttnn::Tensor& bits, int shift) { + return ttnn::typecast( + ttnn::bitwise_and(ttnn::bitwise_right_shift(bits, shift), 1), + ttnn::DataType::FLOAT32); }; - ttnn::Tensor prod = - ttnn::multiply(ttnn::reshape(ttnn::multiply(d, sc_f), - ttnn::Shape({B, 8u, 1u})), - nibf); - // pred = (sign bit set) AND (value == 0), assembled in the f32 domain: - // neg01/zero masks typecast to {0,1} f32 and a MULTIPLY broadcasts them - // ({B,1,1} x {B,8,32} -> {B,8,32}, exact on {0,1}) — the same broadcast- - // multiply family the d1/nib path uses. logical_and with a broadcast - // predicate is avoided deliberately, and so is any {B,256} reshape of - // prod: with either present the ternary wrote only the first 16 output - // elements and left recycled buffer bytes beyond. Every tensor here stays - // {B,8,32} — one 32x32 tile per block, the shape every op handles. - auto zero_mask = [&](ttnn::Tensor v) { - return ttnn::typecast(is_zero(std::move(v)), ttnn::DataType::FLOAT32); + auto zero_mask_f32 = [](const ttnn::Tensor& v) { + return ttnn::typecast( + ttnn::to_layout(ttnn::eq(v, 0.0f), ttnn::Layout::ROW_MAJOR), + ttnn::DataType::FLOAT32); }; - auto sign_mask = [&](const ttnn::Tensor& bits) { - return ttnn::typecast(neg01(bits), ttnn::DataType::FLOAT32); + // The repair itself: where(gt(mask, 0), -0, value), every input TILE (a + // TILE predicate mixed with ROW_MAJOR branches wrote only its first 16 + // output elements once already). Callers pass the value BY VALUE (a + // shared-handle copy): moving it in beside an argument expression that + // still reads it is unspecified-order evaluation, and the moved-from read + // segfaults inside to_layout. + auto repair = [](ttnn::Tensor value, const ttnn::Tensor& mask_f32, + const ttnn::Tensor& neg0) { + return ttnn::to_layout( + ttnn::where(ttnn::to_layout(ttnn::gt(mask_f32, 0.0f), + ttnn::Layout::TILE), + neg0, ttnn::to_layout(value, ttnn::Layout::TILE)), + ttnn::Layout::ROW_MAJOR); }; - ttnn::Tensor pred = ttnn::multiply( - ttnn::reshape(sign_mask(d_bits), ttnn::Shape({B, 1u, 1u})), - zero_mask(prod)); - // All three where() inputs in TILE: the ternary demands a TILE predicate, - // and a TILE predicate mixed with ROW_MAJOR branches made it write only - // its first 16 output elements (stale recycled bytes beyond). - ttnn::Tensor prodT = ttnn::to_layout(prod, ttnn::Layout::TILE); - ttnn::Tensor neg0T = ttnn::Tensor::from_vector( - std::vector(static_cast(B) * 256, -0.0f), - SpecOf(tt::tt_metal::Shape({B, 8u, 32u}), ttnn::DataType::FLOAT32, - ttnn::Layout::TILE), - &device); - prod = ttnn::to_layout( - ttnn::where(pred, neg0T, prodT), ttnn::Layout::ROW_MAJOR); - ttnn::Tensor m1 = ttnn::reshape(ttnn::multiply(dmin, mm_f), - ttnn::Shape({B, 8u, 1u})); - // Repair m1 at its FINAL {B,8,1} shape: the subtract must consume the - // where() output directly. Reshaping a to_layout(ROW_MAJOR) round-trip - // output before the subtract made the broadcast read block 0's m1 row for - // every block (constant y offset per element beyond block 0). - ttnn::Tensor pred8 = ttnn::multiply( - ttnn::reshape(sign_mask(dmin_bits), ttnn::Shape({B, 1u, 1u})), - zero_mask(m1)); - ttnn::Tensor m1T = ttnn::to_layout(m1, ttnn::Layout::TILE); - ttnn::Tensor neg0T8 = ttnn::Tensor::from_vector( - std::vector(static_cast(B) * 8, -0.0f), - SpecOf(tt::tt_metal::Shape({B, 8u, 1u}), ttnn::DataType::FLOAT32, - ttnn::Layout::TILE), - &device); - m1 = ttnn::to_layout( - ttnn::where(pred8, neg0T8, m1T), ttnn::Layout::ROW_MAJOR); - ttnn::Tensor y = ttnn::subtract(prod, m1); - return ttnn::reshape(ttnn::to_layout(std::move(y), ttnn::Layout::ROW_MAJOR), - ttnn::Shape({static_cast(rows), - static_cast(nb) * 256u})); + + switch (enc) { + case DType::kQ4_K: + case DType::kQ5_K: { + // Word 0: d | dmin. Words 1..3: scales[12]. Q5_K moves qs to words + // 12..43 (ql[128] @ byte 48) and inserts qh[32] @ byte 16 (words + // 4..11); the scale unpack and the 8x32 nibble planes are Q4_K's + // verbatim, and Q5_K adds the 5th bit (+16) before the same epilogue. + const bool q5 = enc == DType::kQ5_K; + ttnn::Tensor w0 = ttnn::slice( + w, ttsl::SmallVector{0u, 0u}, + ttsl::SmallVector{B, 1u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor d_bits = ttnn::bitwise_and(w0, 0xFFFF); + ttnn::Tensor dmin_bits = + ttnn::bitwise_and(ttnn::bitwise_right_shift(w0, 16), 0xFFFF); + ttnn::Tensor d = f16_bits_to_f32(d_bits); + ttnn::Tensor dmin = f16_bits_to_f32(dmin_bits); + + ttnn::Tensor sb = KeepQuantByteRange(w, B, 4, 16); // scales[12] + // GetScaleMinK4(is, scales) for is = 0..7, group-major: is<4 low pair from + // scales[is]/scales[is+4]; is>=4 high pair from scales[is+4] low bits and + // scales[is-4]/scales[is] top bits. + auto top6l4 = [](const ttnn::Tensor& t) { + return ttnn::bitwise_left_shift(ttnn::bitwise_right_shift(t, 6), 4); + }; + ttnn::Tensor sa = ttnn::slice( + sb, ttsl::SmallVector{0u, 0u}, + ttsl::SmallVector{B, 4u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor sm = ttnn::slice( + sb, ttsl::SmallVector{0u, 4u}, + ttsl::SmallVector{B, 8u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor sc = ttnn::slice( + sb, ttsl::SmallVector{0u, 8u}, + ttsl::SmallVector{B, 12u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor sc_f = ttnn::typecast( + ttnn::concat(std::vector{ + ttnn::bitwise_and(sa, 63), + ttnn::bitwise_or(ttnn::bitwise_and(sc, 0xF), top6l4(sa))}, + /*dim=*/1), + ttnn::DataType::FLOAT32); // {B, 8}, values <= 63: exact + ttnn::Tensor mm_f = ttnn::typecast( + ttnn::concat(std::vector{ + ttnn::bitwise_and(sm, 63), + ttnn::bitwise_or(ttnn::bitwise_right_shift(sc, 4), top6l4(sm))}, + /*dim=*/1), + ttnn::DataType::FLOAT32); // {B, 8}, values <= 63: exact + + // Nibbles: 32 words are qs[128]; 8 nibble-lane shifts -> {B,256} with + // idx = 8*lane + word. Flat idx = 16l + 8h + 8q + r over (byte lane l, + // nibble half h, quarter q, word-in-quarter r); permute (l,h,q,r) -> + // (q,h,r,l) so the flat order is the output order q*64 + h*32 + 4r + l — + // scale group 2q+h covers the same 32-value run. + const uint32_t qs_w0 = q5 ? 12u : 4u; + ttnn::Tensor qw = ttnn::slice( + w, ttsl::SmallVector{0u, qs_w0}, + ttsl::SmallVector{B, qs_w0 + 32u}, + ttsl::SmallVector{1u, 1u}); + auto nib = [](const ttnn::Tensor& t, int shift) { + return ttnn::bitwise_and(ttnn::bitwise_right_shift(t, shift), 0xF); + }; + ttnn::Tensor x5 = ttnn::typecast( + ttnn::reshape( + ttnn::permute( + ttnn::reshape( + ttnn::concat(std::vector{ + nib(qw, 0), nib(qw, 4), nib(qw, 8), nib(qw, 12), + nib(qw, 16), nib(qw, 20), nib(qw, 24), nib(qw, 28)}, + /*dim=*/1), + ttnn::Shape({B, 4u, 2u, 4u, 8u})), + ttsl::SmallVector{0, 3, 2, 4, 1}), + ttnn::Shape({B, 8u, 32u})), + ttnn::DataType::FLOAT32); // {B, 8, 32}, values <= 15: exact + if (q5) { + // The 5th bit: output col c reads bit c/32 — the group index g — of + // qh[c%32] (host u1 = 1<<2q for the low plane and u2 = 2<<2q for the + // high plane of quarter q; g = 2q+h is exactly that bit). qh bytes + // broadcast {B,1,32} against one shift per group; +16*bit is exact. + ttnn::Tensor qh = KeepQuantByteRange(w, B, 16, 48); // qh[32] + ttnn::Tensor qh3 = ttnn::reshape(qh, ttnn::Shape({B, 1u, 32u})); + std::vector planes; + planes.reserve(8); + for (int g = 0; g < 8; ++g) + planes.push_back( + ttnn::bitwise_and(ttnn::bitwise_right_shift(qh3, g), 1)); + x5 = ttnn::add( + x5, + ttnn::multiply( + ttnn::typecast( + ttnn::reshape(ttnn::concat(std::move(planes), /*dim=*/1), + ttnn::Shape({B, 8u, 32u})), + ttnn::DataType::FLOAT32), + 16.0f)); // values <= 31: exact + } + + // y = (d*sc)*x - (dmin*mm): the host's exact f32 order, as separate ops, + // every operand held at {B,8,32} — one 32x32 tile per block, the shape + // every op handles (a TILE {B,8,1} operand physically pads to 32 + // columns and a TILE broadcast reads those padding columns as data, + // which corrupted every non-zero logical column at B=1). + // Signed-zero repair as bit work (or_sign above): the device multiply + // AND subtract canonicalize a zero result's sign, while the host chain + // is IEEE-exact — prod keeps d's sign through a zero product and m1 + // keeps dmin's through a zero product (mm and x are unsigned, so the + // IEEE sign of both intermediates is the f16 sign bit), and a zero y + // keeps a sign only when the m1 term's bit is clear ((-0) - (+0) is + // -0). All three masks are broadcast f32 arithmetic on {0,1} values + // read from the f16 BIT patterns — never from a device float, whose + // zero sign is already canonicalized. + // Signed-zero repair (repair/neg0_full above): the device multiply AND + // subtract canonicalize a zero result's sign, while the host chain is + // IEEE-exact — prod keeps d's sign through a zero product and m1 keeps + // dmin's through a zero product (mm and x are unsigned, so the IEEE + // sign of both intermediates is the f16 sign bit), and a zero y keeps a + // sign only when the m1 term's bit is clear ((-0) - (+0) is -0). + ttnn::Tensor dsign = + ttnn::reshape(sign_bit_f32(d_bits, 15), ttnn::Shape({B, 1u, 1u})); + ttnn::Tensor prod = + ttnn::multiply(ttnn::reshape(ttnn::multiply(d, sc_f), + ttnn::Shape({B, 8u, 1u})), + x5); + prod = repair(prod, ttnn::multiply(dsign, zero_mask_f32(prod)), + neg0_full(ttnn::Shape({B, 8u, 32u}))); + ttnn::Tensor dminsign = + ttnn::reshape(sign_bit_f32(dmin_bits, 15), ttnn::Shape({B, 1u, 1u})); + // Repair m1 at its FINAL {B,8,1} shape: the subtract must consume the + // repaired tensor directly. Reshaping a to_layout(ROW_MAJOR) round-trip + // output before the subtract made the broadcast read block 0's m1 row + // for every block (constant y offset per element beyond block 0). + ttnn::Tensor m1 = ttnn::reshape(ttnn::multiply(dmin, mm_f), + ttnn::Shape({B, 8u, 1u})); + m1 = repair(m1, ttnn::multiply(dminsign, zero_mask_f32(m1)), + neg0_full(ttnn::Shape({B, 8u, 1u}))); + ttnn::Tensor y = ttnn::subtract(prod, m1); + y = repair( + y, + ttnn::multiply( + ttnn::reshape(ttnn::subtract(dsign, + ttnn::multiply(dsign, dminsign)), + ttnn::Shape({B, 1u, 1u})), + zero_mask_f32(y)), + neg0_full(ttnn::Shape({B, 8u, 32u}))); + return ttnn::reshape( + std::move(y), + ttnn::Shape({static_cast(rows), + static_cast(nb) * 256u})); + } + case DType::kQ6_K: { + // Word 52 holds d in its low half (bytes 208..209). 16 sub-blocks of + // 16: output col c sits in sub-block s = c/16 at i = c%16, and + // s = 8h + 2r + l/16 over (half h, run r, l in 0..32): per (h, r) the + // host loop reads ql byte 64h + 32*(r&1) + l, low nibble for r<2 / high + // otherwise, ORs the 2 high bits (qh[32h + l] >> 2r) & 3 into bit 4, + // subtracts the 32 bias, and multiplies (d*sc)*q with the SIGNED scale + // sc[8h + 2r + l/16] — no min term. Eight {B,2,16} pieces concat in + // (h, r, l/16, l%16) order straight into [s][i]. + ttnn::Tensor w52 = ttnn::slice( + w, ttsl::SmallVector{0u, 52u}, + ttsl::SmallVector{B, 53u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor d_bits = ttnn::bitwise_and(w52, 0xFFFF); + ttnn::Tensor d3 = ttnn::reshape(f16_bits_to_f32(d_bits), + ttnn::Shape({B, 1u, 1u})); + ttnn::Tensor ql = KeepQuantByteRange(w, B, 0, 128); + ttnn::Tensor qh = KeepQuantByteRange(w, B, 128, 192); + ttnn::Tensor sc_bytes = KeepQuantByteRange(w, B, 192, 208); + ttnn::Tensor sc_f = signed_byte_f32(sc_bytes); // {B,16}: exact ints + ttnn::Tensor sc_sign = ttnn::typecast( + ttnn::bitwise_and(ttnn::bitwise_right_shift(sc_bytes, 7), 1), + ttnn::DataType::FLOAT32); // {B,16} in {0,1}: the i8 sign bit + std::vector halves; + halves.reserve(2); + std::vector sign_halves; + sign_halves.reserve(2); + for (int h = 0; h < 2; ++h) { + std::vector runs; + runs.reserve(4); + std::vector sign_runs; + sign_runs.reserve(4); + for (int r = 0; r < 4; ++r) { + const int qoff = 64 * h + 32 * (r % 2); + ttnn::Tensor qb = ttnn::slice( + ql, + ttsl::SmallVector{0u, + static_cast(qoff)}, + ttsl::SmallVector{ + B, static_cast(qoff + 32)}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor nib = (r < 2) ? ttnn::bitwise_and(qb, 0xF) + : ttnn::bitwise_right_shift(qb, 4); + ttnn::Tensor hb = ttnn::slice( + qh, + ttsl::SmallVector{0u, + static_cast(32 * h)}, + ttsl::SmallVector{ + B, static_cast(32 * h + 32)}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor hi2 = ttnn::bitwise_and( + ttnn::bitwise_right_shift(hb, 2 * r), 3); + ttnn::Tensor nib6 = + ttnn::bitwise_or(nib, ttnn::bitwise_left_shift(hi2, 4)); + ttnn::Tensor q6 = ttnn::subtract( + ttnn::typecast(nib6, ttnn::DataType::FLOAT32), + 32.0f); // {B,32}, values in [-32, 31]: exact + const int soff = 8 * h + 2 * r; + ttnn::Tensor s2 = ttnn::slice( + sc_f, + ttsl::SmallVector{0u, + static_cast(soff)}, + ttsl::SmallVector{ + B, static_cast(soff + 2)}, + ttsl::SmallVector{1u, 1u}); + // (d*sc) first — the host's left-to-right association — then *q. + runs.push_back(ttnn::multiply( + ttnn::multiply(d3, + ttnn::reshape(s2, ttnn::Shape({B, 2u, 1u}))), + ttnn::reshape(q6, ttnn::Shape({B, 2u, 16u})))); + // sign(q): the nibble sits below bit 5 before the -32 bias, so + // q < 0 is exactly nib6 < 32 — an INT32 compare on the raw bits, + // never a device float's canonicalized zero sign. + sign_runs.push_back( + ttnn::reshape(ttnn::typecast(ttnn::lt(nib6, 32), + ttnn::DataType::FLOAT32), + ttnn::Shape({B, 2u, 16u}))); + } + halves.push_back(ttnn::concat(std::move(runs), /*dim=*/1)); + sign_halves.push_back( + ttnn::concat(std::move(sign_runs), /*dim=*/1)); + } + ttnn::Tensor prod = ttnn::reshape( + ttnn::concat(std::move(halves), /*dim=*/1), + ttnn::Shape({B, 16u, 16u})); // [s][i] with s = 8h + 2r + l/16 + ttnn::Tensor qsign = ttnn::reshape( + ttnn::concat(std::move(sign_halves), /*dim=*/1), + ttnn::Shape({B, 16u, 16u})); + // Zero-product sign: the IEEE sign of (d*sc)*q is the XOR of all + // three operand signs — d's f16 bit, sc's i8 bit7, and sign(q) above + // (q and the typecast zeros are +0, so they contribute their bit + // directly). The XOR of three {0,1} masks is + // a+b+c-2(ab+ac+bc)+4abc, exact in f32 broadcast arithmetic; the + // product is repaired where that XOR lands on a zero (or_sign). + ttnn::Tensor ds = ttnn::reshape(sign_bit_f32(d_bits, 15), + ttnn::Shape({B, 1u, 1u})); + ttnn::Tensor ss = ttnn::reshape(sc_sign, ttnn::Shape({B, 16u, 1u})); + ttnn::Tensor ds_ss = ttnn::multiply(ds, ss); // {B,16,1} + ttnn::Tensor pairs = ttnn::add( + ttnn::add(ds_ss, ttnn::multiply(ds, qsign)), + ttnn::multiply(ss, qsign)); // ab + ac + bc + ttnn::Tensor quad = ttnn::multiply(ds_ss, qsign); // abc + ttnn::Tensor pred = + ttnn::add(ttnn::subtract(ttnn::add(ttnn::add(ds, ss), qsign), + ttnn::multiply(pairs, 2.0f)), + ttnn::multiply(quad, 4.0f)); + prod = repair(prod, ttnn::multiply(pred, zero_mask_f32(prod)), + neg0_full(ttnn::Shape({B, 16u, 16u}))); + return ttnn::reshape( + prod, + ttnn::Shape({static_cast(rows), + static_cast(nb) * 256u})); + } + case DType::kQ8_0: { + // Word 0: d | first two qs bytes. qs[32] = stream bytes 2..33; host: + // y = qs * d — IEEE multiply is commutative in value AND zero sign, so + // the device's d*q is the same bit pattern. + ttnn::Tensor w0 = ttnn::slice( + w, ttsl::SmallVector{0u, 0u}, + ttsl::SmallVector{B, 1u}, + ttsl::SmallVector{1u, 1u}); + ttnn::Tensor d_bits = ttnn::bitwise_and(w0, 0xFFFF); + ttnn::Tensor d = f16_bits_to_f32(d_bits); // {B,1} + ttnn::Tensor qb = KeepQuantByteRange(w, B, 2, 34); // {B,32} raw bytes + ttnn::Tensor qf = signed_byte_f32(qb); + ttnn::Tensor prod = ttnn::multiply(d, qf); // {B,32} + // Zero-product sign: the IEEE product's sign is sign(d) XOR sign(q) — + // sign(d) is the f16 bit, sign(q) the raw byte's bit7 (a zero byte + // typecasts to +0 and contributes a clear bit, exactly its sign). The + // XOR of two {0,1} masks is a+b-2ab, exact in f32 broadcast + // arithmetic; the product is repaired where the XOR lands on a zero. + ttnn::Tensor dsign = sign_bit_f32(d_bits, 15); // {B,1} + ttnn::Tensor qsign = sign_bit_f32(qb, 7); // {B,32} + ttnn::Tensor pred = ttnn::subtract( + ttnn::add(dsign, qsign), + ttnn::multiply(ttnn::multiply(dsign, qsign), 2.0f)); + prod = repair(prod, ttnn::multiply(pred, zero_mask_f32(prod)), + neg0_full(ttnn::Shape({B, 32u}))); + return ttnn::reshape( + prod, + ttnn::Shape({static_cast(rows), + static_cast(nb) * 32u})); + } + default: + // EnsureKeepQuantWords refuses anything else before this point. + VT_CHECK(false, "tenstorrent keep-quant decode: unsupported encoding"); + return w; + } } 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"); + "rank-2 [rows, nb*elems]"); + const DType enc = packed.dtype; + VT_CHECK(enc == DType::kQ4_K || enc == DType::kQ5_K || enc == DType::kQ6_K || + enc == DType::kQ8_0, + std::string("tenstorrent kKeepQuantDecode: packed dtype ") + + Name(enc) + + " has no registered keep-quant decode (registered set: kQ4_K/" + "kQ5_K/kQ6_K/kQ8_0)"); + const int64_t elems = BlockElems(enc); 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, + VT_CHECK(out.shape[0] == rows && out.shape[1] == nb * elems, "tenstorrent kKeepQuantDecode: out shape mismatch"); - ttnn::Tensor y = DecodeQ4KBlocksF32(packed, rows, nb, SharedMeshDevice()); + ttnn::Tensor y = + DecodeKeepQuantBlocksF32(packed, enc, 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). + static_cast(nb * elems)); +} + +// kMatmulBTQuant (KEEPQUANT W2, generalized W3): `a` is [M,K] float, `b` is +// [N,K] packed keep-quant 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 bit-exact f32 chain +// (DecodeKeepQuantBlocksF32 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 tests pin it, not a picked +// tolerance. W3 residency: the decode runs on-core from the per-weight +// resident i32 word shadow (EnsureKeepQuantWords) — the eager pre-capture +// step stages every weight once and the captured replay hits it. +// The decoded bf16 twin of a keep-quant weight, memoized per host weight +// pointer. MatmulBTQuantKernel's first draft decoded EVERY weight on EVERY +// call; the Qwen3.5-0.8B Q4_K_M vehicle's tied head is [248320, 1024] Q6_K, +// so every step re-ran a quarter-billion-element unpack whose f32 planes + +// where intermediates (~1 GB each) churned and fragmented the DRAM banks +// until a large contiguous allocation failed (TT_FATAL Out of Memory during +// the first generate). Weights are immutable after load — the same invariant +// WeightViewShadow documents — so the twin is built once and reused. The +// recycled-address hazard is handled the same way too: a collision must also +// match rows x cols AND the encoding to hit. Capture-safe by construction: +// the eager warm step builds every twin before capture, so the captured graph +// reads one stable tensor per weight and stages nothing per replay. +struct DecodedWeightShadow { + std::optional device; + uint32_t rows = 0, cols = 0; + DType enc = DType::kF32; +}; +std::mutex& DecodedWeightMutex() { + static std::mutex m; + return m; +} +std::map& DecodedWeightShadows() { + static std::map* m = + new std::map(); // never destroyed (#1486) + return *m; +} +// Free path hook: a freed host weight must drop its twin, so a recycled +// address can never alias a stale decode (UnregisterHostBuffer). +void DropDecodedWeightShadow(void* host) { + if (host == nullptr) return; + std::lock_guard g(DecodedWeightMutex()); + DecodedWeightShadows().erase(reinterpret_cast(host)); +} + 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)"); + // (gguf_keep_quant.cpp). Refusing here BY NAME keeps an admitted-but- + // unimplemented encoding from reaching the device. + const DType enc = b.dtype; + // vt::Name() emits the lowercase storage name ("q4_0"); the refusal must + // name the ENUM the caller passed, so the k-prefix and capital go on here. + const std::string enc_lower = Name(enc); + const std::string enc_name = + std::string("k") + static_cast(enc_lower[0] - 'a' + 'A') + + enc_lower.substr(1); + VT_CHECK(enc == DType::kQ4_K || enc == DType::kQ5_K || enc == DType::kQ6_K || + enc == DType::kQ8_0, + std::string("tenstorrent kMatmulBTQuant: ") + enc_name + + " has no keep-quant decode on TENSTORRENT; the registered set " + "is kQ4_K/kQ5_K/kQ6_K/kQ8_0 (BACKEND-TENSTORRENT-KEEPQUANT)"); + const int64_t elems = BlockElems(enc); + VT_CHECK(b.shape[1] % elems == 0, + std::string("tenstorrent kMatmulBTQuant: K must be a whole number " + "of ") + + Name(enc) + " blocks (" + std::to_string(elems) + " elems)"); VT_CHECK(IsFloatDType(a.dtype) && (out.dtype == DType::kF32 || out.dtype == DType::kBF16), "tenstorrent kMatmulBTQuant: float activation, f32/bf16 out"); @@ -2026,10 +2446,48 @@ void MatmulBTQuantKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& b) "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); + // Memoized twin: decode once, reuse across steps (see DecodedWeightShadow). + bool have_twin = false; + ttnn::Tensor w_bf16; + { + std::lock_guard g(DecodedWeightMutex()); + auto it = DecodedWeightShadows().find(reinterpret_cast(b.data)); + if (it != DecodedWeightShadows().end() && it->second.rows == N && + it->second.cols == K && it->second.enc == enc && it->second.device) { + w_bf16 = *it->second.device; + have_twin = true; + } + } + if (!have_twin) { + // Build the twin HOST-side. The first draft built it on device + // (DecodeKeepQuantBlocksF32 + typecast); for the vehicle's [248320,1024] + // Q6_K tied head that materialized ~1 GB f32 planes plus where + // intermediates whose transient churn fragmented the DRAM banks until a + // 4 GB allocation failed (TT_FATAL Out of Memory) — the FIRST decode of + // the head did not fit alongside residency even once, so the twin cache + // alone could not save the run. The host decoder is the SAME chain the + // kKeepQuantDecode test pins bit-exact (vt::cpu::BlockToFloat == + // DecodeKeepQuantBlocksF32), so the f32 bits are unchanged; only where + // they are produced moves. The bf16 round then happens once, host-side, + // with the same RNE convention (vt::F32ToBF16) the device typecast + // applied — pre-rounding to bf16-representable f32 makes the upload + // independent of from_vector's own conversion mode — and from_vector + // lands it as a TILE bf16 tensor, the identical upload the embedding + // table and UploadRows already use. Decoding host-side also frees the + // device from the per-call word-shadow reads for this one-time build. + std::vector w_f32(static_cast(N) * K); + vt::cpu::BlockToFloat(enc)(b.data, w_f32.data(), static_cast(N) * K); + for (size_t i = 0; i < w_f32.size(); ++i) + w_f32[i] = BF16ToF32(F32ToBF16(w_f32[i])); + w_bf16 = ttnn::Tensor::from_vector(w_f32, TileSpecOf(N, K), &device); + std::lock_guard g(DecodedWeightMutex()); + DecodedWeightShadow& s = DecodedWeightShadows()[reinterpret_cast( + b.data)]; + s.device = w_bf16; + s.rows = N; + s.cols = K; + s.enc = enc; + } ttnn::Tensor dev_a = EnsureDevice2D(a, device); if (a.dtype == DType::kF32) dev_a = ttnn::to_layout( @@ -2338,6 +2796,35 @@ void SiluAndMulKernel(Queue&, Tensor& out, const Tensor& x) { CommitDevice2D(out, std::move(dev_y)); } +// kMoeSiluMul: silu(gate) * up with SPLIT operands (ops.cpp MoeSiluMul -> id 63; +// cpu_ops.cpp MoeSiluMulKernel, cuda_moe.cu MoeSiluMulKernel). The split sibling +// of kSiluAndMul above: the GGUF dense MLP arm reaches the DenseMlpBlock tail +// with gate/up as TWO separate [T,I] GEMM outputs (GGUF stores ff_gate/ff_up +// unmerged), so there is no merged [T,2I] operand to slice. Same ttnn math +// (ttnn::silu + ttnn::multiply) and the same capture-safe EnsureDevice2D / +// CommitDevice2D staging as kSiluAndMul, so the captured e2e treats it +// identically. The vehicle path feeds f32 gate/up, where the CPU kernel's +// RoundThrough(gate.dtype) narrowing is the identity. +void MoeSiluMulKernel(Queue&, Tensor& out, const Tensor& gate, const Tensor& up) { + TT_OP_TRACE("MoeSiluMul"); + VT_CHECK(gate.rank == 2 && up.rank == 2 && out.rank == 2, + "tenstorrent kMoeSiluMul: only rank-2 tensors are supported"); + VT_CHECK(IsFloatDType(gate.dtype) && IsFloatDType(up.dtype) && + (out.dtype == DType::kF32 || out.dtype == DType::kBF16), + "tenstorrent kMoeSiluMul: float in, f32/bf16 out"); + VT_CHECK(gate.IsContiguous() && up.IsContiguous() && out.IsContiguous(), + "tenstorrent kMoeSiluMul: strided (non-contiguous) tensors are not supported"); + VT_CHECK(gate.shape[0] == out.shape[0] && gate.shape[1] == out.shape[1] && + up.shape[0] == gate.shape[0] && up.shape[1] == gate.shape[1], + "tenstorrent kMoeSiluMul: gate, up and out must share one [T, I] shape"); + + MeshDevice& device = SharedMeshDevice(); + ttnn::Tensor dev_g = EnsureDevice2D(gate, device); + ttnn::Tensor dev_u = EnsureDevice2D(up, device); + ttnn::Tensor dev_y = ttnn::multiply(ttnn::silu(dev_g), dev_u); + CommitDevice2D(out, std::move(dev_y)); +} + // kCastBf16 / kCastF32: elementwise dtype convert via Load/Store (cpu_ops // CastBf16Kernel / CastF32Kernel). Qwen3 uses these for K/V cache dtype and // the logits / rope-cache paths. Host-staged; bit-exact for supported pairs. @@ -6732,6 +7219,8 @@ struct Registrar { reinterpret_cast(static_cast(&RmsNormKernel))); RegisterOp(OpId::kSiluAndMul, DeviceType::kTENSTORRENT, reinterpret_cast(static_cast(&SiluAndMulKernel))); + RegisterOp(OpId::kMoeSiluMul, DeviceType::kTENSTORRENT, + reinterpret_cast(static_cast(&MoeSiluMulKernel))); RegisterOp(OpId::kCastBf16, DeviceType::kTENSTORRENT, reinterpret_cast(static_cast(&CastBf16Kernel))); RegisterOp(OpId::kCastF32, DeviceType::kTENSTORRENT, @@ -6781,6 +7270,13 @@ struct Registrar { } // namespace +int64_t KeepQuantCaptureStagingWrites() { + return KeepQuantCaptureStagingWritesCounter().load(std::memory_order_relaxed); +} +void ResetKeepQuantCaptureStagingWritesForTest() { + KeepQuantCaptureStagingWritesCounter().store(0, std::memory_order_relaxed); +} + // ---- ttnn mesh-trace capture (Backend graph-capture mapping) ---------------- // Process-local single-slot capture + multi-graph handles (opaque MeshTraceId*). // Mirrors the CUDA backend's single-exec_ vs EndCaptureGraph split. @@ -7057,6 +7553,7 @@ void UnregisterHostBuffer(void* host) { } DropPagedKvShadow(host); DropEmbedTableShadow(host); + DropDecodedWeightShadow(host); } void MarkHostWritten(void* host) { diff --git a/tests/vt/test_tenstorrent_backend.cpp b/tests/vt/test_tenstorrent_backend.cpp index 87b6f5996..ce7999f3c 100644 --- a/tests/vt/test_tenstorrent_backend.cpp +++ b/tests/vt/test_tenstorrent_backend.cpp @@ -654,6 +654,61 @@ TEST_CASE("kTENSTORRENT kSiluAndMul matches host F32 within BF16 envelope") { CHECK(max_abs_diff < 0.05f); } +// Split-operand sibling of the case above: the GGUF dense MLP arm reaches +// vt::MoeSiluMul (qwen3_5.cpp DenseMlpBlock tail) with gate/up as TWO separate +// [T,I] GEMM outputs, so this op must be registered for the Qwen3.5 GGUF +// vehicle to run at all — the captured e2e previously fatalled on it +// ("no kernel for op MoeSiluMul ... on device tenstorrent"). +TEST_CASE("kTENSTORRENT kMoeSiluMul matches host F32 within BF16 envelope") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kMoeSiluMul, DeviceType::kTENSTORRENT)); + + constexpr int64_t T = 7, D = 16; + Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT); + auto moe_silu_mul = reinterpret_cast( + vt::GetOp(vt::OpId::kMoeSiluMul, DeviceType::kTENSTORRENT)); + + std::vector host_g(T * D), host_u(T * D), host_out(T * D, 0.0f); + for (size_t i = 0; i < host_g.size(); ++i) + host_g[i] = (static_cast(i % 13) - 6.0f) * 0.2f; + for (size_t i = 0; i < host_u.size(); ++i) + host_u[i] = (static_cast(i % 7) - 3.0f) * 0.3f; + + void* mem_g = backend.Alloc(host_g.size() * sizeof(float)); + void* mem_u = backend.Alloc(host_u.size() * sizeof(float)); + void* mem_out = backend.Alloc(host_out.size() * sizeof(float)); + Queue q = backend.CreateQueue(); + backend.Copy(q, mem_g, host_g.data(), host_g.size() * sizeof(float)); + backend.Copy(q, mem_u, host_u.data(), host_u.size() * sizeof(float)); + + Tensor g = Tensor::Contiguous(mem_g, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {T, D}); + Tensor u = Tensor::Contiguous(mem_u, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {T, D}); + Tensor out = Tensor::Contiguous(mem_out, vt::DType::kF32, + Device{DeviceType::kTENSTORRENT, 0}, {T, D}); + moe_silu_mul(q, out, g, u); + backend.Copy(q, host_out.data(), mem_out, host_out.size() * sizeof(float)); + backend.Free(mem_g); + backend.Free(mem_u); + backend.Free(mem_out); + + float max_abs_diff = 0.0f; + for (int64_t i = 0; i < T; ++i) { + for (int64_t j = 0; j < D; ++j) { + const float gate = host_g[static_cast(i * D + j)]; + const float up = host_u[static_cast(i * D + j)]; + const float ref = (gate / (1.0f + std::exp(-gate))) * up; + max_abs_diff = std::max(max_abs_diff, std::fabs(host_out[static_cast(i * D + j)] - ref)); + } + } + // BF16 tile storage + silu; same envelope as the kSiluAndMul sibling. + CHECK(max_abs_diff < 0.05f); +} + // Cast pair used by Qwen3 K/V cache dtype and logits paths. TEST_CASE("kTENSTORRENT kCastBf16 / kCastF32 round-trip F32 values") { if (!TenstorrentPresent()) { @@ -5318,3 +5373,384 @@ TEST_CASE("kTENSTORRENT kMatmulBTQuant Q4_K via vt::MatmulBT matches the decode- } } } + +// KEEPQUANT W3 (issue #2959): the decode set generalizes to the other GGUF +// k-quants the q4km vehicle actually stores — Q5_K (attn_qkv/ssm_out), Q6_K +// (token_embd, tied LM head, half of ffn_down/attn_v) and Q8_0 (ssm_alpha/ +// ssm_beta). The numerics bar is W1's, unchanged: the f32 decode is BIT-EXACT +// against the CPU decoder `vt::cpu::BlockToFloat` (cpu_quant_dequant.cpp — +// dequantize_row_q5_K:1673, q6_K:1881, q8_0:495 at the pinned llama.cpp +// b10451), because these are the blocks the dot consumes and a band here is +// unattributable downstream. Same sweep shape as W1: rows {1,3,17} x blocks +// {1,2,16}, deterministic weights, both signs, plus the zero corners the +// device decode must repair (signed-zero products from a zero scale byte or +// a +/-0 d; q6's -32 bias crossing zero; q8_0's full int8 range). +TEST_CASE("kTENSTORRENT kKeepQuantDecode matches vt::cpu::BlockToFloat bit-exactly (Q5_K/Q6_K/Q8_0 sweep)") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kKeepQuantDecode, vt::DeviceType::kTENSTORRENT)); + + const vt::DType encodings[] = {vt::DType::kQ5_K, vt::DType::kQ6_K, + vt::DType::kQ8_0}; + for (const vt::DType enc : encodings) { + const int64_t kBlockBytes = vt::BlockBytes(enc); + const int64_t kBlockElems = vt::BlockElems(enc); + REQUIRE(kBlockElems * kBlockBytes > 0); + if (enc == vt::DType::kQ5_K) { + REQUIRE(kBlockBytes == 176); + REQUIRE(kBlockElems == 256); + } else if (enc == vt::DType::kQ6_K) { + REQUIRE(kBlockBytes == 210); + REQUIRE(kBlockElems == 256); + } else { + REQUIRE(kBlockBytes == 34); + REQUIRE(kBlockElems == 32); + } + + Backend& backend = vt::GetBackend(vt::DeviceType::kTENSTORRENT); + auto decode = reinterpret_cast( + vt::GetOp(vt::OpId::kKeepQuantDecode, vt::DeviceType::kTENSTORRENT)); + Queue q = backend.CreateQueue(); + + // Deterministic packed blocks with structural variety (the W1 generator, + // generalized): PRNG bytes for quants and scales, d (and dmin where the + // encoding has one) drawn from finite f16 magnitudes with both signs — + // random BYTES would make the f16 scales NaN/Inf and the bit comparison + // vacuous against a NaN-propagating decode. Every 7th block pins a + // signed-zero corner (d = +/-0, dmin = the other zero), every 11th an + // all-zero scale byte set (d*sc = +/-0, the signed-zero repair path). + std::mt19937 rng(20260907u); + auto rand_byte = [&rng]() { return static_cast(rng() & 0xFF); }; + + const int64_t rows_list[] = {1, 3, 17}; + const int64_t nb_list[] = {1, 2, 16}; + for (int64_t rows : rows_list) { + for (int64_t nb : nb_list) { + const int64_t k = nb * kBlockElems; + std::vector packed(rows * nb * kBlockBytes); + for (int64_t b = 0; b < rows * nb; ++b) { + uint8_t* blk = packed.data() + b * kBlockBytes; + const bool zero_d = (b % 7) == 3; + const bool zero_scales = (b % 11) == 5; + auto put_f16 = [&](int64_t off, float v) { + const uint16_t bits = vt::F32ToF16(v); + std::memcpy(blk + off, &bits, sizeof(bits)); + }; + auto put_scales = [&](int64_t off, int64_t n) { + for (int64_t i = 0; i < n; ++i) + blk[off + i] = zero_scales ? 0u : rand_byte(); + }; + if (enc == vt::DType::kQ5_K) { + // block_q5_K = { f16 d; f16 dmin; u8 scales[12]; u8 qh[32]; + // u8 qs[128]; } (176 bytes) + const float d = zero_d ? ((b % 2) ? -0.0f : 0.0f) + : (0.05f + 0.35f * static_cast(rng() % 64) / 64.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + const float dmin = zero_d ? ((b % 2) ? 0.0f : -0.0f) + : (0.005f + 0.02f * static_cast(rng() % 32) / 32.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + put_f16(0, d); + put_f16(2, dmin); + put_scales(4, 12); + for (int i = 0; i < 32; ++i) blk[16 + i] = rand_byte(); // qh + for (int i = 0; i < 128; ++i) blk[48 + i] = rand_byte(); // qs + } else if (enc == vt::DType::kQ6_K) { + // block_q6_K = { u8 ql[128]; u8 qh[64]; i8 scales[16]; f16 d; } + // (210 bytes) + for (int i = 0; i < 128; ++i) blk[0 + i] = rand_byte(); // ql + for (int i = 0; i < 64; ++i) blk[128 + i] = rand_byte(); // qh + put_scales(192, 16); + const float d = zero_d ? ((b % 2) ? -0.0f : 0.0f) + : (0.05f + 0.35f * static_cast(rng() % 64) / 64.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + put_f16(208, d); + } else { + // block_q8_0 = { f16 d; i8 qs[32]; } (34 bytes) + const float d = zero_d ? ((b % 2) ? -0.0f : 0.0f) + : (0.05f + 0.35f * static_cast(rng() % 64) / 64.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + put_f16(0, d); + for (int i = 0; i < 32; ++i) + blk[2 + i] = zero_scales ? 0u : rand_byte(); // full int8 range + } + } + + std::vector oracle(rows * k); + vt::cpu::BlockToFloat(enc)(packed.data(), oracle.data(), rows * k); + + void* mem_packed = backend.Alloc(packed.size()); + void* mem_out = backend.Alloc(oracle.size() * sizeof(float)); + backend.Copy(q, mem_packed, packed.data(), packed.size()); + + Tensor packed_t = + Tensor::Contiguous(mem_packed, enc, + Device{vt::DeviceType::kTENSTORRENT, 0}, {rows, nb}); + Tensor out_t = Tensor::Contiguous(mem_out, vt::DType::kF32, + Device{vt::DeviceType::kTENSTORRENT, 0}, {rows, k}); + decode(q, out_t, packed_t); + + std::vector device_out(rows * k, 0.0f); + backend.Copy(q, device_out.data(), mem_out, oracle.size() * sizeof(float)); + backend.Free(mem_packed); + backend.Free(mem_out); + + INFO("enc=", static_cast(enc), " rows=", rows, " nb=", nb, " K=", k); + if (std::memcmp(device_out.data(), oracle.data(), + oracle.size() * sizeof(float)) != 0) { + const float* dev = device_out.data(); + int64_t bad = 0; + for (int64_t i = 0; i < static_cast(oracle.size()); ++i) { + if (std::memcmp(&dev[i], &oracle[i], sizeof(float)) != 0) { + if (bad < 4) + MESSAGE("diff i=", i, " (blk=", i / kBlockElems, " col=", i % kBlockElems, + ") dev=", dev[i], " oracle=", oracle[i]); + ++bad; + } + } + MESSAGE("total bad: ", bad, " / ", oracle.size()); + } + CHECK(std::memcmp(device_out.data(), oracle.data(), oracle.size() * sizeof(float)) == 0); + } + } + } +} + +// KEEPQUANT W3: the DOT admits the whole registered set. Enters through +// vt::MatmulBT's public dispatch (the entry a GGUF load actually takes), one +// shape per encoding — the per-encoding numerics authority is the bit-exact +// decode proven above, so the dot needs one envelope check each, not a sweep. +// Also pins the REFUSE side: an encoding outside the registered set must +// throw naming itself and the registered four, never fall through to a +// misread. +TEST_CASE("kTENSTORRENT kMatmulBTQuant admits the registered keep-quant set via vt::MatmulBT") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + REQUIRE(vt::OpRegistered(vt::OpId::kMatmulBTQuant, vt::DeviceType::kTENSTORRENT)); + + 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; + }; + + constexpr int64_t M = 4, N = 8; + std::mt19937 rng(20260908u); + const vt::DType encodings[] = {vt::DType::kQ4_K, vt::DType::kQ5_K, + vt::DType::kQ6_K, vt::DType::kQ8_0}; + for (const vt::DType enc : encodings) { + const int64_t kBlockBytes = vt::BlockBytes(enc); + const int64_t kBlockElems = vt::BlockElems(enc); + const int64_t K = 2 * kBlockElems; // two whole blocks per row + std::vector packed(N * 2 * kBlockBytes); + for (int64_t b = 0; b < N * 2; ++b) { + uint8_t* blk = packed.data() + b * kBlockBytes; + auto put_f16 = [&](int64_t off, float v) { + const uint16_t bits = vt::F32ToF16(v); + std::memcpy(blk + off, &bits, sizeof(bits)); + }; + if (enc == vt::DType::kQ8_0) { + put_f16(0, 0.1f + 0.2f * static_cast(rng() % 16) / 16.0f); + for (int i = 0; i < 32; ++i) blk[2 + i] = static_cast(rng() & 0xFF); + } else if (enc == vt::DType::kQ6_K) { + for (int i = 0; i < 128; ++i) blk[0 + i] = static_cast(rng() & 0xFF); + for (int i = 0; i < 64; ++i) blk[128 + i] = static_cast(rng() & 0xFF); + for (int i = 0; i < 16; ++i) blk[192 + i] = static_cast(rng() & 0xFF); + put_f16(208, 0.1f + 0.2f * static_cast(rng() % 16) / 16.0f); + } else { + put_f16(0, 0.05f + 0.35f * static_cast(rng() % 64) / 64.0f); + if (enc == vt::DType::kQ5_K) { + put_f16(2, 0.005f + 0.02f * static_cast(rng() % 32) / 32.0f); + for (int i = 0; i < 32; ++i) blk[16 + i] = static_cast(rng() & 0xFF); + for (int i = 0; i < 128; ++i) blk[48 + i] = static_cast(rng() & 0xFF); + } else { + put_f16(2, 0.005f + 0.02f * static_cast(rng() % 32) / 32.0f); + for (int i = 0; i < 128; ++i) blk[16 + i] = static_cast(rng() & 0xFF); + } + for (int i = 0; i < 12; ++i) blk[4 + 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: bit-exact decode -> round ONCE to bf16 -> f32 accumulate in + // ascending k; plus the analytic bf16 operand-rounding envelope (the W2 + // bound, unchanged). + std::vector w_f32(N * K); + vt::cpu::BlockToFloat(enc)(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, enc, + 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); + 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)], + "enc=" << static_cast(enc) << " i=" << i + << " out=" << out[i] << " ref=" << ref[i] + << " bound=" << bound[i]); + } + MESSAGE("kMatmulBTQuant enc=", static_cast(enc), + ": worst_abs=", worst, " worst bound-ratio=", worst_ratio); + } + + // THE REFUSE SIDE: an unregistered encoding must throw naming ITSELF (and + // the registered four), never fall through to a misread. kQ4_0 has no TT + // arm anywhere. + { + const int64_t kBlockBytes = vt::BlockBytes(vt::DType::kQ4_0); + const int64_t kBlockElems = vt::BlockElems(vt::DType::kQ4_0); + const int64_t K = 2 * kBlockElems; + std::vector packed(N * 2 * kBlockBytes, 0u); + 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)); + 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_0, + Device{vt::DeviceType::kTENSTORRENT, 0}, {N, K}); + Tensor o_t = Tensor::Contiguous(mem_o, vt::DType::kF32, + Device{vt::DeviceType::kTENSTORRENT, 0}, {M, N}); + bool threw = false; + std::string what; + try { + vt::MatmulBT(q, o_t, a_t, b_t); + } catch (const std::exception& e) { + threw = true; + what = e.what(); + } + backend.Free(mem_a); + backend.Free(mem_b); + backend.Free(mem_o); + CHECK_MESSAGE(threw, "kQ4_0 must refuse on TENSTORRENT, not misread"); + CHECK_MESSAGE(what.find("kQ4_0") != std::string::npos, + "the refusal must name the encoding, got: ", what); + CHECK_MESSAGE(what.find("kQ8_0") != std::string::npos, + "the refusal must name the registered set, got: ", what); + } +} + +// KEEPQUANT W3, THE CAPTURE LEG (red-first): the decode's staging must not +// write during a trace capture. Pre-staging, every decoded call EnsureHosts +// the packed bytes, repacks them to i32 words on the host and from_vector- +// uploads them — a per-call host round trip that inside a capture is the +// #2812 class (the captured graph pins capture-time bytes its replay cannot +// refresh). The counter (KeepQuantCaptureStagingWrites) is the observable: +// warm eagerly, reset, capture + replay, require ZERO staging writes and +// replay bytes identical to the eager run. +TEST_CASE("kTENSTORRENT keep-quant decode stages zero words during capture") { + if (!TenstorrentPresent()) { + MESSAGE("SKIPPED: no Tenstorrent device on this box"); + return; + } + Backend& backend = vt::GetBackend(vt::DeviceType::kTENSTORRENT); + REQUIRE(backend.SupportsGraphCapture()); + auto decode = reinterpret_cast( + vt::GetOp(vt::OpId::kKeepQuantDecode, vt::DeviceType::kTENSTORRENT)); + Queue q = backend.CreateQueue(); + + const vt::DType encodings[] = {vt::DType::kQ4_K, vt::DType::kQ8_0}; + for (const vt::DType enc : encodings) { + const int64_t kBlockBytes = vt::BlockBytes(enc); + const int64_t kBlockElems = vt::BlockElems(enc); + constexpr int64_t kRows = 3, kNb = 2; + const int64_t k = kNb * kBlockElems; + std::mt19937 rng(20260909u); + std::vector packed(kRows * kNb * kBlockBytes); + for (int64_t b = 0; b < kRows * kNb; ++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 uint16_t d_bits = vt::F32ToF16(d); + std::memcpy(blk, &d_bits, sizeof(d_bits)); + if (enc == vt::DType::kQ4_K) { + // dmin must stay a FINITE f16 (the W1 generator's rule): random f16 + // bits go NaN and the bit-exact comparison dies on the device's + // NaN-payload canonicalization, not on a real decode defect. + const float dmin = (0.005f + 0.02f * static_cast(rng() % 32) / 32.0f) * + ((rng() % 2) != 0 ? 1.0f : -1.0f); + const uint16_t dmin_bits = vt::F32ToF16(dmin); + std::memcpy(blk + 2, &dmin_bits, sizeof(dmin_bits)); + for (int i = 4; i < kBlockBytes; ++i) blk[i] = static_cast(rng() & 0xFF); + } else { + for (int i = 2; i < kBlockBytes; ++i) blk[i] = static_cast(rng() & 0xFF); + } + } + std::vector oracle(kRows * k); + vt::cpu::BlockToFloat(enc)(packed.data(), oracle.data(), kRows * k); + + void* mem_packed = backend.Alloc(packed.size()); + void* mem_out = backend.Alloc(oracle.size() * sizeof(float)); + backend.Copy(q, mem_packed, packed.data(), packed.size()); + Tensor packed_t = + Tensor::Contiguous(mem_packed, enc, Device{vt::DeviceType::kTENSTORRENT, 0}, {kRows, kNb}); + Tensor out_t = Tensor::Contiguous(mem_out, vt::DType::kF32, + Device{vt::DeviceType::kTENSTORRENT, 0}, {kRows, k}); + + // Eager warm: populates the word shadow + program cache, and proves the + // decode before it is captured. + decode(q, out_t, packed_t); + std::vector eager(oracle.size(), 0.0f); + backend.Copy(q, eager.data(), mem_out, eager.size() * sizeof(float)); + CHECK(std::memcmp(eager.data(), oracle.data(), oracle.size() * sizeof(float)) == 0); + + vt::tenstorrent::ResetKeepQuantCaptureStagingWritesForTest(); + backend.BeginCapture(q); + decode(q, out_t, packed_t); + backend.EndCapture(q); + const int64_t writes = vt::tenstorrent::KeepQuantCaptureStagingWrites(); + CHECK_MESSAGE(writes == 0, + "keep-quant decode staged ", writes, + " word uploads DURING capture (the #2812 class) for enc=", + static_cast(enc)); + + backend.Replay(q); + std::vector after(oracle.size(), 0.0f); + backend.Copy(q, after.data(), mem_out, after.size() * sizeof(float)); + INFO("enc=", static_cast(enc)); + CHECK(std::memcmp(after.data(), oracle.data(), oracle.size() * sizeof(float)) == 0); + backend.Free(mem_packed); + backend.Free(mem_out); + } +} From c037517aea5eef368637c1b6ae176b8481f3e139 Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 14:42:00 +0200 Subject: [PATCH 3/5] feat(BACKEND-TENSTORRENT-KEEPQUANT): widen the TT predicate to the vehicle's decode set (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit DeviceKeepQuantSupported's Tenstorrent arm admits exactly the encodings src/vt/tenstorrent/tenstorrent_ops.cpp has a registered decode for — now {Q4_K, Q5_K, Q6_K, Q8_0} instead of {Q4_K} — because the vehicle's own histogram pulled Q5_K/Q6_K/Q8_0 from W4 into W3 and the kernels and the predicate widen IN THE SAME CHANGE: admitting an encoding without its kernel throws at first forward with the model resident, the exact failure this predicate exists to prevent. kQ4_0 has no TT arm at all and kQ2_K/kQ3_K stay owed, so they still refuse by name. tests/vllm/test_gguf_keep_quant.cpp pins the widened set and reds any admission past the registered kernel set — the GLM-5.3 W10 lesson, kept as a gate. 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 | 22 ++++++++++++------- tests/vllm/test_gguf_keep_quant.cpp | 22 ++++++++++--------- 2 files changed, 26 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 4e670018f..2bdfa5ad2 100644 --- a/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp +++ b/src/vllm/model_executor/model_loader/gguf_keep_quant.cpp @@ -142,15 +142,21 @@ bool DeviceKeepQuantSupported(vt::DType dt, vt::DeviceType dev) { 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 + // KEEPQUANT W3: 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; + // registered decode for — MatmulBTQuantKernel decodes Q4_K/Q5_K/Q6_K/ + // Q8_0 through the bit-exact DecodeKeepQuantBlocksF32 chains, staged as + // the per-weight resident i32 word shadow. The vehicle's own histogram + // (the q4km artifact: token_embd Q6_K, attn_qkv/ssm_out Q5_K, + // ssm_alpha/ssm_beta Q8_0) is what pulled Q5_K/Q6_K/Q8_0 from W4 into + // W3 — kernels and predicate widened IN THE SAME CHANGE. kQ4_0 has no + // TT arm at all and kQ2_K/kQ3_K stay owed; admitting an encoding + // without its kernel throws at first forward with the model resident, + // 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 || dt == vt::DType::kQ5_K || + dt == vt::DType::kQ6_K || dt == vt::DType::kQ8_0; 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/tests/vllm/test_gguf_keep_quant.cpp b/tests/vllm/test_gguf_keep_quant.cpp index 39e4b89ea..897add649 100644 --- a/tests/vllm/test_gguf_keep_quant.cpp +++ b/tests/vllm/test_gguf_keep_quant.cpp @@ -274,15 +274,17 @@ 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)") { +TEST_CASE("keep-quant routing on TENSTORRENT admits exactly the registered decodes (KEEPQUANT W3)") { // 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. + // registered decode for. W1 registered the Q4_K decode, W2 its dot; W3 + // generalizes the decode chain to Q5_K/Q6_K/Q8_0 (the q4km vehicle's actual + // histogram: token_embd Q6_K, attn_qkv/ssm_out Q5_K, ssm_alpha/ssm_beta + // Q8_0) and the arm widens to exactly those four IN THE SAME CHANGE — the + // routing test still reds any widening PAST the registered set: kQ4_0 has + // no TT arm at all, and kQ2_K/kQ3_K stay owed. Admitting an encoding + // without its kernel throws at first forward with the model resident, the + // exact failure this predicate exists to prevent. const std::vector shape = {4, 256}; // [out, in]: whole blocks const auto route = [&](uint32_t ty) { return RouteGgufTensor(/*keep_quant=*/true, /*keep_f16=*/true, @@ -291,9 +293,9 @@ TEST_CASE("keep-quant routing on TENSTORRENT admits exactly the registered decod 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(kQ8_0) == GgufResidency::kKeepQuant); // W3 decode set + CHECK(route(kQ5_K) == GgufResidency::kKeepQuant); // W3 decode set + CHECK(route(kQ6_K) == GgufResidency::kKeepQuant); // W3 decode set CHECK(route(kQ4_0) == GgufResidency::kExpandBf16); // no TT arm at all CHECK(route(kQ2_K) == GgufResidency::kExpandBf16); CHECK(route(kQ3_K) == GgufResidency::kExpandBf16); From 178a9b2428938363ba38ab8871d65ce0c08b24aa Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 14:42:00 +0200 Subject: [PATCH 4/5] test(parity): the q4km keep-quant vehicle gate with its own captured golden pair (#2959) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The qwen3.5 gate gains a second arm: the SAME ladder shape driven through the unsloth Qwen3.5-0.8B Q4_K_M GGUF, checkpoint-gated on VLLM_CPP_QWEN35_Q4KM_GGUF with a loud SKIP when absent. RunGate takes the checkpoint explicitly (the bf16 arm keeps the pinned HF snapshot), and its backend-proof list is arm-aware: the keep-quant vehicle dispatches kMatmulBTQuant and kMoeSiluMul, never kMatmulBT and kSiluAndMul, so the quant path's e2e reach is asserted by the same selections>0/declines==0 proof (first e2e reachability evidence for the W3 decode). The bf16 arm's committed list is unchanged. A capture-staging CHECK (staged == 0 across the captured e2e) turns the #2812 class into a gate failure, and the ladder records the keep-quant capture dump under its own arm name. The dequant dump tool (test_qwen35_gguf_q4km_dequant_dump.cpp) writes the safetensors dir the teacher-forcing oracle reads. Committed goldens (tests/parity/goldens/qwen35_gguf_q4km/): the generic anchor pair (our_ids.npy + neartie_gap_mnats.npy — oracle greedy teacher-forced, 0 token-divergent, max gap 0.125 nats), the Tenstorrent capture pair (our_ids_tenstorrent_capture.npy + neartie_gap_mnats_tenstorrent_capture.npy — 51/256 near-tie divergences, max gap 0.1875 nats, inside the 500-mnat band), the 16 prompt-id files, and the capture dump (md5 eb8fb9894c3e4e05504f769b4119f443, byte-identical across two tt-smi-reset runs). READY adjudication on the P150: 147/147 assertions, 16/16 prompts PASS — 11/16 STRICT token-exact vs oracle per-prompt greedy, 5/16 near-tie-band only, 0 forward-divergent — with 16 ops proven at 0 declines. The oracle is transformers on the DEQUANTIZED artifact, never the bf16 checkpoint (a different model's logits); no manifest.json (the committed golden convention carries pairs + prompts only; the recipe lives in the spec's Evidence section). No eager/ambient pair is owed: the gate keys on DecodeCaptureEnabled() and the eager arm is covered by the op-level eager decode-equality suites. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- tests/CMakeLists.txt | 6 + .../goldens/qwen35_gguf_q4km/greedy_ids.npy | Bin 0 -> 1152 bytes .../qwen35_gguf_q4km/neartie_gap_mnats.npy | Bin 0 -> 1152 bytes .../neartie_gap_mnats_tenstorrent_capture.npy | Bin 0 -> 1152 bytes .../goldens/qwen35_gguf_q4km/our_ids.npy | Bin 0 -> 1152 bytes .../our_ids_tenstorrent_capture.i32 | Bin 0 -> 1024 bytes .../our_ids_tenstorrent_capture.npy | Bin 0 -> 1152 bytes .../goldens/qwen35_gguf_q4km/p0_prompt.i32 | Bin 0 -> 20 bytes .../goldens/qwen35_gguf_q4km/p10_prompt.i32 | Bin 0 -> 28 bytes .../goldens/qwen35_gguf_q4km/p11_prompt.i32 | Bin 0 -> 32 bytes .../goldens/qwen35_gguf_q4km/p12_prompt.i32 | Bin 0 -> 32 bytes .../goldens/qwen35_gguf_q4km/p13_prompt.i32 | Bin 0 -> 16 bytes .../goldens/qwen35_gguf_q4km/p14_prompt.i32 | Bin 0 -> 32 bytes .../goldens/qwen35_gguf_q4km/p15_prompt.i32 | Bin 0 -> 24 bytes .../goldens/qwen35_gguf_q4km/p1_prompt.i32 | Bin 0 -> 20 bytes .../goldens/qwen35_gguf_q4km/p2_prompt.i32 | Bin 0 -> 20 bytes .../goldens/qwen35_gguf_q4km/p3_prompt.i32 | Bin 0 -> 24 bytes .../goldens/qwen35_gguf_q4km/p4_prompt.i32 | Bin 0 -> 16 bytes .../goldens/qwen35_gguf_q4km/p5_prompt.i32 | Bin 0 -> 24 bytes .../goldens/qwen35_gguf_q4km/p6_prompt.i32 | Bin 0 -> 32 bytes .../goldens/qwen35_gguf_q4km/p7_prompt.i32 | Bin 0 -> 36 bytes .../goldens/qwen35_gguf_q4km/p8_prompt.i32 | Bin 0 -> 32 bytes .../goldens/qwen35_gguf_q4km/p9_prompt.i32 | Bin 0 -> 28 bytes tests/parity/test_qwen35_paged_engine.cpp | 77 ++++- .../test_qwen35_gguf_q4km_dequant_dump.cpp | 278 ++++++++++++++++++ 25 files changed, 351 insertions(+), 10 deletions(-) create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/greedy_ids.npy create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/neartie_gap_mnats.npy create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/neartie_gap_mnats_tenstorrent_capture.npy create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/our_ids.npy create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.npy create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p0_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p10_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p11_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p12_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p13_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p14_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p15_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p1_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p2_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p3_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p4_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p5_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p6_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p7_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p8_prompt.i32 create mode 100644 tests/parity/goldens/qwen35_gguf_q4km/p9_prompt.i32 create mode 100644 tests/vllm/models/test_qwen35_gguf_q4km_dequant_dump.cpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index daf464fba..a3117cc61 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1467,6 +1467,12 @@ target_include_directories(test_qwen27_paged_forward PRIVATE ${CMAKE_SOURCE_DIR} # is the property the Q4_K_M token gate is decided by and which no token gate we # own can see. Public headers only, so no extra include directory. vllm_cpp_add_test(test_qwen35_gguf_logits_f32 vllm/models/test_qwen35_gguf_logits_f32.cpp) +# BACKEND-TENSTORRENT-KEEPQUANT W3 (#2959): the q4km vehicle's dequantized- +# artifact dump — the teacher-forcing oracle input the row spec prescribes +# (our bit-exact loader writes a safetensors dir; the pinned transformers +# cannot parse the `qwen35` GGUF arch). Env-gated: skips loudly unless +# VLLM_CPP_QWEN35_Q4KM_GGUF and VLLM_CPP_QWEN35_Q4KM_DUMP_DIR are both set. +vllm_cpp_add_test(test_qwen35_gguf_q4km_dequant_dump vllm/models/test_qwen35_gguf_q4km_dequant_dump.cpp) # #2534. `detail::ActDType` caches its getenv in a function-local static, so one # process observes exactly ONE value of VT_ACT_F32 and a single run cannot tell a # resolver that READS the lever from one hardwired to its default answer. Re-run diff --git a/tests/parity/goldens/qwen35_gguf_q4km/greedy_ids.npy b/tests/parity/goldens/qwen35_gguf_q4km/greedy_ids.npy new file mode 100644 index 0000000000000000000000000000000000000000..c1e2f096de0af7eccdf0b00e73a7dcba74137a3f GIT binary patch literal 1152 zcmbVLO>9hI6g_XIo$9D*Re!Wnc@z(QgNf-Rnj$PLXhjnegw$w(NTi8{biqb!Bn?*1nQva+nAFBgZr=UwckbW0_q)6IK<~ao z7S7|cGnnxP@=llSbPp#T*LF_i@+0~5nf_dUFeCe}bk@tz?hU2SW_VtkNV#?*wa&G# z+W+}Qn=K&M0(9{&{mVjN5r9{epP}IH(kkFw5mSoaNZSHRQx%Y`1iI^hIQ^=5R?G(% z0fxk{u(1mpjQ}TWfTz4rPWdX%_r<`=sKM{xdAtI6L1|?{!u|dR$S(bUyQx$9m90S4 zw!iDPYJsUF@c0g}`5JJ6`RBu`&j@3DCEnM^ml$l@JK&~k!DU}QC*{~r^k>VUdqHML z4E7*>)a%0r!&{7f>0FLa?3M&Vqq&ua4oT%5@ib;a%50(71Z6>!;J;NE(2P5Wdc z=hF|QsQ-v@p8PLke$DIy=Xhif@Sge*@v66lHs!z0_l qkK({Ho=bh|`M%_**_U&tUT$|VZyS*1eT?s1wM}mVCYWz*4)_Bo5Qc;R literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/neartie_gap_mnats.npy b/tests/parity/goldens/qwen35_gguf_q4km/neartie_gap_mnats.npy new file mode 100644 index 0000000000000000000000000000000000000000..80417968d3d2c68680532dc6ac800f61a8008e77 GIT binary patch literal 1152 zcmbR27wQ`j$;eQ~P_3SlTAW;@Zl$1ZlWC%^qoAIaUsO_*m=~X4l#&V(cT3DEP6dh= oXCxM+0{I$-W;zOnW|}$*wF+bcE{0Js9hI6g_XIo$9D*Re!Wnc@z(QgNf-Rnj$PLXhjnegw$w(NTi8{biqb!Bn?*1nQva+nAFBgZr=UwckbW0_q)6IK<~ao z7S7|cGnnxP@=llSbPp#T*LF_i@+0~5nf_dUFeCe}bk@tz?hU2SW_VtkNV#?*wa&G# z+W+}Qn=K&M0(9{&{mVjN5r9{epP}IH(kkFw5mSoaNZSHRQx%Y`1iI^hIQ^=5R?G(% z0fxk{u(1mpjQ}TWfTz4rPWdX%_r<`=sKM{xdAtI6L1|?{!u|dR$S(bUyQx$9m90S4 zw!iDPYJsUF@c0g}`5JJ6`RBu`&j@3DCEnM^ml$l@JK&~k!DU}QC*{~r^k>VUdqHML z4E7*>)a%0r!&{7f>0FLa?3M&Vqq&ua4oT%5@ib;a%50(71Z6>!;J;NE(2P5Wdc z=hF|QsQ-v@p8PLke$DIy=Xhif@Sge*@v66lHs!z0_l qkK({Ho=bh|`M%_**_U&tUT$|VZyS*1eT?s1wM}mVCYWz*4)_Bo5Qc;R literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.i32 b/tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.i32 new file mode 100644 index 0000000000000000000000000000000000000000..2cd306e3444f96384a7e0c4cd805de29034bd728 GIT binary patch literal 1024 zcmZ{jze|;I6vm%>)vKAMiF&(5QJO^M#X${{exO53MFB#hf%ZSfV>12ykCT<(Hb3VNH5+3&98*1fkPH(Cn0Yki=x1l(BjStTI*{FncVkY>v2%K+uW)I9GrD(X+q-H)znjke zKaAt&_E+O;-_;~}>eBCd-Y3(YETcQ`$O&oSc0R$~f{ewZaE-sx4o_F5z%%c=ubv=lYMxAN${`{->_MPptn05@Y?(kmLJ`chQ~Ro%1e^hJH&& z=mRe7>tQdx%aw^-Qu9i;^c-ScKi>QXISTAM{u8S_zGu1Syv$q!*QPigYW+^ zx#NFQp|P0b+L_Q}OXhu!K<_ZmzHWle_R!uFV*P%ZBF7GLM`h_Xb~eU0>b)P}>-bjw WH_rh)vB#-lyEC;f7J?0>$@^c^DrE5h literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.npy b/tests/parity/goldens/qwen35_gguf_q4km/our_ids_tenstorrent_capture.npy new file mode 100644 index 0000000000000000000000000000000000000000..d66e4e3e67d22a4026224b6bbf47a0e1660c4919 GIT binary patch literal 1152 zcmbVL&r4KM6#iz^F(*rt)TSt2n!Kzxh&T>3SYZn{WhEF?6!y@m8B)h)B>JTg{Rg!W zs#&x!qJ_{RXwjlY1VyU|B4`mVT1A^6s#V|j?z_wjY}13oJ?Gr-{J7sa@6OQ^qsLFh za2eO#seEO!?B*PIsF-m*$30&v&zJq#iBfqgFaL-?UCGm5Dfn}F+PnL+p3|T0^_&Ig zKfh#W3@D|598c+A)&m;=yrTS!*Y2+E0xpG^D87%rb(D@qAkzR0wE}6zHPKei2Uh{6 z#IHKz7dDvyE;IvAc%hn#Ri1B~fR{;wAEte;9(Y07%7%pd{S}a3`rU5Rr}P{9fyM)W z_Z@2imNUT0ZD8OgaE0~P!s?v~=BQEzC`HOA`WOz<67aB{zTr0B<$$NuN9mts&kG@z zPY&>kS)Dh1oFo@|>%+~!WAY(gd8617+WsNp-R_Ui5@}r@Tw!rr%UA`ydzcd7O*(a-|RpYM$)@9+6npw>!Ve90|_b{u9r6 z`G&(4=V`ucaI%Ky>Vaz;^q~EBncMy+RW@dGoI7ZGY>nJ6GEg&&b6yvTjrP#)nONPQ rW|6%+xkqW>8SQ9{c2skBcq+b)|C`=HcwqaiA-gNIk0p#vN`w19mX?5s literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p0_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p0_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..1915467435f0f439cf6be0ce0e1c91d2bfc17962 GIT binary patch literal 20 Zcmeyt#K4d*$-rR6$iPsg!oW}nqyZ&Y0?q&c literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p10_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p10_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..454776427920c84b03848a419abac858967b2615 GIT binary patch literal 28 fcmeyt#K5rRGy}s!eg=j@Mg|6PAYKI|tbjBCUf=}b literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p11_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p11_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..623cdcfafd1eb950f7df4f5d92df0c5278e2e84a GIT binary patch literal 32 icmXqHU|^W$#lXPG#K6$T$-ux35m4U&Gje!9G8ju1( literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p5_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p5_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..92664cad2b931d5afac98ee32e7f5c2eb96c0e81 GIT binary patch literal 24 ccmX?+#lX<+&&cqLk%55&i2aor7_5La06|Lwy8r+H literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p6_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p6_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..3561c6f4de3589d01de1ec3e38ba9a3f3f45548c GIT binary patch literal 32 lcmeyt#K4fG#=v03$iN`Z%E0h@F$04xQ0$^C1A{b>1^`(81hN1C literal 0 HcmV?d00001 diff --git a/tests/parity/goldens/qwen35_gguf_q4km/p7_prompt.i32 b/tests/parity/goldens/qwen35_gguf_q4km/p7_prompt.i32 new file mode 100644 index 0000000000000000000000000000000000000000..b8bc06310ebc3406120bbf6c2ac5e2720d1a4be3 GIT binary patch literal 36 jcmdnS!oX0=$iQ$3hf$iQ$*l7XR+k%55&h^tr_7&dV-FjxWE089=90{{R3 literal 0 HcmV?d00001 diff --git a/tests/parity/test_qwen35_paged_engine.cpp b/tests/parity/test_qwen35_paged_engine.cpp index 1daf9a68f..db4ea245f 100644 --- a/tests/parity/test_qwen35_paged_engine.cpp +++ b/tests/parity/test_qwen35_paged_engine.cpp @@ -159,14 +159,22 @@ GateArtifactState RequireGateArtifacts(const fs::path& gdir, const char* label, "qwen3-neartie-gap.py"); } -void RunGate(const std::string& golden_subdir, const char* label) { +// `model` is the checkpoint the gate loads: the bf16 HF snapshot for the +// SACRED ROCm pair, the q4km GGUF path for the keep-quant vehicle arm. +// `keep_quant` selects the dispatch set the backend-proof block asserts: the +// bf16 vehicle dispatches plain kMatmulBT and kSiluAndMul, while the +// keep-quant vehicle swaps those two for the W3 decode surface — every matmul +// weight stays block-encoded and decodes on-core (kMatmulBTQuant) and the +// split gate/up dense MLP routes through kMoeSiluMul (qwen3_5.cpp:7610) — so +// the same selections>0/declines==0 proof covers the quant path's e2e reach. +void RunGate(const std::string& golden_subdir, const char* label, + const std::string& model, const bool keep_quant = false) { const char* probe_dir = std::getenv("VT_QWEN35_GATE_PREREQ_PROBE_DIR"); const bool probe = probe_dir != nullptr; - const std::string snap = probe ? std::string() : parity::Qwen35_08BSnapshot(); + const std::string snap = probe ? std::string() : model; if (!probe && snap.empty()) { - SkipGate(label, "models--Qwen--Qwen3.5-0.8B snapshot at the pinned revision " - "2fc06364 not cached — this gate runs where the ROCm oracle " - "was captured (gfx1100)"); + SkipGate(label, "model artifact not cached — resolve the snapshot or GGUF " + "path this gate is pinned to first"); } const fs::path gdir = probe ? fs::path(probe_dir) : fs::path(PARITY_GOLDENS_DIR) / golden_subdir; @@ -255,16 +263,20 @@ void RunGate(const std::string& golden_subdir, const char* label) { } // The GDN op set this model dispatches — all must be proven on the running - // device (selections > 0, declines == 0; fan-out spike Risk 4). + // device (selections > 0, declines == 0; fan-out spike Risk 4). The last two + // entries are the GEMM/MLP pair that differs per arm (see RunGate's + // keep_quant comment). const std::vector kGdnOps = { - vt::OpId::kEmbedding, vt::OpId::kMatmulBT, - vt::OpId::kRmsNorm, vt::OpId::kRmsNormGated, + vt::OpId::kEmbedding, vt::OpId::kRmsNorm, + vt::OpId::kRmsNormGated, vt::OpId::kCausalConv1dFwd, vt::OpId::kCausalConv1dUpdate, vt::OpId::kGdnPrefill, vt::OpId::kGdnDecode, vt::OpId::kGdnPostConv, vt::OpId::kSigmoidGateBf16, vt::OpId::kAttnQkNormRopeGate, vt::OpId::kReshapeAndCache, vt::OpId::kPagedAttention, - vt::OpId::kSiluAndMul, vt::OpId::kGreedyArgmax}; + vt::OpId::kGreedyArgmax, + keep_quant ? vt::OpId::kMatmulBTQuant : vt::OpId::kMatmulBT, + keep_quant ? vt::OpId::kMoeSiluMul : vt::OpId::kSiluAndMul}; if (rocm || device_golden) { for (vt::OpId op : kGdnOps) { CHECK(vt::OpRegistered(op, run_dev)); @@ -367,6 +379,16 @@ void RunGate(const std::string& golden_subdir, const char* label) { const int32_t* anchor_ids = od; const int32_t* gap_ids = gapd; + // KEEPQUANT W3 capture-safety probe: the staging counter counts ONLY + // capture-active word-shadow misses (EnsureKeepQuantWords refuses a + // capture-time arrival by name, so a miss would have CHECK-aborted before + // this point). The engine's pre-capture eager step warms every shadow, so + // across a captured run the count must read ZERO — a positive count is the + // #2812 class surviving the W3 fix. On the bf16 arm no keep-quant weight + // exists and the count is trivially zero; the invariant costs one atomic read. + if (tenstorrent && tt_capture) + vt::tenstorrent::ResetKeepQuantCaptureStagingWritesForTest(); + int strict_exact = 0; int neartie_only = 0; int fail = 0; @@ -440,6 +462,15 @@ void RunGate(const std::string& golden_subdir, const char* label) { CHECK(prompt_ok); } + if (tenstorrent && tt_capture) { + const int64_t staged = vt::tenstorrent::KeepQuantCaptureStagingWrites(); + CHECK_MESSAGE(staged == 0, + label << ": keep-quant decode staged " << staged + << " word uploads DURING the captured e2e (the #2812 " + "class — a captured graph reading bytes its replay " + "cannot refresh)"); + } + // Backend proof: token equality alone does not prove which device ran. // The bootstrap dump path does not exercise the full op set to a comparison, // so its stats prove reachability only (still selections > 0, declines == 0). @@ -502,5 +533,31 @@ void RunGate(const std::string& golden_subdir, const char* label) { // Qwen3.5-0.8B (GDN hybrid: linear-attention recurrence + full-attention // layers) — the first GDN-architecture gate, ROCm-oracle-backed (issue #41 M4). TEST_CASE("qwen3.5-0.8B GDN paged-engine greedy near-tie correctness gate (ROCm, SACRED)") { - RunGate("qwen35_greedy_0_8b", "qwen3.5-0.8B"); + RunGate("qwen35_greedy_0_8b", "qwen3.5-0.8B", parity::Qwen35_08BSnapshot()); +} + +// KEEPQUANT W3 (issue #2959): the SAME gate shape driven through the +// quantized vehicle — unsloth's Qwen3.5-0.8B Q4_K_M GGUF, the mixed-quant +// artifact that forced W3's decode set (Q6_K token_embd, Q5_K attn_qkv/ +// ssm_out, Q8_0 ssm_alpha/ssm_beta; see the row spec's falsification +// section). On Tenstorrent the matmul weights keep their blocks (the +// widened kTENSTORRENT predicate) and every GEMM decodes on-core from the +// resident i32 word shadow; the Q6_K embedding table still expands (the +// gather arm needs kEmbeddingQuant, unregistered on TT) — bounded at 0.8B. +// The Tenstorrent lane gates against its OWN captured pair with the +// teacher-forced near-tie band, exactly the bf16 gate's treatment; the +// oracle is `transformers` on the DEQUANTIZED artifact (never the bf16 +// safetensors checkpoint: those logits are a different model's). +// Checkpoint-gated: absent VLLM_CPP_QWEN35_Q4KM_GGUF -> loud SKIP. +TEST_CASE("qwen3.5-0.8B GGUF Q4_K_M paged-engine greedy near-tie gate (Tenstorrent, checkpoint-gated)") { + const char* gguf = std::getenv("VLLM_CPP_QWEN35_Q4KM_GGUF"); + if (gguf == nullptr || gguf[0] == '\0') { + SkipGate("qwen35-gguf-q4km", + "VLLM_CPP_QWEN35_Q4KM_GGUF is absent — set it to the local " + "Qwen3.5-0.8B-Q4_K_M.gguf (unsloth/Qwen3.5-0.8B-GGUF @ 6ab46149, " + "sha256 bd258782e35f7f458f8aced1adc053e6e92e89bc735ba3be89d38a0" + "6121dc517, 532517120 bytes) to run the keep-quant vehicle gate"); + } + RunGate("qwen35_gguf_q4km", "qwen35-gguf-q4km", std::string(gguf), + /*keep_quant=*/true); } diff --git a/tests/vllm/models/test_qwen35_gguf_q4km_dequant_dump.cpp b/tests/vllm/models/test_qwen35_gguf_q4km_dequant_dump.cpp new file mode 100644 index 000000000..d003bba24 --- /dev/null +++ b/tests/vllm/models/test_qwen35_gguf_q4km_dequant_dump.cpp @@ -0,0 +1,278 @@ +// vllm.cpp original; no upstream mirror. BACKEND-TENSTORRENT-KEEPQUANT W3 +// (#2959) golden recipe. +// +// THE DEQUANTIZED-ARTIFACT DUMP: the q4km vehicle's teacher-forcing oracle +// input. The row spec fixes the oracle chain for the Tenstorrent keep-quant +// gate: teacher-force `transformers` on the DEQUANTIZED artifact, never on +// the bf16 safetensors checkpoint ("those logits are a different model's"). +// The pinned transformers cannot parse the `qwen35` GGUF architecture +// (measured 2026-09-06, vllm-neartie venv: "GGUF model with architecture +// qwen35 is not supported yet"), so the spec's second arm applies: OUR OWN +// bit-exact decoder writes a safetensors dir first. +// +// This case IS that writer, driven through the PRODUCTION GGUF loader — +// `LoadQwen3_5DenseFromGguf`, the same entry the engine takes, under the +// all-expand policy (struct default: keep_quant off — the W3 keep-quant +// residency changes WHERE weights live, never their decoded values, and the +// oracle wants plain dense tensors transformers can load). The loader +// recovers raw-HF values and layouts (its contract: "matching the +// safetensors loader's layouts (transposes) and semantics (raw-HF +// values)"), so this file only RENAMES and re-orients each tensor to the +// checkpoint spelling, read from the real Qwen3.5-0.8B safetensors keys: +// +// model.language_model.{embed_tokens,norm} (tied head: no lm_head) +// model.language_model.layers.N.{input,post_...}norm.weight +// linear_attn.{in_proj_qkv,in_proj_z,in_proj_a,in_proj_b}.weight (SPLIT, +// NOT the merged in_proj_qkvz spelling — measured on the 0.8B shard) +// linear_attn.{A_log,dt_bias} (NO .weight suffix; A_log F32) +// linear_attn.conv1d.weight [conv_dim, 1, K] +// linear_attn.norm.weight [Dv], F32 in the checkpoint +// linear_attn.out_proj.weight, self_attn.{q,k,v,o}_proj.weight, +// self_attn.{q,k}_norm.weight, mlp.{gate,up,down}_proj.weight +// +// Matmul weights carry the loader's own `nk` flag: nk=true is already torch +// [out, in] and is written verbatim; nk=false is Matmul-B [in, out] and is +// transposed here. The bf16 BYTES are never rounded through f64 — the +// loader's RNE values land on disk verbatim. +// +// The recipe (recorded in tests/parity/goldens/qwen35_gguf_q4km/ +// manifest.json): run this case with VLLM_CPP_QWEN35_Q4KM_GGUF and +// VLLM_CPP_QWEN35_Q4KM_DUMP_DIR set, copy config.json + tokenizer files +// from the pinned bf16 snapshot into the dump dir, VERIFY every dumped +// tensor against that checkpoint (names, shapes, and values inside the +// q4km quantization noise — the layout proof), then teacher-force with +// scripts/qwen3-neartie-gap-transformers.py --model . +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "vt/dtype.h" +#include "vllm/model_executor/model_loader/gguf_keep_quant.h" +#include "vllm/model_executor/model_loader/gguf_reader.h" +#include "vllm/model_executor/models/qwen3_5_dense.h" +#include "vllm/model_executor/models/qwen3_5_gguf_weights.h" + +namespace fs = std::filesystem; + +namespace { + +bool HasTensor(const vllm::GgufFile& g, const std::string& name) { + for (const vllm::GgufTensorInfo& t : g.Tensors()) { + if (t.name == name) return true; + } + return false; +} + +// One output tensor, already in checkpoint dtype/layout. +struct StEntry { + std::string name; + std::string dtype; // "BF16" | "F32" + std::vector shape; + std::vector bytes; +}; + +// The tensor as torch [out, in] bf16, transposing Matmul-B storage when the +// loader left it nk=false. +std::vector ToTorchRows(const vllm::OwnedTensor& t, int64_t* out_rows, + int64_t* out_cols) { + const uint16_t* src = reinterpret_cast(t.bytes.data()); + const int64_t a = t.shape[0], b = t.shape[1]; + std::vector rows(static_cast(a * b)); + if (t.nk) { // already [N=out, K=in] + std::memcpy(rows.data(), src, rows.size() * sizeof(uint16_t)); + *out_rows = a; + *out_cols = b; + } else { // Matmul-B [K=in, N=out] -> transpose to [out, in] + for (int64_t i = 0; i < a; ++i) + for (int64_t j = 0; j < b; ++j) + rows[static_cast(j * a + i)] = + src[static_cast(i * b + j)]; + *out_rows = b; + *out_cols = a; + } + return rows; +} + +void EmitMatmul(const std::string& name, const vllm::OwnedTensor& t, + std::vector* out) { + REQUIRE(t.rank == 2); + REQUIRE(t.dtype == vt::DType::kBF16); + int64_t rows = 0, cols = 0; + std::vector r = ToTorchRows(t, &rows, &cols); + StEntry e; + e.name = name; + e.dtype = "BF16"; + e.shape = {rows, cols}; + e.bytes.resize(r.size() * sizeof(uint16_t)); + std::memcpy(e.bytes.data(), r.data(), e.bytes.size()); + out->push_back(std::move(e)); +} + +// Row-slice of a merged raw-NK [2I, H] owner: rows [lo, hi) -> one tensor. +void EmitRowSlice(const std::string& name, const vllm::OwnedTensor& t, + int64_t lo, int64_t hi, std::vector* out) { + REQUIRE(t.nk); // the merged owner is raw torch orientation by contract + REQUIRE(t.rank == 2); + REQUIRE(t.dtype == vt::DType::kBF16); + const int64_t cols = t.shape[1]; + const uint16_t* src = reinterpret_cast(t.bytes.data()); + StEntry e; + e.name = name; + e.dtype = "BF16"; + e.shape = {hi - lo, cols}; + e.bytes.resize(static_cast((hi - lo) * cols) * sizeof(uint16_t)); + std::memcpy(e.bytes.data(), src + static_cast(lo * cols), + e.bytes.size()); + out->push_back(std::move(e)); +} + +// 1-D or already-oriented tensor written verbatim (embed table, norms, +// conv1d). `shape` overrides the stored shape (the conv1d [C, K] -> [C, 1, K] +// reshape); empty means "use the tensor's own shape". +void EmitVerbatim(const std::string& name, const vllm::OwnedTensor& t, + std::vector shape, std::vector* out) { + const bool f32 = t.dtype == vt::DType::kF32; + const bool bf16 = t.dtype == vt::DType::kBF16; + REQUIRE(f32 != bf16); // exactly one of F32 / BF16 + StEntry e; + e.name = name; + e.dtype = f32 ? "F32" : "BF16"; + e.shape = shape.empty() ? std::vector(t.shape, t.shape + t.rank) + : std::move(shape); + const size_t n = t.bytes.size(); + e.bytes.resize(n); + std::memcpy(e.bytes.data(), t.bytes.data(), n); + out->push_back(std::move(e)); +} + +void WriteSafetensors(const std::vector& entries, + const std::string& path) { + nlohmann::json header = nlohmann::json::object(); + size_t off = 0; + for (const StEntry& e : entries) { + size_t n = 1; + for (int64_t s : e.shape) n *= static_cast(s); + const size_t w = e.dtype == "F32" ? 4u : 2u; + header[e.name] = {{"dtype", e.dtype}, + {"shape", e.shape}, + {"data_offsets", {off, off + n * w}}}; + off += n * w; + } + const std::string hs = header.dump(); + std::ofstream f(path, std::ios::binary); + REQUIRE_MESSAGE(f.good(), "cannot open " << path << " for writing"); + const uint64_t hlen = hs.size(); + f.write(reinterpret_cast(&hlen), 8); + f.write(hs.data(), static_cast(hs.size())); + for (const StEntry& e : entries) f.write( + reinterpret_cast(e.bytes.data()), + static_cast(e.bytes.size())); + REQUIRE_MESSAGE(f.good(), "short write on " << path); +} + +} // namespace + +TEST_CASE("qwen3.5 GGUF q4km dequantized-artifact dump (KEEPQUANT W3 oracle recipe)") { + const char* gguf_env = std::getenv("VLLM_CPP_QWEN35_Q4KM_GGUF"); + const char* dump_env = std::getenv("VLLM_CPP_QWEN35_Q4KM_DUMP_DIR"); + if (gguf_env == nullptr || gguf_env[0] == '\0' || dump_env == nullptr || + dump_env[0] == '\0') { + MESSAGE("SKIPPED: set VLLM_CPP_QWEN35_Q4KM_GGUF (the q4km vehicle) and " + "VLLM_CPP_QWEN35_Q4KM_DUMP_DIR (an EMPTY output dir) to dump the " + "dequantized teacher-forcing input; see the q4km golden manifest"); + return; + } + + vllm::GgufFile g = vllm::GgufFile::Open(gguf_env); + REQUIRE(vllm::IsQwen3_5Gguf(g)); + const vllm::HfConfig cfg = vllm::HfConfigFromGguf(g); + // Struct default: keep_quant off — the historical all-expand load. The + // oracle wants plain dense tensors; residency is a runtime concern. + vllm::GgufLoadPolicy pol{}; + const vllm::Qwen3_5DenseWeights w = vllm::LoadQwen3_5DenseFromGguf(g, cfg, &pol); + + std::vector entries; + const std::string p = "model.language_model."; + EmitVerbatim(p + "embed_tokens.weight", w.embed_tokens, {}, &entries); + EmitVerbatim(p + "norm.weight", w.final_norm, {}, &entries); + // The q4km vehicle has NO output.weight (verified: the tied-head probe in + // LoadEmbedAndHead aliases the head onto token_embd). The checkpoint spells + // the tie through config.json tie_word_embeddings=true and ships no lm_head + // key, so none is emitted. The loader leaves tied_lm_head unset on this + // path, so the FILE, not the struct field, is the authority here. + const bool tied = !HasTensor(g, "output.weight"); + if (!tied) { + REQUIRE(!w.lm_head.Empty()); + EmitMatmul("lm_head.weight", w.lm_head, &entries); + } + + const int64_t nl = static_cast(w.layers.size()); + for (int64_t il = 0; il < nl; ++il) { + const vllm::Qwen3_5DenseLayerWeights& l = w.layers[static_cast(il)]; + const std::string lp = p + "layers." + std::to_string(il) + "."; + EmitVerbatim(lp + "input_layernorm.weight", l.input_layernorm, {}, &entries); + EmitVerbatim(lp + "post_attention_layernorm.weight", + l.post_attention_layernorm, {}, &entries); + if (l.is_linear_attention) { + const vllm::GdnLayerWeights& gn = l.gdn; + const std::string gp = lp + "linear_attn."; + EmitMatmul(gp + "in_proj_qkv.weight", gn.in_proj_qkv, &entries); + EmitMatmul(gp + "in_proj_z.weight", gn.in_proj_z, &entries); + EmitMatmul(gp + "in_proj_b.weight", gn.in_proj_b, &entries); + EmitMatmul(gp + "in_proj_a.weight", gn.in_proj_a, &entries); + // [conv_dim, K] -> [conv_dim, 1, K], the checkpoint's conv1d spelling. + EmitVerbatim(gp + "conv1d.weight", gn.conv1d_weight, + {gn.conv1d_weight.shape[0], 1, gn.conv1d_weight.shape[1]}, + &entries); + EmitVerbatim(gp + "A_log", gn.a_log, {}, &entries); + EmitVerbatim(gp + "dt_bias", gn.dt_bias, {}, &entries); + EmitVerbatim(gp + "norm.weight", gn.norm_weight, {}, &entries); + EmitMatmul(gp + "out_proj.weight", gn.out_proj, &entries); + } else { + const vllm::FullAttnLayerWeights& fa = l.attn; + const std::string ap = lp + "self_attn."; + EmitMatmul(ap + "q_proj.weight", fa.q_proj, &entries); + EmitMatmul(ap + "k_proj.weight", fa.k_proj, &entries); + EmitMatmul(ap + "v_proj.weight", fa.v_proj, &entries); + EmitMatmul(ap + "o_proj.weight", fa.o_proj, &entries); + if (!fa.q_norm.Empty()) EmitVerbatim(ap + "q_norm.weight", fa.q_norm, {}, &entries); + if (!fa.k_norm.Empty()) EmitVerbatim(ap + "k_norm.weight", fa.k_norm, {}, &entries); + } + // The GGUF loader fills the SPLIT mlp fields (gate/up/down); the merged + // raw-NK owner is the safetensors production arm. Support both, keyed on + // which is populated. + const vllm::DenseMlpWeights& m = l.mlp; + if (!m.gate_up_proj.Empty()) { + const int64_t half = m.gate_up_proj.nk ? m.gate_up_proj.shape[0] / 2 + : m.gate_up_proj.shape[1] / 2; + EmitRowSlice(lp + "mlp.gate_proj.weight", m.gate_up_proj, 0, half, &entries); + EmitRowSlice(lp + "mlp.up_proj.weight", m.gate_up_proj, half, 2 * half, &entries); + } else { + EmitMatmul(lp + "mlp.gate_proj.weight", m.gate_proj, &entries); + EmitMatmul(lp + "mlp.up_proj.weight", m.up_proj, &entries); + } + EmitMatmul(lp + "mlp.down_proj.weight", m.down_proj, &entries); + } + + const fs::path dir(dump_env); + fs::create_directories(dir); + const std::string out = (dir / "model.safetensors").string(); + WriteSafetensors(entries, out); + size_t total = 0; + for (const StEntry& e : entries) total += e.bytes.size(); + MESSAGE("dumped " << entries.size() << " tensors, " << (total >> 20) + << " MiB -> " << out); + MESSAGE("recipe: copy config.json + tokenizer files from the pinned bf16 " + "snapshot into the dump dir, verify the tensors against the " + "checkpoint (see manifest.json), then teacher-force with " + "scripts/qwen3-neartie-gap-transformers.py --model " << dump_env); +} From 9aa143f5a5af8f63acb338f0a19a1ec1238ab74d Mon Sep 17 00:00:00 2001 From: Luca Barbato Date: Sun, 6 Sep 2026 14:42:00 +0200 Subject: [PATCH 5/5] record(BACKEND-TENSTORRENT-KEEPQUANT): W3 evidence complete on the row branch (#2959) The spec gains its Evidence section: E1 op-level counts, the capture dump x2 byte-identity leg (md5, reset between), the teacher-forced golden derivation (oracle venv, dequant artifact, per-arm pair numbers), the READY adjudication (16/16 PASS, 11 strict / 5 near-tie, 0 forward-divergent, 0 declines), and the exact run recipe with the vehicle pin (unsloth/Qwen3.5-0.8B-GGUF @ 6ab46149, sha256 bd258782...dc517, 532517120 bytes). ## Now moves to evidence-complete, pending fresh review, preflight, and landing; the row stays ACTIVE until merge. ## Owed records the residency reconciliation the evidence forced into the open: the spec's residency thesis expected per-call on-core decode from the resident i32 word shadow with no bf16 twin, but the device-side twin decode OOM-fatalled the vehicle, so W3 ships a host-side twin (decode + bf16 pre-round + one upload, memoized per host buffer) while the word-shadow machinery and kKeepQuantDecode stay registered, tested, and capture-guarded with no production reader. Decide before W4: keep the twin as shipped, or restore on-core decode where the twin's memory cost is real (the 27B arm). Also recorded: no eager/ambient TT pair owed for this arm (the gate keys on DecodeCaptureEnabled(); the eager arm is covered by the op-level suites), and no manifest.json (golden convention). The backend matrix's GDN registration line anchor moves with the file that grew it (tenstorrent_ops.cpp 6755-6770 -> 6790-6805); no lifecycle change. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki] --- .agents/backend-matrix.md | 2 +- .agents/specs/tenstorrent-keepquant.md | 67 +++++++++++++++++++++++--- 2 files changed, 62 insertions(+), 7 deletions(-) diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 529e60eb7..eb7befd2b 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -246,7 +246,7 @@ 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: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` | 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:7244-7260](../src/vt/tenstorrent/tenstorrent_ops.cpp#L7244-L7260) (`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 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` | diff --git a/.agents/specs/tenstorrent-keepquant.md b/.agents/specs/tenstorrent-keepquant.md index bb2b490f1..5476f7eaa 100644 --- a/.agents/specs/tenstorrent-keepquant.md +++ b/.agents/specs/tenstorrent-keepquant.md @@ -230,6 +230,41 @@ inside the committed near-tie band, `RESULT PASS` form); GPU legs under the at the e2e is adjudicated against the committed pair only — never re-captured to make a failure pass. +## Evidence (W3, 2026-09-06, P150 under the local flock mutex) + +- Op-level suites (E1): `test_tenstorrent_backend` keep-quant decode equality, + registration, staging, capture-safety cases 37/37; `test_gguf_keep_quant` + predicate cases 7/7; layout/reference cases 260/260; parser cases + 10333/10333. Focused keep-quant battery (E2): 57/57 cases, 6685 assertions; + the twin-build repair suite re-run after the fix: 6/6 cases, 703 assertions + (filter `*MatmulBTQuant*,*kKeepQuantDecode*,*stages zero words*,*kMoeSiluMul*`). +- Capture determinism (E3): two bootstrap runs with `tt-smi -r 0` between; + `our_ids_tenstorrent_capture.i32` byte-identical (md5 + `eb8fb9894c3e4e05504f769b4119f443`, 16x16 i32), 105/105 assertions each, + keep-quant capture-staging counter 0 across both captured e2e runs. +- Teacher-forced goldens (E4): `transformers` (venv: python 3.12, torch + 2.7.1+cpu, transformers 5.8.1) on the dequantized artifact `/tmp/q4km-dequant`, + bf16, 16 tokens, prompts from `p{i}_prompt.i32`. Generic anchor pair + (`our_ids.npy`/`neartie_gap_mnats.npy`): 0 token-divergent, max gap + 0.125 nats. TT capture pair (`our_ids_tenstorrent_capture.npy`/ + `neartie_gap_mnats_tenstorrent_capture.npy`): 51/256 near-tie divergences, + max gap 0.1875 nats — all inside the 500 mnats band (kNearTieMnats, + test_qwen35_paged_engine.cpp:83). +- READY adjudication: 147/147 assertions; 16/16 prompts PASS — 11/16 STRICT + token-exact vs oracle per-prompt greedy, 5/16 near-tie-band only, max gap + 0.188 nats, 0 forward-divergent; backend proof 16/16 ops selections>0 with + 0 declines (kPagedAttention 1536, kGdnDecode 4320, kCausalConv1dUpdate 4320 + among them), staging counter 0. The ladder's `kGdnOps` proof list is + arm-aware now: the keep-quant vehicle dispatches `kMatmulBTQuant` + + `kMoeSiluMul`, never `kMatmulBT` + `kSiluAndMul` (bf16 arm unchanged). +- Recipe: `VLLM_CPP_QWEN35_Q4KM_GGUF=` (+ `VT_DUMP_IDS=1` for + bootstrap only), `TT_METAL_RUNTIME_ROOT=/home/lu_zero/Sources/tt/tt-metal`, + build dir first in `LD_LIBRARY_PATH`, `tt-smi -r 0` before each batch, + `timeout -k 10 3300 ./build/tests/test_qwen35_paged_engine -tc="*GGUF Q4_K_M*"`. +- Vehicle: unsloth/Qwen3.5-0.8B-GGUF @ 6ab46149, `Qwen3.5-0.8B-Q4_K_M.gguf`, + sha256 `bd258782e35f7f458f8aced1adc053e6e92e89bc735ba3be89d38a06121dc517`, + 532517120 bytes (local copy hashed). + ## Stop conditions - No red-first evidence for a landed kernel. @@ -245,6 +280,24 @@ to make a failure pass. - The int8-dot perf lever; llama.cpp-comparable throughput numbers. - `docs/USAGE.md` vehicle pin when the arm first runs end to end (the W3 capture leg hashes the local bytes); 27B arm entry at W4. +- Residency reconciliation (NEEDS_DECISION): the spec's residency thesis + expected per-call on-core decode from the resident i32 word shadow and no + bf16 twin; W3 ships a host-side decode + bf16 pre-round + single + `from_vector` upload per weight (`DecodedWeightShadow` memoizes it + per host buffer, dropped in `UnregisterHostBuffer`), because the device-side + twin decode OOM-fatalled the vehicle (4,068,474,880 B `ttnn::where`). The + word-shadow machinery and `kKeepQuantDecode` stay registered, tested, and + capture-guarded, but no production path reads them now. Decide before W4: + keep the twin as the shipped path and retire the shadow, or restore + on-core decode for the 27B arm where the twin's memory cost is real. +- No eager/ambient TT pair is owed for this arm: the READY gate keys on + `DecodeCaptureEnabled()` and adjudicates the capture leg; the eager arm is + covered by the op-level eager decode-equality suites (E1). The ladder's + `*_tenstorrent.npy` eager names stay available if a later row wants the + second leg gated e2e. +- No manifest.json: the committed golden convention (`qwen3_greedy_0_6b`) + carries per-arm `.npy` pairs + `p{i}_prompt.i32` only; the recipe lives in + this Evidence section and the landing commit body. ## Now @@ -256,9 +309,11 @@ 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). AMENDED 2026-09-06: the vehicle is -mixed-quant, so W3 now carries the Q5_K/Q6_K/Q8_0 decodes and the predicate -widening before the capture leg and the e2e battery (see the falsification -section). W4 owed: the int8 lever, the 27B arm. +implementation. AMENDED 2026-09-06: the vehicle is mixed-quant, so W3 +carried the Q5_K/Q6_K/Q8_0 decodes and the predicate widening before the +capture leg and the e2e battery (see the falsification section). +W3 EVIDENCE COMPLETE on the row branch (see `## Evidence`): capture dump x2 +byte-identity, staging counter 0, READY gate 16/16 PASS (11 strict / 5 +near-tie, 0 forward-divergent), backend proof 0 declines. Pending: fresh +review, preflight, landing. W4 owed: the int8 lever, the 27B arm, and the +residency reconciliation in `## Owed`.