Conversation
#169 made the beam encoder geometry-parametric and it does NOT cover K=9. Two `static_assert`s in `QtipGeom` refuse it outright — `K_ <= 8` and `8u % K_ == 0u` ("K must divide 8: the packer emits whole bytes") — and behind them `SYMS_PER_BYTE = 8/K` is 0, so `qtip_packed_bytes_per_row` divides by zero and `qtip_pack_symbol` cannot address a symbol that straddles two bytes. The Rust side has the same hole: `cuda_ops.rs` computed `8 / super::K` and would have panicked before a kernel ran. That is the whole gap. Everything the K=8 work was actually about survives K=9 untouched, which is worth stating precisely because it is also the cost argument: * `LANES = min(2^K, 256 / 2^(L-K))` gives 32 at K=9/L=12 (8 groups, 512-symbol alphabet), so `CAND = 512/32 = 16` — the SAME register array as K=4 (16/1) and K=8 (256/16). No `cand[512]`, no spill. * `n_cand = ng * 2^K` is 2^L-saturated at 4096 at all three geometries, so the per-timestep candidate count, the ~3.87 radix passes, the 32-bit cost key and the `2^K * V * 4 B`-per-group LUT traffic (65,536 B/timestep) are invariant. * `MAX_GROUPS` falls 16 -> 8 and the group table 128 B -> 64 B; the 16 KiB backtrace stage floor is unchanged. * `num_symbols = in_features / V` is K-independent. So the packing is the only thing that moves, and it moves to the format the serving side already reads: symbol `t` at bits `[t*K, t*K + K)`, LSB-first, row `ceil(num_symbols*K/8)` bytes — `trellis_v4l12.rs::Rung::{pack, extract, packed_bytes}`. No second convention is invented; the UQFF geometry section is `[tag=3, K, L, V]`, so K=8 -> K=9 is one byte of artifact change. The straddle bound is `ceil((8 - gcd(K,8) + K)/8)`, not `ceil((7+K)/8)`: the reachable bit offsets are multiples of `gcd(K,8)`, so 7 is unreachable at K=10 and two bytes are EXACT at both K=9 and K=10. The packer writes only the bytes a symbol's bits fall in, which is what keeps the last symbol of a row inside `ceil(num_symbols*K/8)`. K=4/V=2/L=16 STAYS BYTE-IDENTICAL, and not by assertion: the byte-aligned expression is kept verbatim under `if constexpr`, so K=4 and K=8 compile from unchanged source rather than from a general form believed to fold back to it. Five host tests pin the rule the discarded branch would produce (the bitstream reproduces the shipped nibble packing over 4096 symbols; it is a plain byte store at K=8; K=9 round-trips through an independently written extractor at the real 1792-symbol V4 shape including the last symbol; the straddle bound is exact at K=4/8/9/10; a partial trailing byte is refused, and at K=4 that rule is still exactly "num_symbols must be even"). A `const` block asserts at COMPILE time that the general formula reproduces `num_symbols/2` at K=4 and `num_symbols` at K=8 at the real V4 shapes. All five were proven red: M1 (packed_bytes_per_row assumes one byte per symbol) fails to compile on the const assert; M2 (one-byte extract) reds only the K=9 test; M3 (whole-byte guard always true) reds only the partial-byte test; M4 (naive ceil((7+K)/8)) reds only the straddle test at K=10; M5 (nibble order flipped) reds only the K=4 identity test. `arc-tools/qtip_beam_res_usage_check.sh` now asserts SIX beam kernels, not four — a "LOCAL:0" verdict must be about a build that still contains the K=9 instantiation, which is precisely the one a naive parameterisation would spill. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
…ed bytes
"K=4/V=2/L=16 stays byte-identical" was, until this, an argument. The
instruments that could settle it are `cuda_beam_matches_cpu_beam_bit_for_bit`
and `cuda_beam_unpruned_matches_cuda_exhaustive`, and both are
`#[cfg(feature = "cuda")]` and SKIP without a device — so on the Mac dev box
and on every GPU-less runner, nothing checked the packer at all. The packer is
integer arithmetic on a byte array, though: stub `__device__` /
`__forceinline__` and the exact source text nvcc compiles also compiles with
clang++/g++.
`arc-tools/qtip_pack_host_check.sh` #includes the real header and asserts, at
the real V4-Flash expert shapes (in_features 7168 and 2048):
* K=4 emits bytes IDENTICAL to the shipped nibble packer over 4096 symbols —
compared against a transcription of `cpu_reference_packed`, not against the
new wire rule, so it is the shipping guarantee that is being checked.
* K=4/K=8/K=9 all match an independently written bitstream reference.
* `packed_bytes_per_row` equals `ceil(num_symbols*K/8)` at every rung.
Two things the obvious version of this script would have got wrong, both found
by running the mutations rather than reasoning about them:
1. A GUARD BAND CANNOT SEE THIS OVERRUN. Dropping the packer's
`if (8u*b < used)` writes `|= 0` past the row — the byte value does not
change, so zeros or any sentinel look untouched. The row is therefore an
EXACT-size heap allocation and the build asks for `-fsanitize=address`,
which reports it as a heap-buffer-overflow. The script says which mode it
ran in rather than implying coverage it lacks. The overrun is not benign
in the kernel: rows are contiguous and different blocks own different rows
concurrently, so a non-atomic RMW on the next row's first byte can clobber
that row's own store.
2. `QtipGeom<6,3,12>` is instantiated purely to make that guard non-vacuous.
At K=9 and K=10 every reachable offset needs exactly 2 bytes, so the guard
never fires and would be untested. K=6 has `off==0 => used==6`, one byte
against a bound of two.
Proven red: `MAX_BYTES_PER_SYMBOL -> 1` (K=9 and K=6 truncate); nibble order
flipped (K=4 identity fails); `packed_bytes_per_row -> num_symbols` (K=9 length
fails); guard dropped (ASan heap-buffer-overflow).
Also recorded, because a gate's own docs must not overclaim: enlarging
`MAX_BYTES_PER_SYMBOL` to the naive `ceil((7+K)/8)` does NOT go red here and
should not — the per-symbol `used` guard makes an over-large bound a dead
unrolled iteration. The bound's exactness matters where it is a READ width, on
the serving side, and is pinned by `the_straddle_bound_is_exact_at_every_baked_k`.
Wired into the CUDA lane ahead of the spill gate: it is the cheaper failure and
needs no CUDA toolkit.
Co-Authored-By: Claude Opus 5 (1M context) <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 25015 17952 4311 2752 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 490291 349935 88466 51890 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
…oved Three doc-level consequences of the K=9 change, none of them cosmetic: * `QUANTIZATION_PERFORMANCE.md:34` cites `qtip_beam.cu:67-79` for the "3.2e9 symbol positions per V4-Flash layer at V=2" figure. Adding the encode-cost section to that header pushed the cited block to `:83-95`. A stale `path:line` in a doc whose whole premise is "nothing is stated without a label" is worse than no label. * `CUDA_VALIDATION.md` gains gate 1d. Its table is the answer to "what is checked for free"; a gate that is not in it does not exist to the next agent. * The quality figures behind the rung choice (-0.00698 / -0.00307 / +0.00402) are now stated ONCE, in `qtip/trellis_v4l12.rs`'s module docs, and pointed at from the kernel headers instead of restated there. Two copies of a measured number drift. The kernel headers also stop citing `qtip_beam.cu` line numbers for the two lines that falsify `(n/V)*W*2^K` — they name the constructs (`& G::GROUP_MASK` in step 1a, `n_cand = ng * 2^K` at the end of step 2) instead, since those line numbers move every time the header above them grows. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
`set -u` plus an EMPTY bash array is an unbound-variable error in bash <= 4.3
(macOS ships 3.2), so `"${SAN[@]}"` killed the script the moment the
`-fsanitize=address` probe failed — i.e. on exactly the GPU-less runners the
gate exists to serve, and with a shell error rather than a verdict.
Found by forcing the fallback, not by reading the script. The no-sanitizer path
now runs every byte check and says so; only overrun detection is lost.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
…REG:64
`cuobjdump -res-usage`, CUDA 12.4, build.rs's exact flags, both matrix arches,
CI run 32265671233 (green):
sm_80 sm_90
K=9/V=4/L=12 LUT REG:64 SHARED:19632 LOCAL:0 REG:64 SHARED:20656 LOCAL:0
K=9/V=4/L=12 computed REG:61 SHARED:19632 LOCAL:0 REG:63 SHARED:20656 LOCAL:0
K=8/V=4/L=12 LUT REG:62 SHARED:19696 LOCAL:0 REG:63 SHARED:20720 LOCAL:0
K=8/V=4/L=12 computed REG:59 SHARED:19696 LOCAL:0 REG:62 SHARED:20720 LOCAL:0
K=4/V=2/L=16 LUT REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0
K=4/V=2/L=16 computed REG:60 SHARED:38000 LOCAL:0 REG:63 SHARED:39024 LOCAL:0
STACK:0 on all twelve. No spill anywhere: `cand[]` stayed in registers at K=9,
where a naive `cand[2^K]` would have been 512 entries — the exact mis-scoping
behind the retracted "~1,700 s/layer / 20 h / $30" estimate.
Two things the table settles that no argument could:
* The K=4 and K=8 sm_80 numbers are UNCHANGED to the byte from what #169
recorded before K=9 existed. The shipped rung's register allocation and
shared-memory layout did not move when the packer became a bitstream. That
is a compiled fact about codegen, stronger than the `if constexpr`
source-identity argument it corroborates.
* K=9's SHARED is exactly 64 B below K=8's at BOTH arches, which is the group
table falling from 16 entries to 8 at 8 B each. The compiler arrived at the
predicted number independently.
⚠ NEW LANDMINE, and the reason this is not just a green checkmark: K=9/LUT
lands on REG:64 at both arches — exactly the cap `__launch_bounds__(256, 4)`
imposes (65,536/4/256). It did not spill, so 4 blocks/SM holds, but it has ZERO
headroom and is the tightest of the six. The next register added to its hot
path is a spill, silently. Anything touching the K=9 expansion, the radix loop
or the packer must re-run the gate and read the number.
Also corrected while here: the escape-hatch note claimed dropping
QB_MIN_BLOCKS_PER_SM to 3 "reproduces the measured allocation exactly and is a
no-op". It does not — a looser cap (85 registers) leaves nvcc free to take more
than the 59-64 measured at 4 blocks. It is still the right escape hatch; it is
not allocation-preserving.
The kernel source is unchanged since the measured run — comments only, verified
by a comment-stripped diff — so the table describes this tree.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
…suspend the census absolutes Two corrections from the coordinator, both of which change what this branch asserts. 1. OOB READ, not a slow path. The serving side now PADS the packed row stride to 4 bytes and has DELETED its tail clamp, so a decoder reads a row in 4-byte units. `packed_bytes_per_row` returns `ceil(n*K/8)`, which at K=9 is `9*(n/8)` — a multiple of 4 only when `n % 32 == 0`, not `n % 8 == 0`. The whole-bytes-only guard would therefore have let the encoder emit a stride the decoder reads past the end of. The guard now also requires a 4-byte stride, but ONLY for straddling rungs (`K % 8 != 0`). K=4 and K=8 divide 8 and keep exactly the acceptance they ship with — tightening the rung that baked the published artifact to protect a rung that does not exist yet would trade a regression for a hypothetical. Refusing rather than padding is the point: where the bake is allowed to run, padded stride and unpadded stride are the SAME number, so the encoder cannot disagree with the decoder whichever of the two definitions serving is on. Adopting the padding would mean guessing a layout owned by another change, and a wrong guess is the OOB read this is avoiding. Every real V4-Flash shape passes untouched: 1792 and 512 symbols give 2016 B and 576 B at K=9, both already 4-aligned. Proven red: relaxing the new clause back to `is_multiple_of(8)` fails on `!packed_row_is_whole_bytes(8, 9)` — a 9-byte row, whole bytes but not 4-aligned, which is exactly the case that used to slip through. 2. THE CENSUS ABSOLUTES ARE SUSPENDED, so this branch stops quoting them. The `4.375 vs 15.125` / `3.46x` decode inst/weight figures came from the SASS census's differential-over-unroll-depth method, whose K=8 replay control did not reproduce: `#pragma unroll` is a hint, nvcc re-rolled the loops, all three unroll depths returned identical counts, and the linearity guard reads |d2-d1|/d1 — which is 0.00% precisely when all three are equal. It passed hardest when the rig was most broken. Struck from both kernel headers, with the reason recorded so nobody reinstates them from git history. The clamp DELTA (+1.126 K=9, +1.125 K=10) survives, being a difference between two kernels in the same mode. The structural claim — one symbol carries four weights, so fewer instructions per weight — does not depend on the census; the ratio did. The s/layer derivation in `qtip_beam.cu` does NOT inherit the suspension: it divides from a measured A100 bake wall-clock (FACTS.md:990), and the ~3.87 radix passes it references are a runtime histogram from `probe_beam_kernel_cost_drivers`, not a census differential. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
Caught only after the previous push, because the verification chain piped clippy into `tail` — which makes the pipeline's exit status `tail`'s, so a red clippy read as green and the `&&` ran on anyway. Verified here with real exit codes: clippy rc=0, test rc=0, host gate rc=0. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
Retargeted at the integration branchBase changed: The queue is being restructured to the shape the owner asked for: one PR open against Two things had to land on
This PR was not closed and is not considered stale. An audit of the queue found the overwhelming majority of it to be real work that was never merged, not noise. What you need to do: rebase onto |
On the red
|
Stacked on #169 (base
wave-encoder-k8l12). Merge #169 first.Does #169 already bake K=9? No.
QtipGeomrefuses it in twostatic_asserts —K_ <= 8and8u % K_ == 0u("K must divide 8: the packer emits whole bytes") — and behindthem
SYMS_PER_BYTE = 8/K == 0, soqtip_packed_bytes_per_rowdivides by zeroand
qtip_pack_symbolcannot address a straddling symbol. The Rust side has thesame hole:
cuda_ops.rscomputed8 / super::Kand would have panicked beforeany kernel launched. #169's title says "ready to bake K=8/V=4/L=12" and that is
exactly what it is.
What #169 already got right, and why it is also the cost argument
Working the K=9/L=12 case through #169's own rule: groups
2^(L-K) = 8,alphabet
2^K = 512,LANES = min(512, 256/8) = 32, soCAND = 512/32 = 16— the same register array as K=4 (16/1) and K=8(256/16). The
cand[]spill landmine does not fire. Nothing else per-timestepmoves either:
n_cand = ng · 2^Kcand[]num_symbolsk_in/2k_in/4k_in/4The candidate count is
2^L-saturated at every geometry, so(n/V)·W·2^Kwasnever the shape of this kernel's cost — which is the same reason W=256 and W=32
measure within ~1% (
FACTS.md:1320).What changed
qtip_geom.cuh— the packed row is a bitstream: symboltat bits[t·K, t·K+K), LSB-first, rowceil(num_symbols·K/8)bytes. That is theserving side's
Rung::{pack, extract, packed_bytes}(feat(qtip): V=4/L=12 decode family — K as a parametric seam, bit-exact CPU parity #168), not a secondconvention: the UQFF geometry section is
[tag=3, K, L, V], so K=8 → K=9 isone byte of artifact change.
ceil((8 − gcd(K,8) + K)/8), notceil((7+K)/8)— thereachable offsets are multiples of
gcd(K,8), so 7 is unreachable at K=10 andtwo bytes are exact at both K=9 and K=10. The packer writes only the bytes a
symbol's bits fall in, which keeps the last symbol of a row in bounds.
qtip_beam.cu—QtipGeomK9V4L12added to the dispatch table. K=10 isdeliberately not instantiated: a geometry that has not been through
-res-usageis a claim, not a rung.cuda_ops.rs/mod.rs— onepacked_bytes_per_rowformula, plus awhole-byte refusal that is still exactly "num_symbols must be even" at K=4.
qtip_beam_res_usage_check.sh— asserts six beam kernels, not four. ALOCAL:0verdict has to be about a build that still contains the K=9instantiation.
K=4/V=2/L=16 stays byte-identical
Not by assertion. The byte-aligned expression is kept verbatim under
if constexpr, so K=4 and K=8 compile from unchanged source rather than from ageneral form believed to fold back to it. Five host tests pin the rule the
discarded branch would produce, and a
constblock asserts at compile timethat the general formula reproduces
num_symbols/2at K=4 andnum_symbolsatK=8 at the real V4 shapes (7168 and 2048 in-features).
All five proven red: M1 (
packed_bytes_per_rowassumes one byte/symbol)fails to compile on the const assert; M2 (one-byte extract) reds only the
K=9 test; M3 (whole-byte guard always true) reds only the partial-byte test;
M4 (naive
ceil((7+K)/8)) reds only the straddle test at K=10; M5(nibble order flipped) reds only the K=4 identity test.
cuda_beam_matches_cpu_beam_bit_for_bit/cuda_beam_unpruned_matches_cuda_exhaustiveare#[cfg(feature = "cuda")]andskip without a device; they were not run here (no GPU). They remain the
hardware gate for this change.