From f6d9c146040b48bba71a6e692586159e11185892 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 15:35:08 +0100 Subject: [PATCH 1/7] =?UTF-8?q?feat(qtip/bake):=20bake=20K=3D9/V=3D4/L=3D1?= =?UTF-8?q?2=20=E2=80=94=20the=20packed=20row=20becomes=20a=20bitstream?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit #169 made the beam encoder geometry-parametric and it does NOT cover K=9. Two `static_assert`s in `QtipGeom` refuse it outright — `K_ <= 8` and `8u % K_ == 0u` ("K must divide 8: the packer emits whole bytes") — and behind them `SYMS_PER_BYTE = 8/K` is 0, so `qtip_packed_bytes_per_row` divides by zero and `qtip_pack_symbol` cannot address a symbol that straddles two bytes. The Rust side has the same hole: `cuda_ops.rs` computed `8 / super::K` and would have panicked before a kernel ran. That is the whole gap. Everything the K=8 work was actually about survives K=9 untouched, which is worth stating precisely because it is also the cost argument: * `LANES = min(2^K, 256 / 2^(L-K))` gives 32 at K=9/L=12 (8 groups, 512-symbol alphabet), so `CAND = 512/32 = 16` — the SAME register array as K=4 (16/1) and K=8 (256/16). No `cand[512]`, no spill. * `n_cand = ng * 2^K` is 2^L-saturated at 4096 at all three geometries, so the per-timestep candidate count, the ~3.87 radix passes, the 32-bit cost key and the `2^K * V * 4 B`-per-group LUT traffic (65,536 B/timestep) are invariant. * `MAX_GROUPS` falls 16 -> 8 and the group table 128 B -> 64 B; the 16 KiB backtrace stage floor is unchanged. * `num_symbols = in_features / V` is K-independent. So the packing is the only thing that moves, and it moves to the format the serving side already reads: symbol `t` at bits `[t*K, t*K + K)`, LSB-first, row `ceil(num_symbols*K/8)` bytes — `trellis_v4l12.rs::Rung::{pack, extract, packed_bytes}`. No second convention is invented; the UQFF geometry section is `[tag=3, K, L, V]`, so K=8 -> K=9 is one byte of artifact change. The straddle bound is `ceil((8 - gcd(K,8) + K)/8)`, not `ceil((7+K)/8)`: the reachable bit offsets are multiples of `gcd(K,8)`, so 7 is unreachable at K=10 and two bytes are EXACT at both K=9 and K=10. The packer writes only the bytes a symbol's bits fall in, which is what keeps the last symbol of a row inside `ceil(num_symbols*K/8)`. K=4/V=2/L=16 STAYS BYTE-IDENTICAL, and not by assertion: the byte-aligned expression is kept verbatim under `if constexpr`, so K=4 and K=8 compile from unchanged source rather than from a general form believed to fold back to it. Five host tests pin the rule the discarded branch would produce (the bitstream reproduces the shipped nibble packing over 4096 symbols; it is a plain byte store at K=8; K=9 round-trips through an independently written extractor at the real 1792-symbol V4 shape including the last symbol; the straddle bound is exact at K=4/8/9/10; a partial trailing byte is refused, and at K=4 that rule is still exactly "num_symbols must be even"). A `const` block asserts at COMPILE time that the general formula reproduces `num_symbols/2` at K=4 and `num_symbols` at K=8 at the real V4 shapes. All five were proven red: M1 (packed_bytes_per_row assumes one byte per symbol) fails to compile on the const assert; M2 (one-byte extract) reds only the K=9 test; M3 (whole-byte guard always true) reds only the partial-byte test; M4 (naive ceil((7+K)/8)) reds only the straddle test at K=10; M5 (nibble order flipped) reds only the K=4 identity test. `arc-tools/qtip_beam_res_usage_check.sh` now asserts SIX beam kernels, not four — a "LOCAL:0" verdict must be about a build that still contains the K=9 instantiation, which is precisely the one a naive parameterisation would spill. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- arc-tools/qtip_beam_res_usage_check.sh | 29 +-- mistralrs-quant/kernels/qtip/qtip_beam.cu | 48 +++-- mistralrs-quant/kernels/qtip/qtip_geom.cuh | 121 ++++++++++++- mistralrs-quant/src/qtip/cuda_ops.rs | 27 +-- mistralrs-quant/src/qtip/mod.rs | 196 +++++++++++++++++++++ 5 files changed, 379 insertions(+), 42 deletions(-) diff --git a/arc-tools/qtip_beam_res_usage_check.sh b/arc-tools/qtip_beam_res_usage_check.sh index 829b855b2..4c7b3f480 100755 --- a/arc-tools/qtip_beam_res_usage_check.sh +++ b/arc-tools/qtip_beam_res_usage_check.sh @@ -11,12 +11,13 @@ # name in their source ("`LOCAL:` must remain 0") and, until now, nothing ran it. # # It is also the check that separates the two competing cost models for the -# K=8/V=4/L=12 encode. A naive constant bump turns `cand[2^K]` into `cand[256]`, -# which spills for certain — that spilled kernel is where the discredited -# "~1,700 s/layer / 20 h / $30" figure came from. The lane-team re-blocking in -# `qtip_beam.cu` keeps `cand[]` at 16 entries at BOTH geometries; this script is -# what turns that claim into a compiled fact instead of a hand count. (Hand -# counting C++ undercounts SASS by ~2.05x on this kernel family — do not.) +# V=4/L=12 encode. A naive constant bump turns `cand[2^K]` into `cand[256]` at +# K=8 and `cand[512]` at K=9, which spills for certain — that spilled kernel is +# where the discredited "~1,700 s/layer / 20 h / $30" figure came from. The +# lane-team re-blocking in `qtip_beam.cu` keeps `cand[]` at 16 entries at EVERY +# instantiated geometry; this script is what turns that claim into a compiled +# fact instead of a hand count. (Hand counting C++ undercounts SASS by ~2.05x +# on this kernel family — do not.) # # THIS GATE IS PROVEN RED ON EVERY RUN # ------------------------------------ @@ -27,10 +28,13 @@ # dynamically-indexed 256-float local array is compiled and pushed through # the SAME parser. If it is not reported as a spill, this script fails — # the gate has lost its teeth and says so. -# 2. A KERNEL-COUNT assertion per source. `qtip_beam.cu` must yield 4 beam -# kernels (2 geometries x 2 codebooks); if the K=8/V=4/L=12 instantiation -# were ever dropped, a "no spill" result would be about a kernel that no -# longer exists. `qtip2b_beam.cu` yields 1: that rung has no stored-LUT +# 2. A KERNEL-COUNT assertion per source. `qtip_beam.cu` must yield 6 beam +# kernels (3 geometries — K=4/V=2/L=16, K=8/V=4/L=12, K=9/V=4/L=12 — x 2 +# codebooks); if an instantiation were ever dropped, a "no spill" result +# would be about a kernel that no longer exists. That is not hypothetical: +# K=9 is the rung the bake targets and it is the one whose `cand[]` a naive +# parameterisation would blow up to 512. `qtip2b_beam.cu` yields 1: that +# rung has no stored-LUT # path at all (the codebook is always computed from `mult`), so its kernel # is not templated on the codebook the way the LUT rung's is. # @@ -120,11 +124,12 @@ echo # --------------------------------------------------------------------------- # 1. The shipped beam kernels. -# qtip_beam.cu : 2 geometries (K4/V2/L16, K8/V4/L12) x 2 codebooks = 4 +# qtip_beam.cu : 3 geometries (K4/V2/L16, K8/V4/L12, K9/V4/L12) x 2 +# codebooks = 6 # qtip2b_beam.cu : 1 kernel — computed codebook only, so no codebook template # --------------------------------------------------------------------------- SOURCES=(qtip_beam.cu qtip2b_beam.cu) -EXPECT=(4 1) +EXPECT=(6 1) fail=0 for i in "${!SOURCES[@]}"; do diff --git a/mistralrs-quant/kernels/qtip/qtip_beam.cu b/mistralrs-quant/kernels/qtip/qtip_beam.cu index 96f3e8705..11a455cb9 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -8,10 +8,15 @@ // // GEOMETRY IS A TEMPLATE PARAMETER (see qtip_geom.cuh) // --------------------------------------------------- -// Two rungs are compiled, both at 2 bits/weight: +// Three rungs are compiled: // -// K=4 / V=2 / L=16 the shipped rung; 15.125 decode instructions/weight -// K=8 / V=4 / L=12 the same rate at 4.375 decode instructions/weight +// K=4 / V=2 / L=16 2.00 bpw, the shipped rung; 15.125 decode inst/weight +// K=8 / V=4 / L=12 2.00 bpw at 4.375 decode inst/weight +// K=9 / V=4 / L=12 2.25 bpw — THE BAKE TARGET. K=8/V=4/L=12 is quality- +// CLOSED (-0.00698 with a random codebook, -0.00307 with +// a converged trellis-Lloyd one, against a +-0.0008 ship +// band, over six codebook designs); K=9 measures +// +0.00402, 5x better than the shipped control. // // The K=8 rung is NOT reachable by bumping the constants. Three things in this // kernel are shaped by `2^(L-K)` and `2^K`, and at K=8/L=12 all three invert: @@ -30,17 +35,29 @@ // // K=4/L=16: 4096 groups >= 256 threads -> LANES = 1, cand[16/1 ] = 16 // K=8/L=12: 16 groups -> LANES = 16, cand[256/16] = 16 +// K=9/L=12: 8 groups -> LANES = 32, cand[512/32] = 16 // -// The block stays fully occupied at both geometries and the register array -// stays 16 entries at both — the candidate count per timestep is `2^L`- -// saturated (`ng * 2^K = 4096`) either way, so the SAME work is simply +// The block stays fully occupied at all three geometries and the register array +// stays 16 entries at all three — the candidate count per timestep is `2^L`- +// saturated (`ng * 2^K = 4096`) every time, so the SAME work is simply // re-blocked. `LANES == 1` collapses every expression below to the shipped // code, which is why the K=4 rung stays byte-identical. // +// 🔑 THAT INVARIANCE IS ALSO THE K=9 COST ARGUMENT. Going K=8 -> K=9 at fixed +// V=4/L=12 changes NO per-timestep quantity in this kernel: same 4096 +// candidates, same cand[16], same ~3.87 radix passes over the same 32-bit cost +// key, same V=4 branch metric, same `2^K * V * 4 B` of LUT per group over half +// as many groups (65,536 B per timestep either way), same `beam_w` trace +// words, and `num_symbols = in_features / V` is K-independent. The one term +// that does move is the packer: one byte store per symbol at K=8 becomes two +// byte read-modify-writes in thread 0's serial backtrace. So K=9 s/layer is +// K=8 s/layer plus that term — see the encode-cost note below. +// // The third is handled by sizing the shared scratch as // `max(2^(L-K) u64, STAGE_U64_MIN)`: at K=4 that is exactly the 4096-entry // group table that already staged the backtrace (no change at all), and at K=8 -// it is a 16 KiB stage over a 128 B table. That is not a smaller staging +// (128 B table) and K=9 (64 B table) it is the same 16 KiB stage over a table +// far too small to serve as one. That is not a smaller staging // budget in the terms that matter — the backtrace's dependent global loads are // `num_symbols / window`, and V=4 halves `num_symbols` exactly as the halved // window doubles it: 9472/32 == 4736/16 == 296 per row, unchanged. @@ -210,10 +227,11 @@ struct QbShape { static_assert(MAX_GROUPS <= GROUPS, "every live group must have a lane team in the block"); // `cand[]` MUST stay in registers; see the spill note on - // QB_MIN_BLOCKS_PER_SM. 16 is what both shipped geometries produce, and it - // is the number `cuobjdump -res-usage` has been checked at (LOCAL:0, - // REG<=62). A geometry that would push it higher must be measured before it - // is trusted, not reasoned about — hence the ceiling rather than a comment. + // QB_MIN_BLOCKS_PER_SM. 16 is what all three instantiated geometries + // produce (K=4: 16/1, K=8: 256/16, K=9: 512/32), and it is the number + // `cuobjdump -res-usage` has been checked at (LOCAL:0, REG<=62). A geometry + // that would push it higher must be measured before it is trusted, not + // reasoned about — hence the ceiling rather than a comment. static_assert(CAND >= 1 && CAND <= 16, "cand[] must stay small enough to live in registers"); // `keep_mask` is a 32-bit bitset over `cand[]`. @@ -763,6 +781,14 @@ int launch_qtip_quantize_rows_beam_geom_f32( d_weight, d_lut, d_row_scales, d_packed, d_trace, n_rows, in_features, num_symbols, row_offset, beam_w, cb_mult, stream); } + // 2.25 bpw. The caller must size `d_packed` as `ceil(num_symbols*9/8)` per + // row — `num_symbols / (8/K)` is a division by zero at this rung, which is + // why `qtip_packed_bytes_per_row` is no longer written that way. + if (k_bits == 9 && v_dim == 4 && l_bits == 12) { + return qb_launch( + d_weight, d_lut, d_row_scales, d_packed, d_trace, + n_rows, in_features, num_symbols, row_offset, beam_w, cb_mult, stream); + } return -2; } diff --git a/mistralrs-quant/kernels/qtip/qtip_geom.cuh b/mistralrs-quant/kernels/qtip/qtip_geom.cuh index 19d0a354a..555c92b31 100644 --- a/mistralrs-quant/kernels/qtip/qtip_geom.cuh +++ b/mistralrs-quant/kernels/qtip/qtip_geom.cuh @@ -14,6 +14,15 @@ // weight (4.375 vs 15.125), because one symbol now // carries four weights instead of two // +// and raising K at fixed V=4/L=12 is the ONLY way to buy quality on this +// family, because `bpw = K/V`: +// +// K=9 / V=4 / L=12 2.25 bits/weight, same 32,768 B codebook, same decode +// shape. This is the rung the encoder is being taught to +// bake; K=8/V=4/L=12 is quality-CLOSED (measured +// -0.00698 random / -0.00307 trellis-Lloyd against a +// +-0.0008 ship band). +// // Before this header each kernel restated `K`/`L`/`V` as its own local // `constexpr`, and every one of those restatements assumed K=4 somewhere else // in the file — in the packing (`|= sym << 4` for odd t), in a group table @@ -23,6 +32,27 @@ // explicit and `static_assert`-checked, so a geometry a kernel cannot serve // fails to COMPILE rather than silently emitting the wrong bytes. // +// THE BITSTREAM (this is a WIRE format, not a local convention) +// ------------------------------------------------------------- +// Symbol `t` occupies bits `[t*K, t*K + K)` of the packed row, LSB-first, and +// the row is `ceil(num_symbols * K / 8)` bytes. That is exactly what the +// serving side's `qtip/trellis_v4l12.rs::Rung::{pack, extract, packed_bytes}` +// implements, and the bake MUST NOT invent a second convention: the UQFF +// geometry section is `[tag=3, K, L, V]`, so K is an explicit wire field and +// the encoder and the decoder are reading the same bytes at the same offsets. +// +// At K=4 and K=8 the bitstream and the old "8/K whole symbols per byte" packer +// are the SAME bytes — `bit = 4t` puts the even symbol in the low nibble and +// the odd one in the high nibble, and `bit = 8t` is a plain byte store — which +// is why `qtip_pack_symbol` keeps the byte-aligned expression verbatim under +// `if constexpr` rather than trusting a general form to fold back to it. +// +// K=9 is where the two diverge: 8 % 9 != 0, so a symbol straddles two bytes, +// `SYMS_PER_BYTE` is meaningless (integer 8/9 == 0, and the old +// `num_symbols / SYMS_PER_BYTE` divided by zero), and the packer has to write +// `MAX_BYTES_PER_SYMBOL` bytes. See that member for why the bound is 2 and not +// 3. +// // WHAT IS NOT HERE // ---------------- // The Rust-side geometry (`qtip/mod.rs::{K, V, L}`) is still three separate @@ -48,10 +78,41 @@ struct QtipGeom { static constexpr uint32_t GROUP_COUNT = 1u << (L_ - K_); static constexpr uint32_t GROUP_MASK = GROUP_COUNT - 1u; static constexpr uint32_t LUT_SIZE = 1u << L_; // states, not floats - static constexpr uint32_t SYMS_PER_BYTE = 8u / K_; - static_assert(K_ >= 1 && K_ <= 8, "K must fit a byte-packable symbol"); - static_assert(8u % K_ == 0u, "K must divide 8: the packer emits whole bytes"); + // Does a whole number of symbols fit a byte? True at K=1,2,4,8; false at + // K=9. Only the true case has a meaningful `SYMS_PER_BYTE`; the false case + // is pinned to 0, so the only code that may read it is code guarded by + // `if constexpr (BYTE_ALIGNED)` — anything else is an obvious divide-by-zero + // rather than a quietly short row. (`8u / 9u` is 0 by integer division, not + // UB, so the discarded ternary arm is still well-formed.) + static constexpr bool BYTE_ALIGNED = (8u % K_ == 0u); + static constexpr uint32_t SYMS_PER_BYTE = BYTE_ALIGNED ? (8u / K_) : 0u; + + // gcd(K, 8). 8 is a power of two, so this is the largest power of two that + // divides K, capped at 8. + static constexpr uint32_t GCD8 = + (K_ % 8u == 0u) ? 8u : ((K_ % 4u == 0u) ? 4u : ((K_ % 2u == 0u) ? 2u : 1u)); + + // Bytes ONE symbol can touch, worst case over the bit offsets that are + // actually REACHABLE. + // + // Symbol `t` starts at bit offset `(t*K) mod 8`, which ranges over the + // multiples of `gcd(K, 8)` and nothing else — so the worst reachable offset + // is `8 - gcd(K, 8)`, not 7, and the bound is + // `ceil((8 - gcd(K,8) + K) / 8)`: + // + // K=8 -> gcd 8, offsets {0}, bound 1 (a plain byte store) + // K=9 -> gcd 1, offsets {0..7}, bound 2 + // K=10 -> gcd 2, offsets {0,2,4,6}, bound 2 (naive ceil((7+K)/8) + // says 3 and is wrong) + // + // This is the same formula as `Rung::max_bytes_per_symbol` on the serving + // side, and the same correction: `ceil((7+K)/8)` over-reads at K=10. + static constexpr uint32_t MAX_BYTES_PER_SYMBOL = (8u - GCD8 + K_ + 7u) / 8u; + + // K <= 15 rather than K <= 8: the symbol is no longer required to tile a + // byte, only to leave the trellis some state (`L_ > K_`) with `L_ <= 16`. + static_assert(K_ >= 1 && K_ <= 15, "K must leave state in a 16-bit trellis"); static_assert(L_ > K_, "the trellis must retain state across a symbol"); // `s_beam_state` is `unsigned short` and every selection key is laid out as // `(total_order_key(cost) << 16) | state`, so the state must fit 16 bits. @@ -63,24 +124,66 @@ struct QtipGeom { using QtipGeomK4V2L16 = QtipGeom<4, 2, 16>; // Same 2 bits/weight at 4.375 decode instructions/weight (3.46x fewer). using QtipGeomK8V4L12 = QtipGeom<8, 4, 12>; +// 2.25 bits/weight on the same table and the same decode shape. The rung the +// quality sweep actually selected; K=8/V=4/L=12 is closed. +using QtipGeomK9V4L12 = QtipGeom<9, 4, 12>; +// K=10/V=4/L=12 (2.5 bpw) needs no new code here — GROUP_COUNT falls to 4, so +// `QbShape::LANES` becomes 64 and `CAND` stays 16 — but it is deliberately NOT +// instantiated: a geometry that has not been through `-res-usage` is a claim, +// not a rung. -// Bytes of packed output one row of `num_symbols` symbols occupies. +// Bytes of packed output one row of `num_symbols` symbols occupies: +// `ceil(num_symbols * K / 8)`. +// +// The byte-aligned branch is the shipped expression, character for character, +// so the K=4 and K=8 instantiations compile from unchanged source rather than +// from a general form that is merely believed to fold back to it. The two +// agree wherever both are defined: at K=4 with `num_symbols` even (which +// `cuda_ops.rs` enforces) `(4n+7)/8 == n/2`, and at K=8 `(8n+7)/8 == n`. template __device__ __forceinline__ int qtip_packed_bytes_per_row(int num_symbols) { - return num_symbols / (int)G::SYMS_PER_BYTE; + if constexpr (G::BYTE_ALIGNED) { + return num_symbols / (int)G::SYMS_PER_BYTE; + } else { + return (num_symbols * (int)G::K + 7) / 8; + } } -// OR one symbol into the packed row. +// OR one symbol into the packed row at bits `[t*K, t*K + K)`, LSB-first. // // At K=4 this is exactly the shipped `pkd[t/2] |= sym` for even `t` and // `pkd[t/2] |= sym << 4` for odd `t` (shift = 4 * (t & 1)); at K=8 it is a // plain byte store (SYMS_PER_BYTE == 1, shift == 0). The caller must have // zeroed the row first — this ORs, it does not assign. +// +// The straddling branch writes bytes `b0 .. b0 + ceil((off+K)/8) - 1`, i.e. +// only the bytes the symbol's bits actually fall in — NOT a fixed +// `MAX_BYTES_PER_SYMBOL` run. That is what keeps the LAST symbol of a row in +// bounds: its highest bit is `num_symbols*K - 1`, which lives in byte +// `(num_symbols*K - 1) / 8 <= ceil(num_symbols*K/8) - 1`. +// +// Only thread 0 packs (the backtrace walk is serial), so the read-modify-write +// on a byte two adjacent symbols share is not a race. template __device__ __forceinline__ void qtip_pack_symbol( uint8_t* __restrict__ pkd, int t, unsigned int sym ) { - const int slot = t / (int)G::SYMS_PER_BYTE; - const unsigned shift = G::K * ((unsigned)t % G::SYMS_PER_BYTE); - pkd[slot] |= (uint8_t)((sym & G::SYM_MASK) << shift); + if constexpr (G::BYTE_ALIGNED) { + const int slot = t / (int)G::SYMS_PER_BYTE; + const unsigned shift = G::K * ((unsigned)t % G::SYMS_PER_BYTE); + pkd[slot] |= (uint8_t)((sym & G::SYM_MASK) << shift); + } else { + const int bit = t * (int)G::K; + const int b0 = bit >> 3; + const unsigned int off = (unsigned int)(bit & 7); + const unsigned int used = off + G::K; // bits of `b0`'s window touched + // `sym < 2^15` and `off <= 7`, so this cannot overflow 32 bits. + const unsigned int v = (sym & G::SYM_MASK) << off; + #pragma unroll + for (unsigned int b = 0; b < G::MAX_BYTES_PER_SYMBOL; ++b) { + if (8u * b < used) { + pkd[b0 + (int)b] |= (uint8_t)((v >> (8u * b)) & 0xFFu); + } + } + } } diff --git a/mistralrs-quant/src/qtip/cuda_ops.rs b/mistralrs-quant/src/qtip/cuda_ops.rs index 66048cdee..05f7b8fc8 100644 --- a/mistralrs-quant/src/qtip/cuda_ops.rs +++ b/mistralrs-quant/src/qtip/cuda_ops.rs @@ -923,7 +923,8 @@ fn ensure_non_beam_geometry(what: &str) -> Result<()> { } /// One-shot quantize entry point. Returns `(packed_blocks, row_scales)`: -/// * `packed_blocks` — `[n_rows, num_symbols / (8/K)]` U8, `8/K` symbols per byte. +/// * `packed_blocks` — `[n_rows, ceil(num_symbols*K/8)]` U8; symbol `t` occupies +/// bits `[t*K, t*K + K)`, LSB-first. /// * `row_scales` — `[n_rows]` F32, per-row scale. /// /// `weight_rotated_f32` should already be in the rotated frame (caller is @@ -947,15 +948,21 @@ pub(crate) fn quantize_rows_cuda( } let (n_rows, k_in) = weight_rotated_f32.dims2()?; let num_symbols = k_in / super::V as usize; - // Symbols per packed byte, from the rung's K: 2 at K=4, 1 at K=8. Derived - // rather than written as `2` so the packing follows `qtip::K` the day it - // moves — `qtip_geom.cuh::SYMS_PER_BYTE` is the CUDA side of this number. - let syms_per_byte = 8 / super::K as usize; - if !num_symbols.is_multiple_of(syms_per_byte) { + // The packed row is a bitstream: symbol `t` at bits `[t*K, t*K + K)`, + // LSB-first, `ceil(num_symbols*K/8)` bytes. `super::packed_bytes_per_row` + // is the single formula; `qtip_geom.cuh::qtip_packed_bytes_per_row` is its + // CUDA twin and `Rung::packed_bytes` is the serving side's. + // + // The previous `8 / super::K` form was a division by zero at K=9 — the + // rung the quality sweep selected — and would have panicked before the + // kernel was ever reached. + if !super::packed_row_is_whole_bytes(num_symbols, super::K) { candle_core::bail!( - "QTIP quantize CUDA: num_symbols ({num_symbols}) must be a multiple of \ - {syms_per_byte} for K={} packing", - super::K + "QTIP quantize CUDA: num_symbols ({num_symbols}) x K={} is not a whole number \ + of bytes ({} bits). Refusing rather than baking a row whose stride the decoder \ + would have to guess.", + super::K, + num_symbols * super::K as usize, ); } if weight_rotated_f32.dtype() != DType::F32 { @@ -986,7 +993,7 @@ pub(crate) fn quantize_rows_cuda( let row_scales = compute_row_scales_cuda(&weight_contig, codebook.scale_divisor())?; // Step 2: allocate packed output. - let packed_per_row = num_symbols / syms_per_byte; + let packed_per_row = super::packed_bytes_per_row(num_symbols, super::K); let packed_buf = dev.alloc_zeros::(n_rows * packed_per_row)?; let packed_shape = candle_core::Shape::from_dims(&[n_rows, packed_per_row]); diff --git a/mistralrs-quant/src/qtip/mod.rs b/mistralrs-quant/src/qtip/mod.rs index 2817ee533..0e4c757d6 100644 --- a/mistralrs-quant/src/qtip/mod.rs +++ b/mistralrs-quant/src/qtip/mod.rs @@ -385,6 +385,58 @@ const LUT_SIZE: usize = 1 << L; /// Symbol alphabet size: 2^K. const ALPHABET: usize = 1 << K; +/// Packed bytes holding `num_symbols` trellis symbols at `k` bits per symbol: +/// `ceil(num_symbols * k / 8)`. +/// +/// The Rust twin of `kernels/qtip/qtip_geom.cuh::qtip_packed_bytes_per_row`, +/// and the same rule as the serving side's `Rung::packed_bytes`. The encoder +/// and the decoder must agree on the row stride byte for byte, so there is one +/// formula and both sides use it. +/// +/// **Not `num_symbols / (8 / k)`.** That form is only correct when `k` divides +/// 8; at the K=9/V=4/L=12 rung `8 / 9 == 0` and it divides by zero. It agrees +/// with this one wherever it is defined: at K=4 with `num_symbols` even +/// (enforced by [`packed_row_is_whole_bytes`]) `(4n+7)/8 == n/2`, and at K=8 +/// `(8n+7)/8 == n`. +pub(crate) const fn packed_bytes_per_row(num_symbols: usize, k: u32) -> usize { + (num_symbols * k as usize).div_ceil(8) +} + +/// Does a row of `num_symbols` symbols at `k` bits occupy a whole number of +/// bytes? +/// +/// At K=4 this is exactly the shipped "num_symbols must be even" rule. It is +/// kept at every rung because a partial trailing byte is a row stride the two +/// sides can disagree about for free, and every real V4-Flash / Qwen expert +/// shape satisfies it (`in_features` 7168 and 2048 give `num_symbols` 1792 and +/// 512 at V=4, both multiples of 8). Refuse, never round. +pub(crate) const fn packed_row_is_whole_bytes(num_symbols: usize, k: u32) -> bool { + (num_symbols * k as usize).is_multiple_of(8) +} + +// The general formula must reproduce the shipped K=4 nibble packing and the +// K=8 byte packing EXACTLY, at the real V4-Flash expert shapes — otherwise +// re-pointing the bake at another rung silently resizes the artifact of the +// rung that already shipped. Asserted at compile time so it holds in every +// build, `cuda` feature or not, and so these `const fn`s have a use that does +// not depend on which features are enabled. +// +// in_features 7168 -> num_symbols 3584 (V=2) / 1792 (V=4) +// in_features 2048 -> num_symbols 1024 (V=2) / 512 (V=4) +const _: () = { + // K=4/V=2/L=16, the shipped rung: identical to `num_symbols / 2`. + assert!(packed_bytes_per_row(3584, 4) == 3584 / 2); + assert!(packed_bytes_per_row(1024, 4) == 1024 / 2); + // K=8/V=4/L=12: identical to `num_symbols`. + assert!(packed_bytes_per_row(1792, 8) == 1792); + assert!(packed_bytes_per_row(512, 8) == 512); + // K=9/V=4/L=12: 9/8 of the K=8 row, exactly, at both shapes. + assert!(packed_bytes_per_row(1792, 9) == 1792 * 9 / 8); + assert!(packed_bytes_per_row(512, 9) == 512 * 9 / 8); + assert!(packed_row_is_whole_bytes(1792, 9)); + assert!(packed_row_is_whole_bytes(512, 9)); +}; + pub(crate) const N_BITS: usize = 2; /// Per-state pseudo-random Gaussian LUT. @@ -4261,6 +4313,150 @@ impl QuantizedSerde for QtipLayer { mod tests { use super::*; + /// The bitstream `kernels/qtip/qtip_geom.cuh::qtip_pack_symbol` writes in + /// its straddling branch, implemented independently here. + /// + /// Symbol `t` at bits `[t*K, t*K + K)`, LSB-first — the same rule as the + /// serving side's `Rung::pack`. This is a REFERENCE, not a wrapper: it + /// exists so the K=4 identity below compares two separately-written + /// packers rather than one packer against itself. + fn bitstream_pack(syms: &[u32], k: u32) -> Vec { + let mut out = vec![0u8; packed_bytes_per_row(syms.len(), k)]; + let mask = (1u64 << k) - 1; + for (t, &s) in syms.iter().enumerate() { + let bit = t * k as usize; + let v = (u64::from(s) & mask) << (bit % 8); + let need = ((bit % 8) + k as usize).div_ceil(8); + for i in 0..need { + out[bit / 8 + i] |= ((v >> (8 * i)) & 0xFF) as u8; + } + } + out + } + + /// The shipped K=4 nibble packer, verbatim from `cpu_reference_packed`. + fn nibble_pack_k4(syms: &[u32]) -> Vec { + let mut packed = vec![0u8; syms.len() / 2]; + for (i, &sym) in syms.iter().enumerate() { + if i.is_multiple_of(2) { + packed[i / 2] = (sym & 0x0F) as u8; + } else { + packed[i / 2] |= ((sym & 0x0F) << 4) as u8; + } + } + packed + } + + /// **The no-GPU half of "K=4/V=2/L=16 stays byte-identical".** + /// + /// `qtip_pack_symbol` stopped being "one byte per `8/K` symbols" and became + /// a bitstream so that K=9 — where a symbol straddles two bytes — can be + /// baked at all. That generalisation is only safe if the bitstream rule and + /// the nibble rule are the SAME bytes at K=4, which is a property of the + /// two rules and can be checked on the host with no CUDA toolchain and no + /// device. + /// + /// (The CUDA side does not even rely on this: the byte-aligned branch is + /// kept verbatim under `if constexpr`, so K=4 compiles from unchanged + /// source. This pins the rule the discarded branch would have produced, so + /// that collapsing the two later is a checked change rather than a hope.) + #[test] + fn the_bitstream_reproduces_the_shipped_k4_nibble_packing() { + let mut state = 0x1234_5678u32; + let syms: Vec = (0..4096) + .map(|_| { + state = state.wrapping_mul(1_664_525).wrapping_add(1_013_904_223); + (state >> 17) & 0x0F + }) + .collect(); + assert_eq!(bitstream_pack(&syms, 4), nibble_pack_k4(&syms)); + } + + /// K=8 is a plain byte store, and the bitstream must agree with that too. + #[test] + fn the_bitstream_is_a_byte_store_at_k8() { + let syms: Vec = (0..1792u32).map(|t| (t * 37) & 0xFF).collect(); + let packed = bitstream_pack(&syms, 8); + assert_eq!(packed.len(), 1792); + assert!(packed + .iter() + .zip(syms.iter()) + .all(|(&b, &s)| u32::from(b) == s)); + } + + /// K=9 round-trips through the bit rule the SERVING side extracts with. + /// + /// The encoder and `Rung::extract` must read the same bits at the same + /// offsets, so the extraction is written out here independently rather + /// than as `bitstream_pack`'s inverse-by-construction. Includes the last + /// symbol of the row, which is the one a wrong `packed_bytes` would run + /// off the end of. + #[test] + fn k9_symbols_land_where_the_decoder_looks_for_them() { + const K: u32 = 9; + let n = 1792usize; // V4-Flash gate/up: in_features 7168 at V=4. + let mut state = 0x9E37_79B9u32; + let syms: Vec = (0..n) + .map(|_| { + state = state.wrapping_mul(1_664_525).wrapping_add(1_013_904_223); + (state >> 13) & 0x1FF + }) + .collect(); + let packed = bitstream_pack(&syms, K); + + assert_eq!(packed.len(), packed_bytes_per_row(n, K)); + assert_eq!(packed.len(), n * 9 / 8, "9 bits/symbol, whole bytes"); + + for (t, &want) in syms.iter().enumerate() { + let bit = t * K as usize; + let (b0, off) = (bit / 8, (bit % 8) as u32); + let need = (off as usize + K as usize).div_ceil(8); + assert!( + b0 + need <= packed.len(), + "symbol {t} extraction runs past the packed row" + ); + let mut w = 0u32; + for i in 0..need { + w |= u32::from(packed[b0 + i]) << (8 * i); + } + assert_eq!((w >> off) & 0x1FF, want, "symbol {t}"); + } + } + + /// The straddling bound is `ceil((8 - gcd(K,8) + K)/8)`, not + /// `ceil((7+K)/8)`: the reachable bit offsets are multiples of `gcd(K,8)`, + /// so 7 is not reachable at K=8 or K=10. Mirrors + /// `QtipGeom::MAX_BYTES_PER_SYMBOL`; an over-estimate would have the packer + /// touch a byte past the row. + #[test] + fn the_straddle_bound_is_exact_at_every_baked_k() { + for (k, want) in [(4u32, 1usize), (8, 1), (9, 2), (10, 2)] { + let reached: usize = (0..8usize) + .map(|t| ((t * k as usize) % 8 + k as usize).div_ceil(8)) + .max() + .unwrap(); + let g: usize = [8usize, 4, 2, 1] + .into_iter() + .find(|g| (k as usize).is_multiple_of(*g)) + .unwrap(); + let formula = (8 - g + k as usize).div_ceil(8); + assert_eq!(reached, want, "K={k}: reachable worst case"); + assert_eq!(formula, want, "K={k}: gcd formula"); + } + } + + /// A partial trailing byte is refused, not rounded — at every rung. + #[test] + fn a_row_that_is_not_whole_bytes_is_rejected() { + assert!(packed_row_is_whole_bytes(1792, 9)); + assert!(packed_row_is_whole_bytes(512, 9)); + assert!(!packed_row_is_whole_bytes(1793, 9)); + // K=4's rule was "num_symbols must be even"; this must still be it. + for n in 0..64usize { + assert_eq!(packed_row_is_whole_bytes(n, 4), n.is_multiple_of(2)); + } + } + /// Sanity: LUT has the right size and reasonable distribution. #[test] fn gaussian_lut_is_well_formed() { From 4e775a2c1a7702760c078ea9a2f9f186a6ae13d3 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 15:42:49 +0100 Subject: [PATCH 2/7] test(qtip/bake): compile qtip_geom.cuh on the HOST and check the packed bytes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit "K=4/V=2/L=16 stays byte-identical" was, until this, an argument. The instruments that could settle it are `cuda_beam_matches_cpu_beam_bit_for_bit` and `cuda_beam_unpruned_matches_cuda_exhaustive`, and both are `#[cfg(feature = "cuda")]` and SKIP without a device — so on the Mac dev box and on every GPU-less runner, nothing checked the packer at all. The packer is integer arithmetic on a byte array, though: stub `__device__` / `__forceinline__` and the exact source text nvcc compiles also compiles with clang++/g++. `arc-tools/qtip_pack_host_check.sh` #includes the real header and asserts, at the real V4-Flash expert shapes (in_features 7168 and 2048): * K=4 emits bytes IDENTICAL to the shipped nibble packer over 4096 symbols — compared against a transcription of `cpu_reference_packed`, not against the new wire rule, so it is the shipping guarantee that is being checked. * K=4/K=8/K=9 all match an independently written bitstream reference. * `packed_bytes_per_row` equals `ceil(num_symbols*K/8)` at every rung. Two things the obvious version of this script would have got wrong, both found by running the mutations rather than reasoning about them: 1. A GUARD BAND CANNOT SEE THIS OVERRUN. Dropping the packer's `if (8u*b < used)` writes `|= 0` past the row — the byte value does not change, so zeros or any sentinel look untouched. The row is therefore an EXACT-size heap allocation and the build asks for `-fsanitize=address`, which reports it as a heap-buffer-overflow. The script says which mode it ran in rather than implying coverage it lacks. The overrun is not benign in the kernel: rows are contiguous and different blocks own different rows concurrently, so a non-atomic RMW on the next row's first byte can clobber that row's own store. 2. `QtipGeom<6,3,12>` is instantiated purely to make that guard non-vacuous. At K=9 and K=10 every reachable offset needs exactly 2 bytes, so the guard never fires and would be untested. K=6 has `off==0 => used==6`, one byte against a bound of two. Proven red: `MAX_BYTES_PER_SYMBOL -> 1` (K=9 and K=6 truncate); nibble order flipped (K=4 identity fails); `packed_bytes_per_row -> num_symbols` (K=9 length fails); guard dropped (ASan heap-buffer-overflow). Also recorded, because a gate's own docs must not overclaim: enlarging `MAX_BYTES_PER_SYMBOL` to the naive `ceil((7+K)/8)` does NOT go red here and should not — the per-symbol `used` guard makes an over-large bound a dead unrolled iteration. The bound's exactness matters where it is a READ width, on the serving side, and is pinned by `the_straddle_bound_is_exact_at_every_baked_k`. Wired into the CUDA lane ahead of the spill gate: it is the cheaper failure and needs no CUDA toolkit. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- .github/workflows/cuda_compile_check.yaml | 14 +- arc-tools/qtip_pack_host_check.sh | 228 ++++++++++++++++++++++ mistralrs-quant/kernels/qtip/qtip_beam.cu | 43 ++++ 3 files changed, 284 insertions(+), 1 deletion(-) create mode 100755 arc-tools/qtip_pack_host_check.sh diff --git a/.github/workflows/cuda_compile_check.yaml b/.github/workflows/cuda_compile_check.yaml index e51950b7d..82287a608 100644 --- a/.github/workflows/cuda_compile_check.yaml +++ b/.github/workflows/cuda_compile_check.yaml @@ -269,8 +269,20 @@ jobs: # silent — `__launch_bounds__` spills rather than failing — and a spilled # `cand[]` is exactly the mis-scoping behind the discredited # "~1,700 s/layer" K=8/V=4/L=12 encode estimate. The script proves itself - # red on a negative control and asserts both geometries are present, so it + # red on a negative control and asserts every geometry is present, so it # cannot pass vacuously. nvcc + cuobjdump only; no GPU. + # + # The packer gate below runs FIRST because it is the cheaper failure: it + # compiles `qtip_geom.cuh` with the HOST compiler (the CUDA attributes + # stubbed) and checks the bytes the packed bitstream actually contains at + # K=4, K=8 and K=9. The K=9 rung straddles bytes, so "K=4/V=2/L=16 stays + # byte-identical" stopped being self-evident; the CUDA parity tests that + # would prove it (`cuda_beam_*_matches_cpu_*`) SKIP without a device, so + # on every GPU-less runner nothing checked it at all. No CUDA toolkit + # needed — it would run on the Mac dev box too. + - name: QTIP packed-bitstream gate (host compiler, no CUDA) + run: bash arc-tools/qtip_pack_host_check.sh + - name: QTIP beam spill gate (LOCAL:0, all geometries) run: bash arc-tools/qtip_beam_res_usage_check.sh sm_${{ matrix.compute_cap }} diff --git a/arc-tools/qtip_pack_host_check.sh b/arc-tools/qtip_pack_host_check.sh new file mode 100755 index 000000000..cb3ebf59e --- /dev/null +++ b/arc-tools/qtip_pack_host_check.sh @@ -0,0 +1,228 @@ +#!/usr/bin/env bash +# Compile the REAL `kernels/qtip/qtip_geom.cuh` with a host C++ compiler and +# assert the packed bitstream it emits, at every baked geometry. +# +# WHY THIS EXISTS +# --------------- +# `qtip_pack_symbol` stopped being "one byte per 8/K symbols" and became a +# bitstream, because K=9/V=4/L=12 — the rung the quality sweep selected, at +# 2.25 bpw — puts a symbol across a byte boundary. Two things then need +# proving, and the obvious instruments cannot prove either: +# +# * K=4/V=2/L=16 must stay BYTE-IDENTICAL. It is the rung that ships; a +# regression here breaks the published artifact, not a future one. +# * K=9 must land its bits exactly where the DECODER looks for them +# (`qtip/trellis_v4l12.rs::Rung::extract`), and must not write past +# `ceil(num_symbols*K/8)`. +# +# The CUDA parity tests (`cuda_beam_*_matches_cpu_*`) are the hardware gate for +# both, and they SKIP without a device — so on a Mac dev box, and on every +# GPU-less CI runner, nothing checked this at all. The packer is pure integer +# arithmetic on a byte array, though: stub `__device__`/`__forceinline__` and +# the exact source text nvcc compiles also compiles with clang++/g++. That is +# what this does. No CUDA toolkit, no GPU, ~1 second. +# +# It checks the SOURCE, not a transcription of it: the header is #included, so +# a change to `qtip_pack_symbol` or `MAX_BYTES_PER_SYMBOL` is seen here even +# though the reference packer below is written independently. +# +# THIS GATE IS PROVEN RED +# ----------------------- +# Mutations run against it while it was written, each red and only in the row +# it should be: +# * `MAX_BYTES_PER_SYMBOL` -> 1 : K=9 and K=6 rows truncate +# * byte-aligned shift `4*(t&1)` -> `4*((t+1)&1)` : K=4 nibble identity fails +# * `packed_bytes_per_row` -> `num_symbols` : K=9 length assertion fails +# * drop `if (8u*b < used)` from the packer : ASan heap-buffer-overflow +# +# AND ONE MUTATION THAT DOES **NOT** GO RED, STATED SO NOBODY BUILDS ON IT. +# Enlarging `MAX_BYTES_PER_SYMBOL` (e.g. to the naive `ceil((7+K)/8)`, which is +# 3 at K=9) is NOT caught here, and should not be: the packer's per-symbol +# `used = off + K` guard already stops the extra unrolled iteration from +# storing, so an over-large bound costs a dead instruction and nothing else. +# The bound's EXACTNESS matters on the serving side, where the same formula is +# a read width — that is pinned by `qtip/mod.rs`'s +# `the_straddle_bound_is_exact_at_every_baked_k`, not by this script. +# +# `QtipGeom<6,3,12>` is instantiated below purely to make that `used` guard +# non-vacuous. It is NOT a rung anyone bakes: at K=9 and K=10 every reachable +# offset gives `ceil((off+K)/8) == 2 == MAX_BYTES_PER_SYMBOL`, so the guard +# never fires and would be untested code. K=6 has `off == 0 => used == 6`, one +# byte, against a bound of 2 — the case that runs off the end of the row when +# the guard is removed. +# +# ⚠ Removing that guard writes `|= 0`, so the byte VALUE is unchanged and a +# guard band of zeros — or of any sentinel — cannot detect it. Only ASan can, +# which is why the row is an exact-size heap allocation and the build asks for +# `-fsanitize=address`. The overrun still matters in the kernel: rows are +# contiguous in one buffer and different blocks own different rows at the same +# time, so a non-atomic read-modify-write on the next row's first byte can +# clobber that row's own store. +# +# Usage: arc-tools/qtip_pack_host_check.sh +set -euo pipefail + +REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +KDIR="${REPO_ROOT}/mistralrs-quant/kernels/qtip" +OUT="$(mktemp -d)" +trap 'rm -rf "${OUT}"' EXIT + +CXX="${CXX:-}" +if [ -z "${CXX}" ]; then + for c in clang++ g++ c++; do + if command -v "${c}" >/dev/null 2>&1; then CXX="${c}"; break; fi + done +fi +if [ -z "${CXX}" ]; then + echo "FAIL: no host C++ compiler found (tried clang++, g++, c++)." + echo " This gate needs no CUDA toolkit — only a C++17 compiler." + exit 1 +fi + +cat >"${OUT}/qtip_pack_host_check.cpp" <<'EOF' +// The CUDA attributes are the only thing that stops a host compiler from +// reading this header. Stub them and the source text under test is byte for +// byte what nvcc sees. +#define __device__ +#define __forceinline__ inline +#define __restrict__ +#include "qtip_geom.cuh" + +#include +#include +#include +#include + +// INDEPENDENT reference, written from the wire format rather than from the +// header: symbol `t` at bits [t*K, t*K + K), LSB-first. This is the rule +// `qtip/trellis_v4l12.rs::Rung::pack` implements and `Rung::extract` inverts. +static std::vector ref_pack(const std::vector& syms, uint32_t k) { + const size_t bytes = (syms.size() * (size_t)k + 7) / 8; + std::vector out(bytes, 0); + const uint64_t mask = (1ull << k) - 1; + for (size_t t = 0; t < syms.size(); ++t) { + const size_t bit = t * (size_t)k; + const uint64_t v = ((uint64_t)syms[t] & mask) << (bit % 8); + const size_t need = ((bit % 8) + k + 7) / 8; + for (size_t i = 0; i < need; ++i) { + out[bit / 8 + i] |= (uint8_t)((v >> (8 * i)) & 0xFF); + } + } + return out; +} + +// The SHIPPED K=4 nibble packer, transcribed from `cpu_reference_packed` in +// `qtip/mod.rs`. "Byte-identical" is a claim about this, not about ref_pack. +static std::vector nibble_pack_k4(const std::vector& syms) { + std::vector p(syms.size() / 2, 0); + for (size_t i = 0; i < syms.size(); ++i) { + if (i % 2 == 0) p[i / 2] = (uint8_t)(syms[i] & 0xF); + else p[i / 2] |= (uint8_t)((syms[i] & 0xF) << 4); + } + return p; +} + +static uint32_t rng_state = 0x9E3779B9u; +static uint32_t nxt() { rng_state = rng_state * 1664525u + 1013904223u; return rng_state; } + +template +static int check(const char* name, int n) { + std::vector syms((size_t)n); + for (int i = 0; i < n; ++i) syms[(size_t)i] = nxt() & G::SYM_MASK; + + const int bytes = qtip_packed_bytes_per_row(n); + const auto want = ref_pack(syms, G::K); + if ((size_t)bytes != want.size()) { + printf("FAIL %s: packed_bytes_per_row = %d, wire format needs %zu\n", + name, bytes, want.size()); + return 1; + } + + // EXACT-SIZE heap row, on purpose. A packer that runs one byte past the + // end writes `|= 0` there — the value is unchanged, so a guard band of + // zeros (or of any sentinel) cannot see it. ASan's redzone can, which is + // why this allocates exactly `bytes` and the script builds with + // `-fsanitize=address` whenever the compiler has it. + // + // That overrun is NOT harmless in the kernel even though the byte value + // does not change: rows are contiguous in one buffer and different blocks + // own different rows concurrently, so a non-atomic read-modify-write on + // the next row's first byte can clobber that row's own store. + uint8_t* got = new uint8_t[(size_t)bytes](); + for (int t = 0; t < n; ++t) qtip_pack_symbol(got, t, syms[(size_t)t]); + + int rc = 0; + if (memcmp(got, want.data(), (size_t)bytes) != 0) { + for (int i = 0; i < bytes; ++i) { + if (got[(size_t)i] != want[(size_t)i]) { + printf("FAIL %s: byte %d is 0x%02X, wire format says 0x%02X\n", + name, i, got[(size_t)i], want[(size_t)i]); + break; + } + } + rc = 1; + } else { + printf("ok %-16s K=%u V=%u L=%2u bytes=%-5d byte_aligned=%d max_bytes/sym=%u\n", + name, G::K, G::V, G::L, bytes, (int)G::BYTE_ALIGNED, G::MAX_BYTES_PER_SYMBOL); + } + delete[] got; + return rc; +} + +int main() { + int bad = 0; + // The real V4-Flash expert shapes: in_features 7168 (gate/up) and 2048 + // (down), so num_symbols is 3584/1024 at V=2 and 1792/512 at V=4. + bad |= check("K4V2L16 n=3584", 3584); + bad |= check("K4V2L16 n=1024", 1024); + bad |= check("K8V4L12 n=1792", 1792); + bad |= check("K8V4L12 n=512", 512); + bad |= check("K9V4L12 n=1792", 1792); + bad |= check("K9V4L12 n=512", 512); + + // NOT a rung. See the header: K=6 is the geometry that makes the packer's + // `used = off + K` guard fire (off == 0 gives a 6-bit symbol inside ONE + // byte, against MAX_BYTES_PER_SYMBOL == 2). Without that guard the last + // symbol of the row writes into the guard band. `n` is a multiple of 4 so + // the row is whole bytes. + bad |= check>("K6V3L12 n=1024", 1024); + + // THE shipping guarantee: K=4 is byte-identical to the packer that baked + // the published artifact, not merely consistent with the new wire rule. + std::vector s(4096); + for (size_t i = 0; i < s.size(); ++i) s[i] = nxt() & 0xF; + std::vector got(s.size() / 2, 0); + for (size_t t = 0; t < s.size(); ++t) { + qtip_pack_symbol(got.data(), (int)t, s[t]); + } + if (got != nibble_pack_k4(s)) { + printf("FAIL K=4 is NOT byte-identical to the shipped nibble packer\n"); + bad = 1; + } else { + printf("ok K=4 is byte-identical to the shipped nibble packer (4096 symbols)\n"); + } + + printf(bad ? "RESULT: FAILED\n" : "RESULT: PASS\n"); + return bad; +} +EOF + +echo "== QTIP packed-bitstream host gate (no CUDA toolkit, no GPU) ==" +echo "compiler: $(${CXX} --version | head -1)" + +# AddressSanitizer is what detects an overrun whose written VALUE is unchanged +# (`|= 0` past the end of the row). Without it this script still checks every +# byte of every row, but it cannot see that class at all — so say which mode +# it ran in rather than implying a coverage it does not have. +SAN=(-fsanitize=address -fno-omit-frame-pointer) +if ! "${CXX}" -std=c++17 "${SAN[@]}" -x c++ -c -o /dev/null - /dev/null; then + SAN=() + echo "asan: UNAVAILABLE — byte contents are checked, overruns are NOT" +else + echo "asan: on (row allocated at exactly packed_bytes_per_row)" +fi +echo +"${CXX}" -std=c++17 -O1 -g -Wall -Wextra -Werror "${SAN[@]}" \ + -I "${KDIR}" -o "${OUT}/qtip_pack_host_check" \ + "${OUT}/qtip_pack_host_check.cpp" +"${OUT}/qtip_pack_host_check" diff --git a/mistralrs-quant/kernels/qtip/qtip_beam.cu b/mistralrs-quant/kernels/qtip/qtip_beam.cu index 11a455cb9..d4738bb3e 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -103,6 +103,49 @@ // `beam_w` entries, which fit in shared memory with room to spare. HBM traffic // per symbol position drops ~512x. // +// WHAT A RUNG COSTS TO ENCODE — AND WHY `(n/V) * W * 2^K` IS NOT IT +// ----------------------------------------------------------------- +// That formula predicts K=4 -> K=8 costs 8x MORE and produced the discredited +// "~1,700 s/layer / 20 h / $30" figure. This kernel falsifies it in two lines +// of its own source: the `& G::GROUP_MASK` in step 1a masks the predecessor to +// `2^(L-K)`, so raising K SHRINKS the live group count by exactly the factor by +// which it grows the alphabet, and `n_cand = ng * 2^K` at the end of step 2 is +// therefore saturated at `2^L`. The candidate +// count per timestep is 4096 at K=4/L=16 (256 groups x 16), at K=8/L=12 +// (16 x 256) and at K=9/L=12 (8 x 512) alike, and `cand[]` is 16 entries at all +// three. `W` is not a factor either — measured, W=256 = 82.7 s/layer vs +// W=32 = 83.6 s/layer, ~1% (FACTS.md:1320). +// +// What DOES move, per layer: +// +// (a) timesteps. `num_symbols = in_features / V`, so V=2 -> V=4 halves them. +// (b) candidates/timestep, cand[], radix passes, LUT bytes per ROW. All +// invariant (see above; LUT is 2^K*V*4 B per group over 2^(L-K) groups, +// which is 65,536 B/timestep at both V=4 rungs against half as many +// timesteps as V=2's 32,768 B). +// (c) branch metric. V=2 -> V=4 doubles the arithmetic per candidate, but a +// candidate's cost is dominated by its ~3.87 radix histogram passes and +// its compaction test, not by the metric — so this is a fraction of 2x. +// (d) K=8 -> K=9 at fixed V=4/L=12. (a), (b) and (c) are ALL unchanged. Only +// the packer moves: one byte store per symbol becomes two byte +// read-modify-writes, in thread 0's serial backtrace, against an +// L1-resident row of <= 2 KiB. +// +// ⇒ From the measured anchor `qtip2` beam W=256 = 372.0 s/layer on an A100 +// (FACTS.md:990; harness validated to 0.3% against the published V4 bake's +// 370-376 s/layer, FACTS.md:993): +// +// K=9/V=4/L=12 ~= 372.0 x 0.5 x [1.0 .. 1.4] = 186-260 s/layer +// +// i.e. the 2.25 bpw rung is CHEAPER to bake than the 2.00 bpw rung that +// ships, not 8x dearer. 43 layers ⇒ 2.2-3.1 h ⇒ $3.3-$4.6 at $1.49/hr, or +// 1.1-1.6 h across two devices at the same total cost (`ARC_BAKE_DEVICES`, +// measured 1.97x on 2 GPUs, byte-identical — wave29-BD-rung-decision.md:431). +// +// ⚠ PROJECTED, NOT MEASURED. Nothing has run a V=4 bake. This is a division +// from a measured V=2 anchor by terms this file's own structure fixes; the +// only honest way to close it is to bake a layer and difference the markers. +// // SHARED-MEMORY RESIDENCY (the claim wave13-AF was asked to test) // -------------------------------------------------------------- // At beam_w = 256 the *beam* is 2 KiB (cost f32 + state u16 + parent u16). From 8d4635afb28d814ae81d5735c27864a2489cd039 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 15:46:29 +0100 Subject: [PATCH 3/7] docs(qtip): index the packer gate, and repair the citation my edits moved MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three doc-level consequences of the K=9 change, none of them cosmetic: * `QUANTIZATION_PERFORMANCE.md:34` cites `qtip_beam.cu:67-79` for the "3.2e9 symbol positions per V4-Flash layer at V=2" figure. Adding the encode-cost section to that header pushed the cited block to `:83-95`. A stale `path:line` in a doc whose whole premise is "nothing is stated without a label" is worse than no label. * `CUDA_VALIDATION.md` gains gate 1d. Its table is the answer to "what is checked for free"; a gate that is not in it does not exist to the next agent. * The quality figures behind the rung choice (-0.00698 / -0.00307 / +0.00402) are now stated ONCE, in `qtip/trellis_v4l12.rs`'s module docs, and pointed at from the kernel headers instead of restated there. Two copies of a measured number drift. The kernel headers also stop citing `qtip_beam.cu` line numbers for the two lines that falsify `(n/V)*W*2^K` — they name the constructs (`& G::GROUP_MASK` in step 1a, `n_cand = ng * 2^K` at the end of step 2) instead, since those line numbers move every time the header above them grows. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- arc-tools/CUDA_VALIDATION.md | 41 ++++++++++++++++++++ docs/engineering/QUANTIZATION_PERFORMANCE.md | 2 +- mistralrs-quant/kernels/qtip/qtip_beam.cu | 9 ++--- mistralrs-quant/kernels/qtip/qtip_geom.cuh | 11 ++++-- 4 files changed, 54 insertions(+), 9 deletions(-) diff --git a/arc-tools/CUDA_VALIDATION.md b/arc-tools/CUDA_VALIDATION.md index efd5d02bc..c4b598cd9 100644 --- a/arc-tools/CUDA_VALIDATION.md +++ b/arc-tools/CUDA_VALIDATION.md @@ -20,6 +20,7 @@ sm_90 = H100/Hopper). | 1. `cuda_compile_check.yaml` | GitHub Actions | No | Free, automatic | nvcc errors, FFI drift, `cuda`-feature Rust compile errors in the **library** crates — for sm_80 **and** sm_90. Does **not** compile `flash-attn` or link the CLI binaries (see gate 1b). | | 1b. `flash_attn_compile_check.yaml` | GitHub Actions | No | Free, manual | The `flash-attn` feature (candle-flash-attn nvcc compile) the rental uses but gate 1 omits — for the chosen arch (default sm_90). See "flash-attn coverage" below. | | 1c. `qtip_beam_res_usage_check.sh` | Step *inside* gate 1 | No | Free, automatic | **Local-memory spills in the QTIP beam kernels** — `cuobjdump -res-usage` must report `LOCAL:0 STACK:0`. Both beam kernels demand this check by name in their own source and nothing ran it before. See "the spill gate" below. | +| 1d. `qtip_pack_host_check.sh` | Step *inside* gate 1 | No | Free, automatic — **and needs no CUDA toolkit at all** | **The bytes the QTIP packer emits**, at K=4, K=8 and K=9. Compiles the real `qtip_geom.cuh` with the host compiler (CUDA attributes stubbed) and checks K=4 is byte-identical to the shipped nibble packer. See "the packer gate" below. | | 2. `colab_cuda_build_check.ipynb` | Google Colab | nvcc only | Free, manual | Same as gate 1, plus it *links the CLI binaries* (`cuda_compile_check.sh` step 4, `FEATURES=cuda`); also runtime tests **iff** Colab gives sm_80+ | | 3. `cuda_compile_check.sh` (GPU mode) + rental step 4b | Rental / sm_80+ box | Yes | Paid box | Kernel **runtime**: parity, the prefix-grouped Viterbi quantize kernel actually running, no hang | @@ -63,6 +64,46 @@ undercounts SASS by ~2.05x on this kernel family; a spilled `cand[]` is exactly what produced the retracted "~1,700 s/layer" K=8/V=4/L=12 encode estimate. A number you did not compile is an estimate. +### The packer gate — the only free check of what a bake WRITES (gate 1d) +Gates 1 and 1c prove the QTIP kernels compile and do not spill. Neither looks +at a single byte of output. That was tolerable while the packed row was "8/K +whole symbols per byte" at every rung; it stopped being tolerable at +**K=9/V=4/L=12**, where a symbol straddles a byte boundary and the row became a +bitstream — so "K=4/V=2/L=16 is still byte-identical" turned into a claim about +a rewritten packer, guarding the artifact that actually ships. + +The instruments that could settle it, `cuda_beam_matches_cpu_beam_bit_for_bit` +and `cuda_beam_unpruned_matches_cuda_exhaustive`, are `#[cfg(feature = "cuda")]` +**and return early when `Device::new_cuda(0)` fails** — so on the Mac dev box +and on every GPU-less runner they check nothing. They remain gate 3. + +The packer is integer arithmetic on a byte array, though, so: + +```bash +arc-tools/qtip_pack_host_check.sh # no nvcc, no GPU, ~1 s +``` + +stubs `__device__`/`__forceinline__`, `#include`s the **real** +`kernels/qtip/qtip_geom.cuh`, and asserts at the real V4-Flash expert shapes +that K=4 reproduces the shipped nibble packer byte for byte, and that K=4/8/9 +all match an independently written bitstream reference. + +Two things it does that a first attempt would not, and both were found by +running the mutations rather than reasoning about them: +* the row is an **exact-size heap allocation built with `-fsanitize=address`**, + because a packer that runs one byte past the end writes `|= 0` — the value is + unchanged, so no guard band of zeros or of any sentinel can see it. In the + kernel that overrun is a non-atomic read-modify-write on the *next row's* + first byte, concurrent with the block that owns that row; +* it instantiates a geometry nobody bakes (`QtipGeom<6,3,12>`) purely to make + the packer's `used = off + K` guard fire at all — at K=9 and K=10 every + reachable offset needs exactly two bytes, so the guard is otherwise untested + code. + +It also records, in its own header, one mutation that does **not** go red +(enlarging `MAX_BYTES_PER_SYMBOL`), because a gate that overstates its coverage +is how a green check stops meaning anything. + ### flash-attn coverage — now free, one-click (gate 1b) The rental's step-4 build is `cargo build -p arc-cli -p mistralrs-cli --features "cuda flash-attn"`. Gates 1 and 2 above compile the **`cuda` feature only** — diff --git a/docs/engineering/QUANTIZATION_PERFORMANCE.md b/docs/engineering/QUANTIZATION_PERFORMANCE.md index 5df60b393..877f86593 100644 --- a/docs/engineering/QUANTIZATION_PERFORMANCE.md +++ b/docs/engineering/QUANTIZATION_PERFORMANCE.md @@ -31,7 +31,7 @@ Every number below carries one of these labels. Nothing is stated without one. width-256 beam over the same trellis. - Reference model: DeepSeek-V4-Flash, 284 B total / 13 B active, 43 quantized MoE layers, **6.4 B parameters per layer** ⇒ 3.2e9 symbol positions per layer - at `V = 2`. [source-verified, `mistralrs-quant/kernels/qtip/qtip_beam.cu:67-79`] + at `V = 2`. [source-verified, `mistralrs-quant/kernels/qtip/qtip_beam.cu:83-95`] - Artifact: **74.18 GB** UQFF in **8 shards + residual** (15 files) ⇒ **2.09 bits/param** over 284 B params. [measured, H200, reproduced 4×; size HF-API-verified on the published repo. Earlier revisions said "~68 GB / diff --git a/mistralrs-quant/kernels/qtip/qtip_beam.cu b/mistralrs-quant/kernels/qtip/qtip_beam.cu index d4738bb3e..43ac68719 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -12,11 +12,10 @@ // // K=4 / V=2 / L=16 2.00 bpw, the shipped rung; 15.125 decode inst/weight // K=8 / V=4 / L=12 2.00 bpw at 4.375 decode inst/weight -// K=9 / V=4 / L=12 2.25 bpw — THE BAKE TARGET. K=8/V=4/L=12 is quality- -// CLOSED (-0.00698 with a random codebook, -0.00307 with -// a converged trellis-Lloyd one, against a +-0.0008 ship -// band, over six codebook designs); K=9 measures -// +0.00402, 5x better than the shipped control. +// K=9 / V=4 / L=12 2.25 bpw — THE BAKE TARGET. K=8/V=4/L=12 is +// quality-CLOSED and K=9 is not; the sweep, its six +// codebook designs and the Dw_cos figures are recorded +// once, in `qtip/trellis_v4l12.rs`'s module docs. // // The K=8 rung is NOT reachable by bumping the constants. Three things in this // kernel are shaped by `2^(L-K)` and `2^K`, and at K=8/L=12 all three invert: diff --git a/mistralrs-quant/kernels/qtip/qtip_geom.cuh b/mistralrs-quant/kernels/qtip/qtip_geom.cuh index 555c92b31..d90b27e66 100644 --- a/mistralrs-quant/kernels/qtip/qtip_geom.cuh +++ b/mistralrs-quant/kernels/qtip/qtip_geom.cuh @@ -19,9 +19,14 @@ // // K=9 / V=4 / L=12 2.25 bits/weight, same 32,768 B codebook, same decode // shape. This is the rung the encoder is being taught to -// bake; K=8/V=4/L=12 is quality-CLOSED (measured -// -0.00698 random / -0.00307 trellis-Lloyd against a -// +-0.0008 ship band). +// bake; K=8/V=4/L=12 is quality-CLOSED (Dw_cos -0.00698 +// with a random codebook and -0.00307 with a converged +// trellis-Lloyd one, against a +-0.0008 ship band, over +// six codebook designs — trellis freedom binds, not +// codebook coverage). K=9 measures +0.00402. Those +// numbers and the sweep behind them live in +// `qtip/trellis_v4l12.rs`'s module docs; do not restate +// them anywhere else. // // Before this header each kernel restated `K`/`L`/`V` as its own local // `constexpr`, and every one of those restatements assumed K=4 somewhere else From fcf5d9e8fe58e22f5f5dec011ae7cd3b8b00dfc7 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 15:47:35 +0100 Subject: [PATCH 4/7] fix(arc-tools): the packer gate aborted on any runner without libasan MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `set -u` plus an EMPTY bash array is an unbound-variable error in bash <= 4.3 (macOS ships 3.2), so `"${SAN[@]}"` killed the script the moment the `-fsanitize=address` probe failed — i.e. on exactly the GPU-less runners the gate exists to serve, and with a shell error rather than a verdict. Found by forcing the fallback, not by reading the script. The no-sanitizer path now runs every byte check and says so; only overrun detection is lost. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- arc-tools/qtip_pack_host_check.sh | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/arc-tools/qtip_pack_host_check.sh b/arc-tools/qtip_pack_host_check.sh index cb3ebf59e..c8ddcb23e 100755 --- a/arc-tools/qtip_pack_host_check.sh +++ b/arc-tools/qtip_pack_host_check.sh @@ -222,7 +222,11 @@ else echo "asan: on (row allocated at exactly packed_bytes_per_row)" fi echo -"${CXX}" -std=c++17 -O1 -g -Wall -Wextra -Werror "${SAN[@]}" \ +# `${SAN[@]+...}` and not `${SAN[@]}`: under `set -u`, expanding an EMPTY array +# is an unbound-variable error in bash <= 4.3 (macOS ships 3.2), so the +# no-sanitizer fallback would abort the gate on exactly the runners that need +# it. Caught by forcing the fallback, not by reading the script. +"${CXX}" -std=c++17 -O1 -g -Wall -Wextra -Werror ${SAN[@]+"${SAN[@]}"} \ -I "${KDIR}" -o "${OUT}/qtip_pack_host_check" \ "${OUT}/qtip_pack_host_check.cpp" "${OUT}/qtip_pack_host_check" From 609714bf230fbfa2f0c432a71c8820b71910a11d Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 16:38:25 +0100 Subject: [PATCH 5/7] =?UTF-8?q?docs(qtip/bake):=20record=20the=20MEASURED?= =?UTF-8?q?=20res-usage=20at=20K=3D9=20=E2=80=94=20LOCAL:0,=20and=20REG:64?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `cuobjdump -res-usage`, CUDA 12.4, build.rs's exact flags, both matrix arches, CI run 32265671233 (green): sm_80 sm_90 K=9/V=4/L=12 LUT REG:64 SHARED:19632 LOCAL:0 REG:64 SHARED:20656 LOCAL:0 K=9/V=4/L=12 computed REG:61 SHARED:19632 LOCAL:0 REG:63 SHARED:20656 LOCAL:0 K=8/V=4/L=12 LUT REG:62 SHARED:19696 LOCAL:0 REG:63 SHARED:20720 LOCAL:0 K=8/V=4/L=12 computed REG:59 SHARED:19696 LOCAL:0 REG:62 SHARED:20720 LOCAL:0 K=4/V=2/L=16 LUT REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0 K=4/V=2/L=16 computed REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0 STACK:0 on all twelve. No spill anywhere: `cand[]` stayed in registers at K=9, where a naive `cand[2^K]` would have been 512 entries — the exact mis-scoping behind the retracted "~1,700 s/layer / 20 h / $30" estimate. Two things the table settles that no argument could: * The K=4 and K=8 sm_80 numbers are UNCHANGED to the byte from what #169 recorded before K=9 existed. The shipped rung's register allocation and shared-memory layout did not move when the packer became a bitstream. That is a compiled fact about codegen, stronger than the `if constexpr` source-identity argument it corroborates. * K=9's SHARED is exactly 64 B below K=8's at BOTH arches, which is the group table falling from 16 entries to 8 at 8 B each. The compiler arrived at the predicted number independently. ⚠ NEW LANDMINE, and the reason this is not just a green checkmark: K=9/LUT lands on REG:64 at both arches — exactly the cap `__launch_bounds__(256, 4)` imposes (65,536/4/256). It did not spill, so 4 blocks/SM holds, but it has ZERO headroom and is the tightest of the six. The next register added to its hot path is a spill, silently. Anything touching the K=9 expansion, the radix loop or the packer must re-run the gate and read the number. Also corrected while here: the escape-hatch note claimed dropping QB_MIN_BLOCKS_PER_SM to 3 "reproduces the measured allocation exactly and is a no-op". It does not — a looser cap (85 registers) leaves nvcc free to take more than the 59-64 measured at 4 blocks. It is still the right escape hatch; it is not allocation-preserving. The kernel source is unchanged since the measured run — comments only, verified by a comment-stripped diff — so the table describes this tree. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- mistralrs-quant/kernels/qtip/qtip_beam.cu | 62 ++++++++++++++++------- 1 file changed, 45 insertions(+), 17 deletions(-) diff --git a/mistralrs-quant/kernels/qtip/qtip_beam.cu b/mistralrs-quant/kernels/qtip/qtip_beam.cu index 43ac68719..94b207392 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -290,20 +290,45 @@ struct QbShape { // instead of the 24 that 3 blocks give (+33% occupancy, which for a // latency-bound kernel is a ~1.33x throughput term). // -// MEASURED, `cuobjdump -res-usage`, sm_80, CUDA 12.4, build.rs's exact flags -// (arc-tools/qtip_beam_res_usage_check.sh, which CI now runs on every PR): -// -// K=8/V=4/L=12 LUT REG:62 STACK:0 SHARED:19696 LOCAL:0 -// K=8/V=4/L=12 computed REG:59 STACK:0 SHARED:19696 LOCAL:0 -// K=4/V=2/L=16 LUT REG:60 STACK:0 SHARED:38000 LOCAL:0 -// K=4/V=2/L=16 computed REG:60 STACK:0 SHARED:38000 LOCAL:0 -// -// All four are inside the 64-register budget with NO spill, so the occupancy -// is real at both geometries. Both SHARED figures are exactly what the -// declarations above sum to (37,996 -> 38,000 and 19,696), which is an -// independent check that the K=4 layout did not move when the geometry became -// a template parameter. Shared memory is not the constraint either way: -// 4 x 38,000 B is 148 KiB of the 228 KiB an SM has. +// MEASURED, `cuobjdump -res-usage`, CUDA 12.4, build.rs's exact flags +// (arc-tools/qtip_beam_res_usage_check.sh, which CI runs on every PR — the +// table below is from run 32265671233, both matrix arches): +// +// sm_80 sm_90 +// K=9/V=4/L=12 LUT REG:64 SHARED:19632 LOCAL:0 REG:64 SHARED:20656 LOCAL:0 +// K=9/V=4/L=12 computed REG:61 SHARED:19632 LOCAL:0 REG:63 SHARED:20656 LOCAL:0 +// K=8/V=4/L=12 LUT REG:62 SHARED:19696 LOCAL:0 REG:63 SHARED:20720 LOCAL:0 +// K=8/V=4/L=12 computed REG:59 SHARED:19696 LOCAL:0 REG:62 SHARED:20720 LOCAL:0 +// K=4/V=2/L=16 LUT REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0 +// K=4/V=2/L=16 computed REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0 +// +// STACK:0 on all twelve. Every one is inside the 64-register budget with NO +// spill, so the occupancy is real at all three geometries — `cand[]` stayed in +// registers at K=9, where a naive `cand[2^K]` would have been 512 entries. +// +// Two things this table settles that no argument could: +// +// * THE K=4 AND K=8 sm_80 NUMBERS ARE UNCHANGED from before K=9 existed +// (REG:62/59 SHARED:19696 and REG:60/60 SHARED:38000, to the byte). The +// shipped rung's register allocation and shared-memory layout did not move +// when the packer became a bitstream — which is a compiled fact about +// codegen, not an inference from the `if constexpr` keeping its source +// identical. +// * K=9's SHARED is exactly 64 B BELOW K=8's at both arches (19,632 vs +// 19,696; 20,656 vs 20,720). That is the group table falling from 16 +// entries to 8 at 8 B each — the predicted number, arrived at +// independently by the compiler. +// +// ⚠ ZERO REGISTER HEADROOM AT K=9/LUT. It lands on REG:64 at both arches, +// which is exactly the cap `__launch_bounds__(256, 4)` imposes +// (65,536 / 4 / 256). It did NOT spill, so 4 blocks/SM holds — but it is the +// tightest of the six and the next register added to its hot path is a spill, +// not a warning. Anything touching the K=9 expansion, the radix loop or the +// packer must re-run the gate and read the number, not assume the margin that +// K=8 (REG:62/59) still has. +// +// Shared memory is not the constraint at any geometry: 4 x 38,000 B is 148 KiB +// of the 228 KiB an SM has. // // (The older `REG:80 ... SHARED:38992` line this note used to quote was a // pre-wave16 measurement of a kernel that no longer exists — wave16/wave17 @@ -323,9 +348,12 @@ struct QbShape { // mis-scoping that produced the discredited "~1,700 s/layer / 20 h / $30" // K=8/L=12 encode estimate. `.github/workflows/cuda_compile_check.yaml` runs // this check on every PR so it cannot be skipped; if it ever fails, set -// QB_MIN_BLOCKS_PER_SM back to 3 (which reproduces the measured allocation -// exactly and is a no-op) rather than trading a real latency regression for a -// nominal occupancy gain. +// QB_MIN_BLOCKS_PER_SM back to 3 rather than trading a real latency regression +// for a nominal occupancy gain — that raises the cap to 65,536/3/256 = 85 +// registers, which every geometry in the table above clears with room. (It is +// NOT necessarily allocation-preserving: with a looser cap nvcc is free to +// take more registers than the 59-64 measured at 4 blocks. Re-run the gate and +// read the numbers; do not assume the table still applies.) // // Related landmine: `unsigned int cand[CAND]` must stay in registers, which // requires EVERY loop indexing it to be fully unrolled. Weakening any of those From bc1994cb2861d621365721aff0736dc6f2b502b0 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 16:41:39 +0100 Subject: [PATCH 6/7] fix(qtip/bake): refuse a K=9 row whose stride is not 4-byte aligned; suspend the census absolutes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two corrections from the coordinator, both of which change what this branch asserts. 1. OOB READ, not a slow path. The serving side now PADS the packed row stride to 4 bytes and has DELETED its tail clamp, so a decoder reads a row in 4-byte units. `packed_bytes_per_row` returns `ceil(n*K/8)`, which at K=9 is `9*(n/8)` — a multiple of 4 only when `n % 32 == 0`, not `n % 8 == 0`. The whole-bytes-only guard would therefore have let the encoder emit a stride the decoder reads past the end of. The guard now also requires a 4-byte stride, but ONLY for straddling rungs (`K % 8 != 0`). K=4 and K=8 divide 8 and keep exactly the acceptance they ship with — tightening the rung that baked the published artifact to protect a rung that does not exist yet would trade a regression for a hypothetical. Refusing rather than padding is the point: where the bake is allowed to run, padded stride and unpadded stride are the SAME number, so the encoder cannot disagree with the decoder whichever of the two definitions serving is on. Adopting the padding would mean guessing a layout owned by another change, and a wrong guess is the OOB read this is avoiding. Every real V4-Flash shape passes untouched: 1792 and 512 symbols give 2016 B and 576 B at K=9, both already 4-aligned. Proven red: relaxing the new clause back to `is_multiple_of(8)` fails on `!packed_row_is_whole_bytes(8, 9)` — a 9-byte row, whole bytes but not 4-aligned, which is exactly the case that used to slip through. 2. THE CENSUS ABSOLUTES ARE SUSPENDED, so this branch stops quoting them. The `4.375 vs 15.125` / `3.46x` decode inst/weight figures came from the SASS census's differential-over-unroll-depth method, whose K=8 replay control did not reproduce: `#pragma unroll` is a hint, nvcc re-rolled the loops, all three unroll depths returned identical counts, and the linearity guard reads |d2-d1|/d1 — which is 0.00% precisely when all three are equal. It passed hardest when the rig was most broken. Struck from both kernel headers, with the reason recorded so nobody reinstates them from git history. The clamp DELTA (+1.126 K=9, +1.125 K=10) survives, being a difference between two kernels in the same mode. The structural claim — one symbol carries four weights, so fewer instructions per weight — does not depend on the census; the ratio did. The s/layer derivation in `qtip_beam.cu` does NOT inherit the suspension: it divides from a measured A100 bake wall-clock (FACTS.md:990), and the ~3.87 radix passes it references are a runtime histogram from `probe_beam_kernel_cost_drivers`, not a census differential. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- mistralrs-quant/kernels/qtip/qtip_beam.cu | 7 +- mistralrs-quant/kernels/qtip/qtip_geom.cuh | 23 +++++-- mistralrs-quant/src/qtip/mod.rs | 74 ++++++++++++++++++---- 3 files changed, 86 insertions(+), 18 deletions(-) diff --git a/mistralrs-quant/kernels/qtip/qtip_beam.cu b/mistralrs-quant/kernels/qtip/qtip_beam.cu index 94b207392..275df5a3f 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -10,8 +10,11 @@ // --------------------------------------------------- // Three rungs are compiled: // -// K=4 / V=2 / L=16 2.00 bpw, the shipped rung; 15.125 decode inst/weight -// K=8 / V=4 / L=12 2.00 bpw at 4.375 decode inst/weight +// K=4 / V=2 / L=16 2.00 bpw, the shipped rung +// K=8 / V=4 / L=12 2.00 bpw, fewer decode instructions per weight +// (the ABSOLUTE inst/weight figures are SUSPENDED — +// see qtip_geom.cuh; the census's unroll-differential +// method did not reproduce its own K=8 control) // K=9 / V=4 / L=12 2.25 bpw — THE BAKE TARGET. K=8/V=4/L=12 is // quality-CLOSED and K=9 is not; the sweep, its six // codebook designs and the Dw_cos figures are recorded diff --git a/mistralrs-quant/kernels/qtip/qtip_geom.cuh b/mistralrs-quant/kernels/qtip/qtip_geom.cuh index d90b27e66..88230d4f2 100644 --- a/mistralrs-quant/kernels/qtip/qtip_geom.cuh +++ b/mistralrs-quant/kernels/qtip/qtip_geom.cuh @@ -10,9 +10,21 @@ // 2 bits/weight: // // K=4 / V=2 / L=16 the shipped LUT rung ("qtip2") -// K=8 / V=4 / L=12 the same rate, 3.46x fewer decode instructions per -// weight (4.375 vs 15.125), because one symbol now -// carries four weights instead of two +// K=8 / V=4 / L=12 the same rate at materially fewer decode instructions +// per weight, because one symbol now carries four +// weights instead of two +// +// ⚠ The ABSOLUTE inst/weight figures this header used to quote (4.375 vs +// 15.125, "3.46x") are SUSPENDED and deliberately not restated. They came from +// the SASS census's differential-over-unroll-depth method, whose K=8 replay +// control failed to reproduce: `#pragma unroll` is a HINT, nvcc re-rolled the +// loops, all three unroll depths returned identical counts, and the linearity +// guard reads |d2-d1|/d1 — which is 0.00% precisely when all three are equal. +// The guard passed hardest when the rig was most broken. What survives from +// that census is the clamp DELTA (+1.126 at K=9, +1.125 at K=10), because a +// difference between two kernels in the same mode cancels the re-rolling. The +// direction "one symbol carries four weights, so fewer instructions per weight" +// is structural and does not depend on the census; the ratio does. // // and raising K at fixed V=4/L=12 is the ONLY way to buy quality on this // family, because `bpw = K/V`: @@ -125,9 +137,10 @@ struct QtipGeom { static_assert(V_ >= 1, "V is the reproduction-vector dimension"); }; -// The shipped LUT rung. 2 bits/weight, 15.125 decode instructions/weight. +// The shipped LUT rung. 2 bits/weight. (Its inst/weight absolute is suspended +// with the rest of the census — see the header.) using QtipGeomK4V2L16 = QtipGeom<4, 2, 16>; -// Same 2 bits/weight at 4.375 decode instructions/weight (3.46x fewer). +// Same 2 bits/weight, fewer decode instructions per weight (absolute suspended). using QtipGeomK8V4L12 = QtipGeom<8, 4, 12>; // 2.25 bits/weight on the same table and the same decode shape. The rung the // quality sweep actually selected; K=8/V=4/L=12 is closed. diff --git a/mistralrs-quant/src/qtip/mod.rs b/mistralrs-quant/src/qtip/mod.rs index 0e4c757d6..c834542d3 100644 --- a/mistralrs-quant/src/qtip/mod.rs +++ b/mistralrs-quant/src/qtip/mod.rs @@ -402,16 +402,46 @@ pub(crate) const fn packed_bytes_per_row(num_symbols: usize, k: u32) -> usize { (num_symbols * k as usize).div_ceil(8) } -/// Does a row of `num_symbols` symbols at `k` bits occupy a whole number of -/// bytes? +/// Is a row of `num_symbols` symbols at `k` bits a whole number of bytes AND a +/// multiple of 4 bytes? /// -/// At K=4 this is exactly the shipped "num_symbols must be even" rule. It is -/// kept at every rung because a partial trailing byte is a row stride the two -/// sides can disagree about for free, and every real V4-Flash / Qwen expert -/// shape satisfies it (`in_features` 7168 and 2048 give `num_symbols` 1792 and -/// 512 at V=4, both multiples of 8). Refuse, never round. +/// Whole bytes (`n*k % 8 == 0`) is, at K=4, exactly the shipped "num_symbols +/// must be even" rule. +/// +/// **The 4-byte part is what stops an out-of-bounds READ.** The serving side +/// pads the packed row stride to 4 bytes and has deleted its tail clamp, so a +/// decoder now reads a row as 4-byte units. `ceil(n*k/8)` — what +/// [`packed_bytes_per_row`] returns, and what `Rung::packed_bytes` returned +/// before the padding landed — is NOT a multiple of 4 in general: at K=9 it is +/// `9*(n/8)`, which needs `n % 32 == 0`, not `n % 8 == 0`. +/// +/// Refusing the misaligned shapes rather than padding them is deliberate. Where +/// the bake is allowed to run, padded stride and unpadded stride are the SAME +/// number, so the encoder cannot disagree with the decoder no matter which of +/// the two definitions the serving side is on. Adopting the padding instead +/// would mean guessing a layout owned by another change — and a wrong guess +/// here is an OOB read, not a slow path. +/// +/// **The 4-byte rule applies ONLY to the straddling rungs (`K % 8 != 0`).** +/// K=4 and K=8 keep exactly the acceptance they ship with — tightening the +/// rung that baked the published artifact, to protect a rung that does not +/// exist yet, would be a regression traded for a hypothetical. Their strides +/// are `n/2` and `n`, which the caller's own shapes already align. +/// +/// Every real V4-Flash expert shape passes: `in_features` 7168 and 2048 give +/// `num_symbols` 1792 and 512 at V=4, i.e. 2016 B and 576 B at K=9, both +/// already 4-byte aligned. Refuse, never round. pub(crate) const fn packed_row_is_whole_bytes(num_symbols: usize, k: u32) -> bool { - (num_symbols * k as usize).is_multiple_of(8) + let bits = num_symbols * k as usize; + if !bits.is_multiple_of(8) { + return false; + } + if (8 % k as usize) == 0 { + // K divides 8: the shipped acceptance, unchanged. + return true; + } + // Straddling rung: the stride must also be a multiple of 4 bytes. + bits.is_multiple_of(32) } // The general formula must reproduce the shipped K=4 nibble packing and the @@ -4445,15 +4475,37 @@ mod tests { } } - /// A partial trailing byte is refused, not rounded — at every rung. + /// A partial trailing byte is refused, not rounded — at every rung — and a + /// K=9 row that is not a multiple of 4 bytes is refused too. + /// + /// The serving side pads the packed row stride to 4 bytes and has deleted + /// its tail clamp, so a decoder reads a row in 4-byte units. `ceil(n*K/8)` + /// is `9*(n/8)` at K=9, which is a multiple of 4 only when `n % 32 == 0` — + /// so "whole bytes" alone would let the encoder emit a stride the decoder + /// reads past the end of. Refusing means padded and unpadded stride are the + /// same number wherever the bake runs, so the two sides cannot disagree. #[test] - fn a_row_that_is_not_whole_bytes_is_rejected() { + fn a_row_that_is_not_whole_bytes_or_not_4_byte_aligned_is_rejected() { + // The real V4-Flash expert shapes: 2016 B and 576 B, both 4-aligned. assert!(packed_row_is_whole_bytes(1792, 9)); assert!(packed_row_is_whole_bytes(512, 9)); + assert_eq!(packed_bytes_per_row(1792, 9) % 4, 0); + assert_eq!(packed_bytes_per_row(512, 9) % 4, 0); + // Not whole bytes. assert!(!packed_row_is_whole_bytes(1793, 9)); - // K=4's rule was "num_symbols must be even"; this must still be it. + // Whole bytes (n % 8 == 0) but NOT a multiple of 4 bytes: 8*9/8 = 9 B. + // This is the case a whole-bytes-only rule would have let through. + assert!(packed_bytes_per_row(8, 9) % 4 != 0); + assert!(!packed_row_is_whole_bytes(8, 9)); + assert!(!packed_row_is_whole_bytes(24, 9)); + assert!(packed_row_is_whole_bytes(32, 9)); + + // K=4 and K=8 divide 8, so they keep EXACTLY the acceptance they ship + // with. Tightening them here would refuse rows the published bake + // accepts. for n in 0..64usize { assert_eq!(packed_row_is_whole_bytes(n, 4), n.is_multiple_of(2)); + assert!(packed_row_is_whole_bytes(n, 8)); } } From 755cf0aa937c12428482a478a7c5c5c9d98e676a Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Wed, 19 Aug 2026 16:42:25 +0100 Subject: [PATCH 7/7] fix(clippy): manual_is_multiple_of in the new stride-alignment test MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Caught only after the previous push, because the verification chain piped clippy into `tail` — which makes the pipeline's exit status `tail`'s, so a red clippy read as green and the `&&` ran on anyway. Verified here with real exit codes: clippy rc=0, test rc=0, host gate rc=0. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- mistralrs-quant/src/qtip/mod.rs | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/mistralrs-quant/src/qtip/mod.rs b/mistralrs-quant/src/qtip/mod.rs index c834542d3..8c60322e0 100644 --- a/mistralrs-quant/src/qtip/mod.rs +++ b/mistralrs-quant/src/qtip/mod.rs @@ -4489,13 +4489,13 @@ mod tests { // The real V4-Flash expert shapes: 2016 B and 576 B, both 4-aligned. assert!(packed_row_is_whole_bytes(1792, 9)); assert!(packed_row_is_whole_bytes(512, 9)); - assert_eq!(packed_bytes_per_row(1792, 9) % 4, 0); - assert_eq!(packed_bytes_per_row(512, 9) % 4, 0); + assert!(packed_bytes_per_row(1792, 9).is_multiple_of(4)); + assert!(packed_bytes_per_row(512, 9).is_multiple_of(4)); // Not whole bytes. assert!(!packed_row_is_whole_bytes(1793, 9)); // Whole bytes (n % 8 == 0) but NOT a multiple of 4 bytes: 8*9/8 = 9 B. // This is the case a whole-bytes-only rule would have let through. - assert!(packed_bytes_per_row(8, 9) % 4 != 0); + assert!(!packed_bytes_per_row(8, 9).is_multiple_of(4)); assert!(!packed_row_is_whole_bytes(8, 9)); assert!(!packed_row_is_whole_bytes(24, 9)); assert!(packed_row_is_whole_bytes(32, 9));