feat(qtip/bake): bake K=9/V=4/L=12 — the packed row becomes a bitstream

.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 }}
 

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** —

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

arc-tools/qtip_pack_host_check.sh

@@ -0,0 +1,232 @@
+#!/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 <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <vector>
+
+// 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<uint8_t> ref_pack(const std::vector<uint32_t>& syms, uint32_t k) {
+    const size_t bytes = (syms.size() * (size_t)k + 7) / 8;
+    std::vector<uint8_t> 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<uint8_t> nibble_pack_k4(const std::vector<uint32_t>& syms) {
+    std::vector<uint8_t> 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 <class G>
+static int check(const char* name, int n) {
+    std::vector<uint32_t> 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<G>(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<G>(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<QtipGeomK4V2L16>("K4V2L16 n=3584", 3584);
+    bad |= check<QtipGeomK4V2L16>("K4V2L16 n=1024", 1024);
+    bad |= check<QtipGeomK8V4L12>("K8V4L12 n=1792", 1792);
+    bad |= check<QtipGeomK8V4L12>("K8V4L12 n=512",  512);
+    bad |= check<QtipGeomK9V4L12>("K9V4L12 n=1792", 1792);
+    bad |= check<QtipGeomK9V4L12>("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<QtipGeom<6, 3, 12>>("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<uint32_t> s(4096);
+    for (size_t i = 0; i < s.size(); ++i) s[i] = nxt() & 0xF;
+    std::vector<uint8_t> got(s.size() / 2, 0);
+    for (size_t t = 0; t < s.size(); ++t) {
+        qtip_pack_symbol<QtipGeomK4V2L16>(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 2>/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
+# `${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"

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 /