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/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/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/arc-tools/qtip_pack_host_check.sh b/arc-tools/qtip_pack_host_check.sh new file mode 100755 index 000000000..c8ddcb23e --- /dev/null +++ b/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 +#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 +# `${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" 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 96f3e8705..275df5a3f 100644 --- a/mistralrs-quant/kernels/qtip/qtip_beam.cu +++ b/mistralrs-quant/kernels/qtip/qtip_beam.cu @@ -8,10 +8,17 @@ // // 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 +// 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 +// 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: @@ -30,17 +37,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. @@ -86,6 +105,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). @@ -210,10 +272,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[]`. @@ -230,20 +293,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): +// 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): // -// 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 +// 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 // -// 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. +// 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 @@ -263,9 +351,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 @@ -763,6 +854,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..88230d4f2 100644 --- a/mistralrs-quant/kernels/qtip/qtip_geom.cuh +++ b/mistralrs-quant/kernels/qtip/qtip_geom.cuh @@ -10,9 +10,35 @@ // 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`: +// +// 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 (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 @@ -23,6 +49,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 +95,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. @@ -59,28 +137,71 @@ 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. +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..8c60322e0 100644 --- a/mistralrs-quant/src/qtip/mod.rs +++ b/mistralrs-quant/src/qtip/mod.rs @@ -385,6 +385,88 @@ 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) +} + +/// Is a row of `num_symbols` symbols at `k` bits a whole number of bytes AND a +/// multiple of 4 bytes? +/// +/// 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 { + 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 +// 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 +4343,172 @@ 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 — 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_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!(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).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)); + + // 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)); + } + } + /// Sanity: LUT has the right size and reasonable distribution. #[test] fn gaussian_lut_is_well_formed() {