Repository navigation
feat(qtip): V=4/L=12 decode family — K as a parametric seam, bit-exact CPU parity - #168
Merged
Merged
Conversation
…reference
Stage 1 of making K8/V4/L12 a selectable rung: the decode geometry, its
reproduction table, the serving kernel, and the reference the GPU is held to.
No format, no dispatch, no quantizer yet — those are stages 2 and 3.
`bpw = K/V`, so this is 2 bits per weight, the same rate the shipped
K=4/V=2/L=16 rung produces. Nothing here compresses harder. What changes is
decode cost: the table is 2^12 x 4 bf16 = 32,768 B instead of 2^16 x 2 f32 =
524,288 B, which fits static __shared__ with no cudaFuncSetAttribute opt-in,
and K=8 makes a symbol exactly one byte so the nibble unpack disappears.
Added:
* mistralrs-quant/src/qtip/k8v4l12.rs — geometry, bf16 table, CPU decode
reference, and `gemv_row_gpu_model`: a bit-exact model of the kernel
INCLUDING its reduction tree (per-thread slicing, warmup seeding, warp
XOR butterfly, cross-warp butterfly). Modelling the tree is what lets the
CUDA gate compare bits with `==` instead of a tolerance — a tolerance
there would be where a mis-seeded thread hides behind benign
reassociation.
* kernels/qtip/qtip_gemv_k8v4l12.cu — the serving kernel, grid-strided over
rows so the 32 KiB stage is amortised per block rather than per row.
* FFI + `fused_gemv_k8v4l12_cuda`, which refuses a K=4/V=2 artifact three
ways (F32 table, 2^16x2 table, nibble-packed rows) rather than decoding it
into plausible garbage.
RowScaleHoist is its own switch because it is the only lever that costs
bit-exactness: it reassociates the sum. Parity runs with it Off.
NOT MEASURED: the 5.375 / 4.375 inst-per-weight figures came from a standalone
compiled probe, not from this kernel. This kernel has never been compiled with
nvcc and never run. Its instruction count is unestablished.
Guards were each shown red before being trusted (mutations M1-M15). Three
holes found and closed that way: the table's contents were unpinned (a mutation
rounding every value to the nearest integer passed all 16 tests — now a
digest), every fixture divided evenly by the thread count so a floored
sym_per_thread dropped row tails undetected, and nothing tied KERNEL_THREADS
to the .cu (now a whitespace-normalised source guard that catches drift from
either side).
kernels/EXPECTED_KERNEL_COUNT 39 -> 40, as its own guard instructed.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
Code Metrics Report━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Language Files Lines Code Comments Blanks ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ C Header 5 305 210 52 43 CSS 2 1181 1036 34 111 CUDA 73 25129 18043 4319 2767 Dockerfile 1 39 22 8 9 JavaScript 16 3546 2676 482 388 Jinja2 7 694 656 5 33 JSON 74 4600 4597 0 3 Makefile 1 6 5 0 1 Metal Shading Lan| 33 12224 9431 1142 1651 PowerShell 1 300 227 30 43 Python 145 15139 12482 811 1846 Shell 39 9777 6538 2596 643 Plain Text 4 3801 0 2479 1322 TOML 33 1498 1294 54 150 YAML 3 25 23 2 0 ───────────────────────────────────────────────────────────────────────────────── HTML 4 2687 2604 43 40 |- CSS 2 543 479 37 27 |- JavaScript 1 1233 1215 12 6 (Total) 4463 4298 92 73 ───────────────────────────────────────────────────────────────────────────────── Jupyter Notebooks 4 122 83 23 16 |- Markdown 1 60 30 22 8 |- Python 1 122 113 1 8 (Total) 304 226 46 32 ───────────────────────────────────────────────────────────────────────────────── Markdown 203 44777 0 34735 10042 |- BASH 72 1654 1202 331 121 |- C 3 17 17 0 0 |- CUDA 2 84 56 16 12 |- JSON 18 708 708 0 0 |- PowerShell 1 1 1 0 0 |- Python 23 1008 787 113 108 |- Rust 66 2051 1716 77 258 |- TOML 6 207 164 0 43 |- YAML 5 41 36 5 0 (Total) 50548 4687 35277 10584 ───────────────────────────────────────────────────────────────────────────────── Rust 672 328294 283104 16426 28764 |- Markdown 490 28532 471 24619 3442 (Total) 356826 283575 41045 32206 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1320 490405 350026 88474 51905 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
…raction K=8/V=4/L=12 is quality-closed: CPU sweeps put it at Delta w_cos -0.00698 (random codebook) and -0.00307 (converged trellis-Lloyd) against a +/-0.0008 ship threshold, and six codebook designs failed to recover it. K=9 at 2.25 bpw measured +0.00402. So K has to move, and this restructures the rung so that it can move as a specialisation rather than a rewrite. The seam is symbol extraction, and only symbol extraction. L=12 and V=4 fix the table at 32,768 B, and the table does not read K at all — so K changes the bit-extract, the packed row length, the state shift width and the warmup depth, and touches nothing else. `Rung` (Rust) and `QtipSymExtract<K>` (CUDA) are the only K-dependent code. k8v4l12.rs -> trellis_v4l12.rs qtip_gemv_k8v4l12.cu -> qtip_gemv_v4l12.cu kernel is now template <typename T, int THREADS, int K, bool ROW_SCALE_HOIST> launchers take a runtime `k`, dispatching to K=8/9/10 instantiations K=8 stays as the byte-aligned CONTROL with its bit-exact gate intact: a symbol IS a byte, so `QtipSymExtract<8>` is a real specialisation (one LDG, no shift, no mask, no clamp) and is the floor the others are measured against. K=9 and K=10 are carried so the alignment penalty can be measured on TWO non-byte-aligned points, which is what separates "9 is unlucky" from "non-alignment costs C". Bit layout is now stated and pinned: symbol t is bits [tK, tK+K) LSB-first, which at K=4 reproduces the shipped rung's "sym 2b is the low nibble of byte b" exactly. `pack`/`extract` round-trip at every K, and a fixture-free test drives the K=4 rule directly against hand-packed bytes. NOT COMMITTING TO K=9 CODEGEN. The K=9/K=10 instantiations exist for correctness coverage and to give the SASS pass a real serving kernel to measure instead of a probe. No performance claim: this kernel has never been compiled with nvcc or run, at any K. One thing the SASS measurement must not be misread on: the multi-byte extract reads a COMPILE-TIME byte count and clamps the index so it cannot run past the last row. The clamp is correct (bytes beyond ceil((off+K)/8) land above the K-bit mask window, so a substituted in-row byte cannot change the symbol) but it costs a `min` per byte read and would DISAPPEAR if the row stride were padded by MAX_BYTES-1. A K=9 number that includes it is an upper bound on the alignment penalty, not the penalty. Padding is a format decision deliberately not taken while K is unsettled. Flagged to the SASS agent. Also states plainly, per the coordinator, what the 32 KiB table is FOR: the shipped default (`QtipCodebook::DEFAULT = Gaussian`) gathers from a 524,288 B f32 table that does not fit shared memory at any occupancy, so every decoded symbol pays a dependent scattered L2 load — measured 388 GB/s, ~8% of H200 HBM, and it is the decode limiter. Guards shown red before being trusted (E1-E15). Two holes found: nothing tied a `case k:` label to the template argument it dispatches to (`case 9: LAUNCH_K(T, 8, ..)` would decode K=9 at K=8, in bounds and no fault), and nothing refused a kernel case for a K with no CPU model. One of my own mutations was again a no-op — a `default: return;` placed before `case 9:` does not make it unreachable in C++ — and was redone as a real deletion. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
Two findings from the K9 alignment census, acted on before its counts land.
Both are UNMEASURED as speedups — they reduce operation counts that can be
derived exactly, and the derivations are pinned as tests.
1. THE WARMUP TAX WAS THE BIGGER PROBLEM, AND IT IS K-INDEPENDENT
Every lane replays `warmup_syms` symbols before decoding its own slice, so
extractions per weight are (S + W)/(S*V) against an ideal of 1/V — a tax of
1 + W/S, set entirely by slice length S = num_symbols / lanes_per_row. The
kernel split a row across all 128 threads, which makes S short:
in_features 128 lanes/row 32 lanes/row
4096 S=8, 1.25x S=32, 1.06x
1024 S=2, 1.99x S=8, 1.24x
512 S=1, 2.98x S=4, 1.48x
So a WARP now owns a row and a block owns THREADS/32 rows at once. Same
parallelism, four times the slice. It also deletes the cross-warp butterfly,
the `warp_sums` shared array and BOTH `__syncthreads` from the row loop, since
a warp reduces with shuffles alone — and it simplifies the bit-exact CPU model
to a single 32-lane butterfly.
This multiplies the extraction count at every K, so it scales the K=9 alignment
penalty by the same factor rather than competing with it.
Correcting one number in the census: it assumed THREADS=256 and reported 1.5x
at in_features=4096. This kernel is THREADS=128, so it was 1.25x, now 1.06x.
The direction and the argument are unchanged.
2. K=10 NEEDS TWO BYTES, NOT THREE
The bit offset of symbol t is (t*K) mod 8, which only ever takes the multiples
of gcd(K, 8). The worst REACHABLE offset is 8 - gcd(K,8), not 7. At K=10 the
offsets are {0, 2, 4, 6}, so off+K <= 16 and two bytes always suffice; the
naive ceil((7+K)/8) said three and had every symbol read, shift and discard a
wasted byte. The bound is now gcd-based on both sides, and the test proves it
by enumerating reachable offsets rather than trusting the formula.
Still not committing to K=9 codegen; still no compiled measurement of this
kernel at any K. The tail clamp remains and remains removable by padding the
row stride — the census is pricing it separately, which is the right call.
Guards shown red (N1-N9), including three that would let the model and the
kernel drift apart: a model that splits by block size, a kernel that reverts to
splitting a row across the block, and a kernel that reintroduces a cross-warp
reduction the model no longer has.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
heydryft
changed the title
… separately
The census measures inst/weight as a DIFFERENTIAL OVER UNROLL DEPTH. That
cancels prologue and epilogue exactly — which is what makes it robust, and is
also why it is structurally blind to per-row and per-block cost.
So the warp-per-row reshape's reduction half is invisible to it: the deleted
cross-warp butterfly, the `warp_sums` array and both `__syncthreads` can
neither be credited nor charged by that instrument. **A flat census number
therefore does not mean the reshape did nothing**, and this commit makes that
impossible to misread six weeks from now:
* The module docs state the blindness and where the reshape DOES show up
(indirectly, through the extraction rate).
* `structural_per_row_overhead()` gives that half an accounting of its own —
(barriers, warp butterflies, shared round-trips) = (0, 1, 0), against
(2, 2, 2) before. STRUCTURAL COUNTS FROM SOURCE, never presented as
instruction counts; nothing here has been compiled.
* `route_cost_per_weight_x10000()` pins the arithmetic the incoming census
result gets run through, so substituting one of its illustrative endpoints
(0.25 / 0.375 / 0.50) for this kernel's real rate cannot happen silently.
This kernel runs at 0.2651 at in_features=4096 and 0.3105 at 1024.
Also acts on the census's confirmation that the gcd bound makes two bytes
EXACT, not merely sufficient, for K=9 and K=10 — so the 3-byte extraction route
is unreachable for every K anyone has. That is now enforced at compile time
(`static_assert(QtipSymExtract<K>::MAX_BYTES <= 2)`) so a future K that needs
three is a deliberate decision rather than something a new `case` quietly
turns on.
MY OWN BUG, CAUGHT BY THE GUARD I WAS WRITING: the new "exactly one
__syncthreads" check counted raw occurrences, which also counts the two places
the header *describes* barriers — it reported 2 against a kernel that has 1.
The obvious "fix" is to assert 2, after which it would never catch a real
barrier again. Fixed properly by stripping comments first, and pinned by a
mutation pair: a real barrier in the row loop goes red, a barrier merely
MENTIONED in a new comment stays green.
Guards shown red (C1, C3, C4, C5) with C2 verified green-by-design.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
…solutes FORMAT DECISION, TAKEN ON THE ONE NUMBER THAT SURVIVES SCRUTINY Rows are now allocated at `row_stride = align_up(data_bytes + MAX_BYTES-1, 4)` and the kernel's tail clamp is gone. Measured at sm_90 by FULL-KERNEL differencing, the clamp cost +1.126 inst/weight at K=9 and +1.125 at K=10 (padded 6.312 -> clamped 7.438, and 6.250 -> 7.375) — about +4.50 per extraction. Two independent geometries agreeing to 0.1%. It is trustworthy where the absolute figures are not, because it is a DIFFERENCE BETWEEN TWO KERNELS IN THE SAME MODE and therefore does not depend on a baseline. Run through this kernel's own rates that is +1.193 inst/weight at in_features=4096 and +1.397 at 1024. Cost of padding: at most 4 bytes per row — 0.35% at 4096, 0.20% at 7168. The round-up to 4 is NOT tail safety, which needs only MAX_BYTES-1. It makes every row base 4-byte aligned, which the cheapest measured route (the funnel route) requires. The stride is format; changing it twice is the expensive outcome, so it is chosen once to support both. K=8 IS UNTOUCHED. `max_bytes_per_symbol() == 1` means no tail overrun to pad against and no multi-byte funnel to align for, so the control keeps byte-for- byte the layout it was probed with and the padding decision cannot perturb it. `packed_bytes` is split into `data_bytes` (governed by the bit rate) and `row_stride` (what a tensor is allocated at and what the UQFF validates), because conflating them is how a bpw claim ends up 4 bytes wrong. ABSOLUTE INSTRUCTION COUNTS ARE NOW MARKED PROVISIONAL — INCLUDING OURS The isolated micro-census was withdrawn as invalid: its K=8 control measured 100% non-linear, because K=8's contiguous-byte symbols let ptxas merge loads while K=9's (9t)>>3 indices cannot, so every "vs K=8" per-symbol delta was against a moving baseline. The full-kernel rig has not yet reproduced the K=8 control's published 5.375 / 4.375 either. So this family's own headline figures — 5.375, 4.375, and the 15.125 they are compared against — are now labelled unreproduced in the module and kernel docs, not just the new route figures. APPLYING THEIR FAILURE TO MY OWN GUARDS Their rig reported 0.000 inst/weight with verdict "OK", including the control that exists to prove it: nvcc re-rolled the loops, and the linearity check read |d2-d1|/d1, which is 0.00% when all three counts are equal. The guard passed HARDEST when the instrument was most broken. My reassociation guard had the same shape — `|a-b| <= 1e-5 * L1` is satisfied by `0 == 0`, so an all-zero decoder would have passed it comfortably. It now asserts non-degeneracy BEFORE agreement: there must be something before it can agree. Guards shown red (P1-P6), including a stride that pads for the tail but drops alignment, one that aligns but forgets the tail (an out-of-bounds read), and a kernel that reintroduces the clamp. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Stage 1 of 3. The V=4/L=12 decode family — its table, its serving kernel, and the CPU reference the GPU is held to — with K as a parametric seam. No UQFF format (#170), no dispatch (#172).
Every absolute inst/weight figure attached to this family is currently PROVISIONAL, including the ones that justified it.
(9t)>>3indices cannot. Every "vs K=8" delta from it was against a moving baseline.32264842287.nvccor run on a GPU by its author, at any K.Exactly one measured number here is trusted, and the format decision below rests on it alone.
Why the family
L=12andV=4fix the table at2^12 × 4bf16 = 32,768 B, under the 48 KiB static__shared__limit, so a kernel stages it per block and every lookup is an LDS. The shipped rung's table is2^16 × 2f32 = 524,288 B — it does not fit shared memory at any occupancy, so every decoded symbol pays a dependent, scattered L2 load. That is the shipped default path (QtipCodebook::DEFAULTisGaussian, which gathers). Killing that gather is what the 32 KiB table is for.The table does not read K at all. So K changes symbol extraction, the packed row length, the state shift width and the warmup depth — nothing else.
RungandQtipSymExtract<K>are the only K-dependent code.K=8 is the control, not the shipping candidate — quality-closed at Δw_cos −0.00698 / −0.00307 against a ±0.0008 band. K=9 measured +0.00402, 5× better than the shipped control. That evidence is CPU-side and is not in doubt.
The format decision: rows are padded, the clamp is gone
row_stride = align_up(data_bytes + MAX_BYTES − 1, 4).Taken on the one number that survives scrutiny: compiled at sm_90 by full-kernel differencing, the tail clamp cost +1.126 inst/weight at K=9 and +1.125 at K=10 — two independent geometries agreeing to 0.1%. It is trustworthy where the absolutes are not because it is a difference between two kernels in the same mode, so it does not ride on a baseline. Through this kernel's own rates: +1.193 inst/weight @
in_features=4096, +1.397 @ 1024.Padding costs at most 4 bytes per row — 0.35% at 4096, 0.20% at 7168.
The round-up to 4 is not tail safety (which needs only
MAX_BYTES − 1); it makes every row base 4-byte aligned, which the cheapest measured route (funnel) requires. The stride is format, and changing it twice is the expensive outcome, so it is chosen once to support both.K=8 is untouched:
MAX_BYTES == 1⇒ no pad, so the control keeps byte-for-byte the layout it was probed with and this decision cannot perturb its replay.packed_bytesis split intodata_bytes(bit-rate governed) androw_stride(allocated, validated), because conflating them is how a bpw claim ends up 4 bytes wrong.Two other operation-count fixes
One warp owns a row, not one block. Warmup replay makes extractions/weight
(S+W)/(S·V)— a tax of1 + W/Sset by slice length. 128 lanes/row gave 1.25× at 4096 and 2.98× at 512; 32 lanes/row gives 1.06× and 1.48×. It also deletes the cross-warp butterfly,warp_sums, and both__syncthreadsfrom the row loop.structural_per_row_overhead()gives it a separate accounting —(0,1,0)vs(2,2,2)— as structural counts from source, never instruction counts.Two bytes is exact, not merely sufficient. Offsets are multiples of
gcd(K,8), so the bound isceil((8 − gcd(K,8) + K)/8): K=8→1, K=9→2, K=10→2. The 3-byte route is unreachable for every K anyone has, now enforced bystatic_assert.Parity: no tolerance
gemv_row_gpu_modelmodels the arithmetic and the reduction tree, so the CUDA gate compares bits with==. Bit layout is pinned:[tK, tK+K)LSB-first, verified against the shipped K=4 nibble order.Guards shown red (M1–M15, E1–E15, N1–N9, C1–C5, P1–P6)
Eight holes found this way. The newest: my reassociation guard had the same shape as the failure that invalidated the census rig —
|a−b| ≤ 1e-5·L1is satisfied by0 == 0, so an all-zero decoder passed it comfortably. (Their rig reported 0.000 inst/weight with verdict "OK" because|d2−d1|/d1is 0.00% when all counts are equal: steadiness checked without growth.) It now asserts non-degeneracy before agreement.Earlier ones: the table's contents were unpinned; every fixture divided evenly by the thread count; nothing tied
KERNEL_THREADSto the.cu; nothing tied acase k:label to its template argument; a barrier guard counted the word__syncthreadsin a comment and reported 2 against a kernel with 1 — whose tempting "fix" would have blinded it permanently.Scoped clippy clean;
cargo test -p mistralrs-quant323 + 5 passing.What remains
nvcclanes queued.