Conversation
Code Metrics Report━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Language Files Lines Code Comments Blanks ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ C Header 5 305 210 52 43 CSS 2 1181 1036 34 111 CUDA 81 29475 20161 6275 3039 Dockerfile 1 39 22 8 9 JavaScript 16 3546 2676 482 388 Jinja2 7 694 656 5 33 JSON 75 4896 4893 0 3 Makefile 1 6 5 0 1 Metal Shading Lan| 33 12224 9431 1142 1651 PowerShell 1 300 227 30 43 Python 147 15371 12677 824 1870 Shell 42 10304 6858 2769 677 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 213 48932 0 38081 10851 |- BASH 72 1655 1203 331 121 |- C 4 19 19 0 0 |- CUDA 2 84 56 16 12 |- JSON 19 779 779 0 0 |- PowerShell 1 1 1 0 0 |- Python 23 1008 787 113 108 |- Rust 68 2063 1727 78 258 |- TOML 6 207 164 0 43 |- YAML 5 41 36 5 0 (Total) 54789 4772 38624 11393 ───────────────────────────────────────────────────────────────────────────────── Rust 681 340988 293252 18096 29640 |- Markdown 504 32562 471 28063 4028 (Total) 373550 293723 46159 33668 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1353 516771 363188 99077 54506 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
Contributor
Author
Batch sweep on landed master
|
| B | per-req tok/s | ms/tok | agg derived | agg measured | spread | MEM-CTRL % (med/max) | conc | VRAM free |
|---|---|---|---|---|---|---|---|---|
| 1 | 36.68 | 27.3 | 36.68 | 36.39 | 0.8% | 10 / 10 | 1 | 51,748 |
| 8 | 6.45 | 155.0 | 51.62 | 43.70 | 15.3% | 3 / 4 | 7 | 51,748 |
| 32 | 2.43 | 411.6 | 77.76 | 62.71 | 19.3% | 3 / 12 | 28 | 51,748 |
| 64 | 1.61 | 622.0 | 102.89 | 68.63 | 33.3% | 3 / 21 | 53 | 51,140 |
| 128 | 0.82 | 1213.1 | 105.52 | 72.63 | 31.2% | 3 / 38 | 106 | 44,804 |
| 256 | 0.36 | 2802.3 | 91.35 | 79.86 | 12.6% | 3 / 70 | 216 | 38,116 |
| 512 | 0.18 | 5471.2 | 93.58 | 87.32 | 6.7% | 10 / 100 | 478 | 15,876 |
There is no knee
Aggregate rises 36.39 → 87.32 tok/s (×2.40) for 512× the users. Per-request
retention falls to 0.5%. The curve is flat-then-flatter; the amortisation regime
never starts anywhere up to B=512.
The memory controller is the finding
Median memory-controller utilisation is 3% at B=8 through B=256, and 10% at
B=512 — while GPU utilisation reads 99–100%. It does not climb with batch. At
b=1 it is 10%, and adding users makes it lower, not higher.
Per the stated criterion — "if it stays near 4% at B=256, batching is structurally
broken" — it stays at 3%. A GPU at 100% utilisation with a 3%-busy memory
controller is saturated with small kernels, not with weight traffic. 256 or 512
concurrent users are not sharing the 74 GB read; the engine is re-doing per-token
work that batching should have amortised.
What it is not
- Not an OOM or a capacity cap. 15.9 GB still free at B=512; no row failed.
- Not scheduler admission. Achieved concurrency tracks B (478/512).
- Not the kernels being individually slow. They are individually tiny.
derived vs measured diverges most at B=64/128 (33%, 31%) precisely where
achieved concurrency lags B (53/64, 106/128) — which is why the concurrency
column has to be reported next to any aggregate number.
Against the gate
The target is 14,000 tok/s aggregate. Measured 87.32 at B=512 — 0.6% of the
gate, ~160× short, and ~0.3% of the ~33,000 tok/s roofline at that batch. On
this evidence 14K is not reachable by kernel tuning; it needs the expert path to
actually batch.
Harness: arc-tools/batch_sweep.py (this PR). Zero-token rows fail loudly and are
never averaged in — no number here came from the requested-token formula.
Contributor
Author
The missing knee has a named cause: the #167 gate never selects the grouped GEMM in decodeLanded master 1. The gate's arithmetic makes the grouped path unreachable in decode
Decode never has 683 tokens in a step. 2. Proven from the mangled symbol, not from re-reading the condition
The gather kernel is 57.6% of all GPU time and 25× more launches. The grouped kernel 3. Forcing the grouped path is worth up to 1.66×, free
Per-request improves too: B=512 goes 0.18 → 0.34 tok/s (1.9×). The b=1 and B=32 rows show That is a gate condition, and it is the cheapest 1.5× on the board. 4. But the memory controller still does not climb — so this is not the whole answerMedian memory-controller utilisation is 1–2% on the forced-grouped arm, versus 3% on Per the stated decision rule: the grouped path is not grouping either. Fixing the gate Two named causes now, in priority order:
Harness: |
This was referenced Aug 21, 2026
Merged
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…hould permit (#216) Seven open PRs are red for their ADDRESSING, not their quality. The `base-branch` lane hard-fails any PR whose base is not literally `master`, and the PR queue is being restructured so that agents branch off a single integration branch, `release/openrouter-ready`, and PR back into it. The lane's intent is kept intact. It was never really asking "is the base literally master" — it was asking "is the base a stable branch that will itself land on master through a PR this lane has gated". `master` satisfies that. So does exactly one other branch, by construction: `release/openrouter-ready` reaches master through exactly one pull request (#194), whose own base IS `master`, so it takes the `master` arm of this same check and must be green on a tree whose parent is proven master. Master's protection is therefore unchanged: nothing enters master without a full run whose base is master, and `CI complete` still aggregates this lane. The original incident was PRs based on another PR's TRANSIENT branch — a base that can be rebased, force-pushed or abandoned underneath them, so their green described a tree that might never exist. An integration branch is the opposite: long-lived, the tree that will exist, and re-proven against master by #194 before any of it ships. Implementation notes: - Exact-match allowlist via `case`, not a `release/*` glob. A prefix match would let any newly pushed `release/anything` declare itself a landing branch, which is the same unproven-base hole under a new name. - `BASE_REF` still arrives via `env:`, never interpolated into the script. - The job name (`Base branch`) is unchanged, so branch protection and the `ci-complete` needs list keep matching. Verified: the allowlist accepts `master` and `release/openrouter-ready` and refuses `main`, `release/something-else`, `release/openrouter-ready-evil`, `xrelease/openrouter-ready`, `MASTER`, `master*`, `master; echo pwned`, the empty base, and every current PR base branch in the queue. Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
One branch, one PR, one merge decision. Merging it means Arc can serve and sell tokens on OpenRouter. 16 items across blocking / product / trust / commercial, each with its measured provenance. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Jish, direct: 1,400 aggregate is break-even and therefore pointless. 14,000 is $50/GPU-hr against a $4.92 card - 10x margin, a business. And 'OpenRouter-releasable' means Arc as we envisioned it, not a defect list. Rewritten accordingly: batching that amortises, the engine near the silicon, prefill as a product, the memory system, the multipliers we built and never switched on, and instruments we can trust. Key arithmetic: 14K is 84% of roofline at B=256 but only 42% at B=512, because memory traffic per step is flat as batch grows. 14K is a batching problem, not a kernel-speed problem - and batching is broken precisely because we are instruction-bound at 4% memory utilisation. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…-controller Measures what a capacity claim actually rests on: fire B concurrent streaming requests and time the inter-token intervals of each. Two aggregate anchors that must be read together, because they disagree exactly when the batch never formed: * derived = per_request * B -- assumes all B ran at the median rate * measured = cohort tokens / cohort span -- what the box actually delivered Achieved concurrency is reported beside them, integrated over the span. At B=64 the scheduler only ever held ~53 streams, so `derived` overstates by 33%; without the concurrency column that gap is invisible and the higher number is the one that gets quoted. Samples `utilization.memory` DURING the decode, not after. That is the memory-controller busy percentage, and it is the diagnostic that separates "batching is amortising the weight reads" from "the GPU looks busy running tiny kernels". A sampler started after the run writes an empty file and looks green, so this one starts before the requests and stamps every sample. A row where any request returns zero tokens is printed as FAILED and never averaged into the table. Measured on landed master d774267, H200, V4-Flash qtip2 (see PR discussion): aggregate rises only 36.4 -> 87.3 tok/s from B=1 to B=512 (x2.40 for 512x the users) while the memory controller sits at 3% median through B=256. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
`Qtip2bLayer::gather_forward` gated the trellis grouped GEMM on tile occupancy: switch only once each woken expert draws `GROUPED_TILE_M` (=16) pairs. At V4's top-6-of-256 routing that needs `6n/256 >= 16`, i.e. **n >= 683** — a token count decode never reaches. The clause therefore pinned every decode step, and every batch up to 682, on the per-pair gather-GEMV arm. A profile of the shipped default shows the consequence directly: `qtip_gather_gemv_warp_kernel` at 57.6% of GPU time with 25x the launches of the grouped kernel. Two independent measurements retire it. Decode on an H200, clean rows only (>=97% achieved concurrency, <7% derived-vs-measured spread): grouped wins at B=48 1.10x, 64 1.14x, 80 1.22x, 96 1.29x, 128 1.66x, 512 1.51x. Prefill over prompt length: +16% at 16 tokens, +27% at 32, +50% at 64, +78% at 128, +181% at 512. The gate's own justifying A/B -- "1.00x at N=128 and N=512" -- does not reproduce. **Tile fill was the wrong quantity.** At n=16 the tile is 7.5% full and the grouped kernel still wins by 16%, so a predicate demanding a full tile cannot be the gate. The win is that the grouped kernel stages a woken expert's bytes once per m-tile instead of once per (token, expert) pair -- the GEMV arm's reads are exactly linear in n_tokens, 12.00x redundant at n=512, while the grouped arm holds flat at one pass over the expert set -- and that it runs on tensor cores. Occupancy was never the mechanism being traded against. `DECODE_REGIME_MAX_TOKENS` survives as a floor: below it there is nothing to group (b=1 measures 0.47x) and it is the RUN-161 capture-safety rule. 9..=15 is unmeasured, bounded by that floor and by the +16% at 16, and rides with the grouped arm. Fast path by default; `ARC_NO_QTIP_GROUPED_MOE` remains the kill switch and `ARC_QTIP_ONDEVICE_MOE_MAX_TOKENS` still pins the GEMV arm for A/B. The decision moves into `prefer_gather_gemv` so it can be tested: the dispatch site is inside `#[cfg(feature = "cuda")]`, and a non-CUDA `cargo test` reports a full green over code it never compiled -- which it did to me once during this change. The tests pin every measured batch size to the grouped arm and carry a control asserting the retired clause would have excluded all of them. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…rung measured The previous commit removed the tile-fill conjunct from the bitshift rung (`Qtip2bLayer`, --isq qtip2b). The LUT rung (`QtipLayer`, --isq qtip2) carries the identical clause at mod.rs:3652, and the published V4-Flash artifact this was measured against is **qtip2**, not qtip2b: the profile's kernels are `qtip_lut_grouped_gemm_kernel` and `qtip_gather_gemv_warp_kernel`. So every decode number in the sweep — GEMV 78.46 vs grouped 89.37 at B=64, 72.63 vs 120.36 at B=128 — came from this rung, and fixing only the bitshift rung changed nothing on the box: the default still read 81.17 at B=128 after that build. Fixing the rung the measurement actually exercised. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
heydryft
force-pushed
the
release/openrouter-ready
branch
from
August 21, 2026 15:36
17597e6 to
d665956
Compare
heydryft
changed the title
This was referenced Aug 21, 2026
Open
Open
…#170) * feat(qtip): UQFF geometry discriminator — K8/V4/L12 artifacts declare themselves Stage 2 of the K=8/V=4/L=12 rung: an artifact can now say which trellis geometry it was baked at, the loader reads it, and a build that cannot decode it refuses instead of guessing. WHY THIS NEEDS A DISCRIMINATOR AT ALL Both geometries are 2 bits per weight (`bpw = K/V`), so a K=8/V=4 row and a K=4/V=2 row of the same `in_features` occupy the SAME number of packed bytes. A mislabelled artifact therefore indexes in bounds at every symbol and returns plausible garbage rather than faulting. The table is the only tensor that differs — `[4096, 4]` BF16 vs `[65536, 2]` F32 — which is why `validate_shapes` checks the table's dtype AND its length, and why both checks are load-bearing rather than belt-and-braces. WIRE FORMAT A trailing section `[tag=3, K, L, V]`, written ONLY for a non-default geometry and written BEFORE the codebook section. Both of those are deliberate: * Default writes nothing, so every artifact Arc has already produced stays byte-identical and no checksum moves. Pinned as an exact suffix, not as a vague "unchanged": serializing the same layer with the tag flipped appends exactly [3, 8, 12, 4] and nothing else. * BEFORE the codebook, because a build that predates this field parses the trailing region by handing its first byte to `QtipCodebook::from_wire`, which refuses every tag it does not know. Tag 3 first means an old Arc fails closed. Tag 3 last would let it consume the codebook section, stop, and decode K=8/V=4 symbols as K=4/V=2 without faulting. The trailing region is now a section loop; a repeated tag is refused rather than letting the last one win. `QtipGeometry` has no per-site default: adding the field broke all 11 construction sites at compile time and each one now states its geometry, the same discipline `search` and `search_detail` already follow. `from_stacked_parts` takes it explicitly and validates it. `stack_experts` and the 3-D quantize path refuse mixed-geometry stacks — only one table survives a stack, so at most one expert could decode. The computed `sum2` codebook is refused at V=4 on both the write and the read side: it produces a PAIR of values per state and has no V=4 form. Guards were mutation-tested (F1-F9). Three passed on broken code and are now covered: `from_stacked_parts` validated nothing, the table-dtype half of the discriminator was dead (the element-count check masked it), and `stack_experts`' mixed-geometry refusal was untested. `assert_tensor_bits_eq` grew a BF16 arm — it panicked on the new table rather than comparing it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H * fix(qtip): refuse an unsupported geometry at every decode entry point Stage 2 made a K=8/V=4/L=12 artifact loadable. On its own that is a half-applied change: every decoder behind `forward`, `gather_forward`, `dequantize_weights`, `dequantize_expert` and `qtip_packed` unpacks nibbles and indexes a [2^16, 2] table, and handed a K=8 row none of them would fault — at 2 bits per weight the packed byte count is identical and every index lands in bounds. They would serve garbage. So each entry point now states the geometries it implements. `qtip_packed` returns None rather than a `QtipPackedView`, which carries no geometry field and would therefore be read as K=4/V=2 nibbles by any consumer. The refusal names the entry point, and the test asserts that name. An entry point that merely inherits a callee's guard is not guarded: mutation G1 removed `forward`'s own check and every test stayed green, because the error still arrived from `dequantize_weights` further down. All six mutations (G1-G6) are red now. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H * fix(qtip): UQFF validates the PADDED row stride, not the data length Follows the padding decision one branch down. `QtipGeometry::packed_len` now means the allocated stride (what a tensor is shaped at and what this format validates) and `data_bytes` is the bit-rate-governed part, both delegating to `trellis_v4l12::Rung` so there is still one implementation of each. Three tests changed premise rather than value and were rewritten, not patched: the bit-rate ratio claim moved to `data_bytes` where it is exact, and the stride got its own bounds (never smaller than the data, never more than 4 bytes larger). A K=9 row of `in_features=36` is 11 data bytes in a 12-byte stride. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H * fix(qtip): the CUDA bake fast path lost the geometry wire tag `quantize_with_options_cuda` builds a QtipLayer directly and was the one initializer the geometry discriminator missed. It is behind `#[cfg(feature = "cuda")]`, so the default Check jobs compiled fine and only the CUDA lane went red (E0063 at qtip/mod.rs:2318). Set the same K4V2L16 tag its CPU sibling sets — the CUDA bake kernels run the identical K=4/V=2/L=16 trellis, so a GPU-baked artifact must carry the identical tag or it would deserialize as a different geometry. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…lackwell (#94) * feat(turboquant): instantiate the CUDA kernels at head_dim 64/128/256/512 TurboQuant's paged kernels were hardcoded to head_dim 128 — `constexpr int HS = 128`, a single `SGN[128]` sign table, `rotate128`, and six `extern "C"` launchers each opening with `if (hs != 128) return;`. That early return is a silent no-op: the output buffer is left uninitialized, so a mismatched head dim produces garbage *fast* rather than failing. Kernels - `tq_attn`, `tq_attn_clocked`, `tq_cache_k`, `tq_cache_v` and the two clocked cache variants now take HS as a template parameter. NT stays 128 for every HS, so the warp-reduction topology is unchanged; wider heads give each thread DPT = HS/NT output dims instead of adding threads. At HS=128 the emitted work is the same as before. - The per-lane packed-K gather is now one vector load (`tq_load_klane`) instead of BPL scalar byte loads. The x=16 cache swizzle keeps each lane's run contiguous and naturally aligned, which `packed_k_bytes_are_swizzle_aligned` asserts. - Every `if (hs != 128) return;` is replaced by a dispatch over the instantiated widths. D16 dual-arch - `TQ_PREFETCH` software-pipelines the K gather on SM90/SM100/SM103 and runs flat below Hopper — same arithmetic, scheduled for the part it runs on. - Launches now opt in to the arch's real dynamic shared-memory budget. The logits array scales with context length and CUDA caps dynamic shared memory at 48 KB without an explicit opt-in, so contexts past ~12k tokens could not launch at all on any arch. - `ARC_CUDA_ARCHS=90,100,103` emits cubins for Hopper *and* Blackwell rather than only the build box's GPU. Tables - Codebooks and sign diagonals are dimension-dependent, so 512 needs its own. All 20 tables are emitted by `emit_turboquant_cuda_tables`, which calls the same `generate_signs`/`get_codebook` the Rust compressor uses — correct by construction, not by transcription. The generator reproduces the shipped d=128 table byte-for-byte. - `turboquant::cuda_tables` re-derives every table at test time and diffs the checked-in header, and pins the Rust head-dim list against both the kernel's dispatch switch and the FFI wrapper's list. Nothing previously asserted that the tree's four copies of the sign table agreed with the generator or with each other. Gates - `TURBOQUANT_HEAD_DIM = 128` exact-match becomes membership in `TURBOQUANT_CUDA_HEAD_DIMS`; K and V are compressed independently and no longer have to be equal widths. Not measured on a GPU: the kernels are written and gated but this box has no nvcc, so nothing here is a claim about hardware behaviour. * test(gpu): wave63 gate — dual-arch compile proof for the head_dim 512 kernels Proves the four head-dim instantiations compile for sm_90a AND sm_100a/sm_103a (D16), not just the build box's GPU, and that the shipping decode path did not regress. Deliberately contains no quality A/B: TurboQuant quality is settled by prior measurement and re-proving it would spend GPU budget on a known number. * fix(build): watch turbo_paged_attention.cu for changes Every other .cu in this crate had a rerun-if-changed line; this one did not. Emitting any rerun-if-changed disables cargo's watch-the-whole-package default, so edits to the largest CUDA file in the crate were relying entirely on cudaforge's object cache to trigger a rebuild. * fix(turboquant): 64-bit cache offsets, per-device smem opt-in, clocked-kernel parity Three problems from review of the head_dim generalization. 1. 32-bit overflow on cache byte offsets. `pb*kbs`, `pb*vbs`, `bi*cbs` and `bi*vbs` were all `int`. At head_dim 512 with nkvh=8 and block_size=32, `kbs` is 65536 B, so `pb*kbs` wraps once the K cache passes 2 GB — an entirely reachable allocation on an H200, and 4x sooner than at 128. Now computed in `long long` at all eight sites. 2. The shared-memory opt-in cached a process-wide `granted` flag per kernel instantiation. The requestable maximum differs by architecture, so one flag cannot be right across a mixed fleet, and a failed request left `granted` at 0 and fell through to a launch that then failed with no diagnostic. The cache is gone; the call is cheap and now always attempted. 3. `tq_attn_clocked` kept the scalar byte gather while `tq_attn` moved to a vector load plus the arch-gated pipeline. Identical results, but the phase-2 stamps would have described a kernel nobody launches — at head_dim 512 that is 8 scalar loads measured against the 1 vector load actually executed. The loop now mirrors `tq_attn`. Verified without nvcc by type-checking the kernel bodies against a CUDA shim (clang++ -Wall -Wextra -Wshadow, both __CUDA_ARCH__ branches): clean. Still not run on a GPU. * fix(gate): D18 exit codes, branch-independent preflight, source-landmine guard The gate exited 1 when gpu_box_preflight.sh was absent. That read as a code failure for a couple of minutes when it was really 'this machine cannot answer'. Per D18 rule 2: 0 pass, 1 genuine failure (the only signal to act on), 2 environment could not answer. Nine conditions now exit 2 — missing/failing/incomplete preflight, no repo, no nvcc, unreachable remote, un-checkout-able branch, no archive to inspect, no cuobjdump. Four exit 1, all genuine: nvcc compile failure, a native build failure, a test failure, and the archive missing an arch it claims to carry. The preflight is now looked up at ARC_PREFLIGHT, then /root/arc-tools, then /usr/local/lib/arc, and only last inside the repo. It lives on wave61/box-preflight-shared-prefix, so a repo-relative path vanishes the moment a runner checks out any other branch — which is exactly how this gate came to refuse to run. Guarded the sourced-preflight landmine directly. gpu_box_preflight.sh ends on [ "$_arc_pf_sourced" = "1" ] || exit 1 so when SOURCED on a *failing* box that test succeeds, short-circuits the '|| exit 1', and becomes the last command — 'source pf || handler' returns 0 and the handler never fires. The gate ignores the return value entirely and reads _ARC_PF_FAILED, treating an unset flag as 'did not reach a verdict' rather than as a pass. Also reordered so the sm_90a+sm_100a+sm_103a compile is STEP 1 and needs no full workspace build, and it now asserts on cuobjdump output that both cubins are actually in the archive — a build that merely succeeds is not the claim. Exit paths verified locally against a stub reproducing the real preflight's sourced-exit structure: all four environment cases return 2. * fix(paged): keep the TurboQuant DEFAULT at head_dim 128; widen only what is asked for Widening the acceptance gate from head_dim == 128 to every instantiated width was correct for what a user may REQUEST, but PagedCacheType:: TurboQuant carries #[default] (cache_engine.rs:22), so it also silently moved standard-layout models at head_dim 64/256/512 onto kernels that have never executed — and silently dropped prefix caching with them, since supports_prefix_cache() is false for every TurboQuant variant. Two regressions, neither visible at the call site, for users who never asked for TurboQuant at all. This repo has paid for that exact shape once: FP8 KV shipped default-on and unmeasured (wave43-BU) and every V4 request died. #98 cites that incident by name as its own reason not to default itself on; this takes the same choice. resolve_for_model already distinguished explicit from ambient-default in order to pick error-vs-fallback, so the fix is to let that same 'forced' flag also pick WHICH head-dim set applies: compiled -> widens what you may ask for (TURBOQUANT_CUDA_HEAD_DIMS) measured -> widens what you get unasked (TURBOQUANT_DEFAULT_HEAD_DIMS) TURBOQUANT_DEFAULT_HEAD_DIMS is [128] and moves when a hardware gate passes, not when a kernel compiles. Nothing about the kernels changes: the six 'if (hs != 128) return;' no-ops (D18 instance 3) and the 48 KB dynamic smem cap fix are untouched, and explicit --pa-cache-type turboquant still reaches every instantiated width including V4's 512. The unsupported-geometry message now separates 'no kernel exists' from 'a kernel exists but is unmeasured, so the default will not choose it for you' — the second tells the user to opt in rather than to give up. Tests: the old fixture asserted the regression (its expected-to-fall-back list moved [64,96,192,256] -> [96,192,320,1024] precisely because 64 and 256 stopped falling back). Replaced with an explicit-path test over every instantiated width, plus two pins on the default path. Both pins were mutation-checked by reverting the const to the full set: both fail, and the first revision of one passed vacuously because widening the const emptied its loop, so it now asserts the loop is non-empty before iterating. * docs: correct the public record for the widened TurboQuant kernels This PR branched before #101 ("correct the public record") landed, so it had never seen that text. Merging master in makes four README claims false — three falsified by this PR's own kernels, one that was already wrong when it was written. Falsified by this PR: - L31 "the paged kernel exists at head_dim 128 only" - L31 "there is no kernel at head_dim 512, so DeepSeek V4 cannot use it" DeepSeek V4 reports KvCacheLayout::Standard at head_dim 512 (normal_loaders.rs), so the 512 instantiation genuinely reaches it. - L141 an explicit `--pa-cache-type turboquant` off-128 "is a hard error" It is now accepted at any instantiated width; the hard error moved out to the uninstantiated set. Already wrong before this PR: - L141 "TurboQuant is not the default anywhere" `defaults::PAGED_CACHE_TYPE` is `PagedCacheType::TurboQuant` and the CLI's `--pa-cache-type` has no clap default, so leaving it unset on CUDA gives a standard-layout head_dim-128 model TurboQuant KV with no flag — and silently drops prefix caching, which no TurboQuant variant supports. That was true on master; the narrowing in a3a5faa bounds it to 128 but does not remove it. Stating the opposite understated the risk to anyone reading the README to decide whether to opt out, so it is corrected in place rather than quietly reworded, and the Rust example's comment now says `Auto` opts *out* rather than implying TurboQuant is off until asked for. Also refreshes the `--pa-cache-type` help text, which still described the explicit path as 128-or-error. No behaviour change: documentation and one clap doc comment. * fix(build): stop emitting a duplicate -gencode for the build box's own arch `mistralrs-paged-attn/build.rs` appended one `-gencode` per entry in `ARC_CUDA_ARCHS` on top of the one cudaforge derives itself, and applied the same `a` suffix for cap >= 90 that cudaforge's `GpuArch::auto_suffix` does. Verified against cudaforge 0.1.5 rather than assumed: compute_cap.rs auto_suffix(90).to_gencode_arg() -> "-gencode=arch=compute_90a,code=sm_90a" builder.rs:401 let gencode_arg = gpu_arch.to_gencode_arg(); builder.rs:405 command.arg(&gencode_arg)... builder.rs:412 for arg in &self.extra_args { command.arg(arg); } cudaforge emits its own gencode first and then appends `extra_args` verbatim, so on an H200 building `ARC_CUDA_ARCHS=90,100,103` nvcc received `-gencode=arch=compute_90a,code=sm_90a` twice, byte-identical. Fixed by hoisting the existing `get_compute_cap()` binding above the loop (it is what decides which requested arches are genuinely extra) and skipping the one cudaforge already covers. The later duplicate binding is deleted; there is now one. This is the dedupe half of #108's `arc_target::build::split_primary()`, done locally because `arc-target` does not exist on this branch. The other half — `verify_and_export()`, the `cuobjdump` gate that fails a build whose archive is missing a requested arch — genuinely needs that crate and is left to #108, which should replace this block with the two calls. Until it lands, `arc-tools/wave64_v4_turboquant_kv_gate.sh` STEP 2 asserts the arch list externally with `cuobjdump --list-elf`, so a missing arch is caught by the gate even though it is not yet caught by the build. Not verified here: whether nvcc rejects or tolerates the duplicate flag. macOS cannot run the cuda build path. Emitting it once is correct either way, which is why this is worth doing without waiting on that answer.
* feat(qtip): Stage 3 — dispatch and selection for K=9, default OFF
A K=9/V=4/L=12 layer is now SERVED rather than refused, and there is a way to
ask for the rung. K=8 stays exactly as it was: the compiled control, nothing
spent on serving it.
WIRED
* CPU decode for the whole V=4/L=12 family, routed through
`trellis_v4l12::Rung` — which already owns the extraction arithmetic and is
bit-exactly gated against the GPU kernel. Not restated here: a format with
two implementations has none.
* `dequantize_weights` / `dequantize_w` / `forward` (2-D) for the family.
* `forward_v4l12_gemv_cuda`: the family's single GPU path, the single-token
fused decode+gemv the kernel exists for. Row-scale hoist OFF — it is the
one lever that costs bit-exactness against the CPU reference, and there is
no device measurement to trade that for.
STILL REFUSED, LOUDLY
Every path with no V=4 kernel: 3-D `gather_forward`, `dequantize_expert`,
3-D `dequantize_weights`, and `qtip_packed` (whose `QtipPackedView` carries no
geometry field, so a consumer would read K=9 bytes as K=4 nibbles). A
half-wired dispatch that falls through is worse than one that refuses — at
2 bpw a K=4 misread does not fault, it just serves wrong weights.
SELECTION
`ARC_QTIP_GEOMETRY` = `k4v2l16` | `k<K>v4l12`, mirroring `ARC_QTIP_CODEBOOK`.
Unknown values are refused, never resolved.
DEFAULT OFF, and this is the load-bearing line: `QtipGeometry::DEFAULT` stays
`K4V2L16`. Nothing in this family has touched a GPU — the kernel has never been
compiled at any K — and the precedent is specific: the last KV-storage change
that shipped default-on with only CPU validation killed every request on the
first V4 forward that met a real device. Its arithmetic was proven and its
device-time cost was not. Ship the path, leave the default alone.
Guards shown red (S1-S7). THREE HOLES, all mine:
* S2/S5 sit behind `#[cfg(feature = "cuda")]`, which no test lane compiles,
so both mutations passed. Now pinned over the SOURCE — a guard that only
exists on hardware nobody has is not a guard.
* S7: the env test reimplemented the parse rule inline and asserted the
replica, so the real reader could silently resolve an unsupported K to the
default and stay green. `from_env` now delegates to a pure `from_spec` and
the test calls that.
I also deleted my own five new tests mid-edit — a second replacement span
swallowed what the first had just inserted, leaving two dangling doc citations
and a suite that looked green. Restored and verified present by name.
And the source guard fired on itself: `include_str!("mod.rs")` includes the
test module, so a negative assertion naming the string it forbids matched its
own text. Scoped to the serving half. That is the second time this class bit
me here; the note is in the code.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
* fix(qtip): Stage 3 CPU decode reads the padded row stride
The K=9 CPU decode sized its rows at the data length. Rows now carry tail
padding so the kernel's compile-time-width extraction needs no clamp, and that
padding is part of what the artifact stores — so a decoder that assumes the
data length disagrees with the loader about where row N starts.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H
* fix(ci): the CUDA dispatch source-guard died on Windows CRLF
`include_str!("mod.rs")` yields CRLF on a Windows checkout, so
`split_once("#[cfg(test)]\nmod tests {")` never matched and the guard hit
its `expect` instead of asserting anything — the guard was not failing,
it was not running. Normalise line endings before the split; every
assertion after it already goes through the whitespace-squeezer.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…d refuse to instantiate on a miss (#181) Capture recorded and died on the FIRST cuGraphLaunch with 700 (CUDA_ERROR_ILLEGAL_ADDRESS), one line after: [alloc-cache] WARNING: MISS during capture, size 393216 bytes (not pre-warmed) That miss is the cause, not a bystander. An allocation that misses the warm pool while a capture is open is served by cuMemAllocAsync from the graph's PRIVATE pool and recorded as a graph memory node whose address the first launch has to back; it does not get backed, the launch faults, and the poisoned context surfaces later as glibc heap corruption. The allocator's 1 GiB cap is not involved — bounded and unbounded legs are byte-identical during capture (alloc/step 44.1, free/step 78.7, hit-rate 0.9954, held 315.4 MiB) and die on the same line. Why the warm pass missed it: `tail` in the rolling compressor is rebuilt every step at width `tokens - base`, and `base` jumps a whole `ratio` at a group boundary while `tokens` climbs by one, so the size cycles through `ratio` consecutive values (measured `4096 × {18,19,20,21}`). One warm pass warms one phase of that cycle. The captured step lands on another phase and asks for a size the pool has never held. Three changes, none of them a hardcoded size: 1. Warm until the profile stops growing. `ARC_GRAPH_DEFERRED_PASSES` defaults to 4 (was 1), and a pass that teaches the allocator a new size buys another, from a bounded budget (`ARC_GRAPH_DEFERRED_MAX`, 24). A cycle of any period is covered without anyone tuning a constant. 2. Top the pool up from the measured profile before capture opens — `prewarm_alloc_cache(slack)`, `ARC_GRAPH_PREWARM_SLACK` default 4 buffers per size. This covers the residual case the warm passes cannot: the captured step allocating MORE of a size than any warm step did. 3. Gate instantiation on zero capture-time misses. The allocator now keeps a counted, capture-scoped miss ledger (the old `missed` set is a log dedup — it silences a size's SECOND miss, which is exactly the one that matters). If any miss occurred while recording, the capture is cancelled and the run falls back to eager instead of instantiating a graph already known to fault. (3) is the part that holds even if (1) and (2) are incomplete: it converts a poisoned context and a process death into a named list of sizes. If those sizes turn out to grow monotonically rather than cycle, no amount of warmup is the answer and the next step is a shape-constant buffer or allocator size-class bucketing — the refusal message says so. Needs the candle fork's capture-miss ledger + profile-driven pre-warm (`prewarm_alloc_cache`, `reset_capture_misses`, `capture_misses`). This commit originally pinned Cargo.toml straight at 1fa534db, which was correct when it was written and is WRONG now: since then the integration branch moved to 89ab14ef, and the two revs have DIVERGED. 1fa534db is missing three commits 89ab14ef has — "perf(cuda): bound the caching allocator", its cap-contract test, and "fix(cuda): eviction must not run inside the capture window" — the first of which is what #213 (bound the alloc cache against real headroom) rests on. Pinning at 1fa534db would therefore have silently reverted the bounded allocator and a capture-correctness fix. Pinned instead at 9586979db (`arcgraph/bounded-alloc-plus-prewarm`), the merge of the two, which is ahead of 89ab14ef by 6 and of 1fa534db by 8 and behind neither — verified with `gh api repos/aeonmindai/candle/compare/...` and by reading all four required `pub fn`s out of `candle-core/src/cuda_backend/device.rs` at that rev. `cargo check --workspace` is green on the new pin; mistralrs-core's warning count is unchanged at 23. The re-indent of the same-step-probe block in normal.rs is the counterpart of the one 6fcb87f did on the integration branch: that commit de-indented it by 4 precisely because this commit's `if !misses.is_empty()` nesting level was not yet present. It is now, so the block goes back 4 deeper. Co-authored-by: Nirupam Bhowmick <support@runcrate.ai>
…t prices per-sequence advance — 4.46x at spread B=8 [MEASURED] (#121) * measure(sched+arcspec): the throughput ladder #116 unlocks, and the confound it has to be read against is worth anything: MTP was measured at 1.93 tok/step at one user collapsing to 1.06 at 128, and that collapse is what this stack targets. Two pieces, because the number is unreadable without the second. 1. `scheduler::bucket_telemetry` — a `SCHED[agg]` marker carrying buckets_per_step, running_bucket_size and offered_per_step, emitted on the SAME log fence as `MTP[agg]`. Both schedulers bucket the running set by exact cache length and run one bucket per step, preempting the rest. So a cell labelled B=128 can be a 3-wide step in the engine, and an aggregate number measured inside that is a measurement of the scheduler rather than of KV advance. Ragged admission exists precisely to collapse those buckets, so the two effects are confounded by construction — without these counters "aggregate did not move" is unattributable. running_bucket_size is the width that actually ran; offered_per_step is the width the harness thinks it asked for. 2. `arc-tools/arcspec_perseq_ladder.sh` — B = 1, 8, 32, 128, both arms, one server per arm, counters differenced across each cell's own wall-clock fence. It reports aggregate tok/s AND tok_per_step AND the decomposition (tok_per_batch_step x batch_steps/s, plus ms/step), because tok/step rising while aggregate falls has already happened on this chain and either number alone is unreadable. Prompt lengths are RAGGED by construction — 24..320 words cycled across workers — which is the deliberate difference from `arcspec_perseq_ab.sh`'s fixed 40. Uniform prompts hide the failure mode this stack addresses: the `xs` window defect could not even be reached with equal-length prompts. 320 keeps a >3x margin to the ~1,055-word serving cliff this branch does not carry the fix for. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * measure(arcspec): add the uniform arm — the only regime where the batch is whole The scheduler A/B came back while this was being built and it changes what the ladder can conclude. The bucketing law is now measured, not hypothesised: running bucket size = B / (distinct cache lengths) holding at 8/8=1 and 32/8=4, with `1 running, 7 waiting` sustained, and B=8 on spread lengths measuring 7.91 tok/s against B=1's 15.36 — batching is NEGATIVE on realistic traffic. So a spread-only ladder measures per-sequence KV advance inside a scheduler that is running one sequence at a time, and a flat aggregate there is unreadable: it cannot distinguish "the fix does not pay" from "the batch never existed". Both regimes are now run, per arm, mean-matched: spread 144 24 320 64 260 40 200 96 words (8 distinct lengths) uniform 144 words — the spread's mean AND its first element, so B=1 is byte-identical between regimes and the only thing that differs at B>1 is the spread itself Uniform is where a whole batch actually forms, so it is the only place this stack's effect on aggregate throughput is visible without the serialisation swamping it. Spread is still the regime the stack exists for. The report now prints the bucketing law's prediction beside the measured `running_bucket_size` per cell, and refuses to let a spread cell be read as a batch result when its running bucket is ~1. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(v4): stop reallocating the xs window every decode step — restore the bound the type already documents `tail` is described at the top of this file as "the raw rows behind `tokens`, bounded by `span_groups * ratio + margin` and independent of context length". That is an invariant the type has always claimed. The allocation did not honour it: `advance` rebuilt the buffer with `cat` + `narrow` every step at width exactly `tokens - base`, which is not stable — `tokens` climbs by one per token while `base` jumps a whole `ratio` at a group boundary — so the buffer cycled through `ratio` consecutive sizes and was reallocated on every decode step, on every one of V4's 41 compressed layers. This is not "add a pin". It is restoring a documented invariant that one code path broke. THE BOUND, WHICH IS THE REVIEWABLE PART keep_from = ((tokens - margin) / ratio + 1 - span_groups) * ratio base = max(keep_from, previous base), capped at tokens Write m = tokens - margin and q = floor(m / ratio). Then q * ratio > m - ratio, so keep_from = (q + 1 - span_groups) * ratio > m - span_groups * ratio, giving W = tokens - base <= tokens - keep_from < margin + span_groups * ratio so W <= span_groups * ratio + margin - 1, ALWAYS. Pinning to span_groups * ratio + margin is provably sufficient and provably never truncates. It is a function of the layer's geometry, not of context length: 24 for CSA (ratio 4, span 2, margin 16), 144 for HCA. That is the difference between this and a capacity that turns out to be a million-token context. Derived twice independently and agreeing, and checked a third way: the steady band is measured at [capacity - ratio, capacity - 1] = {20..23}. THE {18..21} vs {20..23} DISCREPANCY, SETTLED RATHER THAN ASSERTED ArcGraph measured 4096 x {18,19,20,21} from outside the engine; the retention rule predicts {20,21,22,23}. Both ratio-consecutive, both containing 21, offset by 2. The pre-committed reconciliation was that theirs is the pre-saturation ramp (base still at 0 while tokens climbs) and mine the steady state. `the_window_ramps_then_settles_to_ratio_consecutive_sizes` checks exactly that: the early band sits below the steady one, the steady one is `ratio` consecutive sizes at [cap-ratio, cap-1], and no width over 128 steps reaches `cap`. A bound that is right for the wrong reason is a trap for the next context length. WHY DEFAULT ON WITHOUT A THROUGHPUT NUMBER FIRST The measure-then-default-on rule exists because FP8 KV shipped on and changed VALUES. This changes only the size of an allocation: the compressor is handed a slice covering the same absolute tokens either way, because every offset is derived from the row's own token count rather than from the buffer's width. `pinning_the_window_is_numerically_inert` runs both settings against one stream for 40 steps and requires exact equality of the compressed rows and both time bases — and asserts the unpinned widths actually varied, so the equality is not vacuous. `ARC_V4_XS_PIN_WINDOW=0` restores the resizing buffer, which is also the A/B arm. split_row keeps the pinned buffer rather than re-narrowing it: `clone_out_cache` calls it once per layer per sequence on EVERY engine step, so narrowing there would undo the pin exactly on the hot path. The resume point does not move. Two fixtures now run explicitly unpinned, because pinning removes their discriminator rather than their subject: `ragged_xs_tail_is_refused_by_name_not_panicked` needs the widths 18/22 that only a resizing buffer produces, and `splitting_a_batched_row_restores_the_per_sequence_window` needs a row narrower than the shared window. Both keep testing what they were written for, and both gained a pinned-side counterpart. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * measure(arcspec): price the window pin in the same session — a third arm ON vs ON_UNPINNED differ in exactly one thing: whether the compressor's raw window is reallocated every decode step on all 41 compressed layers, or held at the bound the type documents. Same per-sequence flags in both, so the pin is isolated rather than confounded with the KV mechanism, and it is priced in the session that was already going to queue for the box. On the width discrepancy the pin rests on: this harness does not need to settle it. `max_tokens` is 4096 and each cell drives a 45 s steady window, so every request is thousands of tokens past saturation and the ramp is invisible here anyway. The settling is measured directly instead, on CPU, over 128 steps, by `the_window_ramps_then_settles_to_ratio_consecutive_sizes` — the early band sits below the steady one, the steady one is `ratio` consecutive sizes at [cap-ratio, cap-1], and no width ever reaches `cap`. That is a stronger instrument than a throughput leg and it costs no card time. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * instr(v4): name the window mode once per process, and assert the pin A/B is real D18, applied to the arm I am about to run. `ARC_V4_XS_PIN_WINDOW=0` is the only thing separating ON from ON_UNPINNED, and if that name were wrong the two arms would be the same build, produce a 1.000x ratio, and read as "the pin costs nothing" — the granted-but-inert failure with a throughput number attached. So the flag names itself once per process on first read ("xs rolling window is PINNED / RESIZING"), following the same convention as "per-sequence KV advance is ON", and the ladder summary now requires PINNED in the ON log and RESIZING in the ON_UNPINNED log before either arm's ratio may be read. If it cannot find both it prints PIN A/B IS VOID and says the comparison is of one thing with itself, rather than printing a ratio. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * measure(arcspec): make the pin control provably the control, not the default The window pin defaults ON, so the TREATMENT is the default and the CONTROL is the arm carrying `ARC_V4_XS_PIN_WINDOW=0`. That inverts the usual failure: a typo in the flag name does not disable the feature, it produces a control that silently ran the treatment and a clean-looking 1.000x ratio. Three things, all of which had to be right and only one of which was obvious: * The assertion reads the runtime LOG, never the binary. Both mode strings are compiled in unconditionally, so `strings <binary> | grep RESIZING` succeeds in every arm and proves nothing. Only the line a process emits says which branch it took. * `RUST_LOG=info` is set once inside `run_arm`, shared by every arm, so a log filter cannot suppress the line in one arm while leaving it in another. An assertion that configuration can mute is a guard with an undocumented off switch. * The guard now also fails when the control's log contains PINNED, not just when it lacks RESIZING — the leak is the thing being tested for, so it is checked directly rather than inferred from an absence. On VOID it refuses to let the ratio be read at all rather than printing it with a caveat. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcKV): one xs window pin, two triggers; scope the unpinned tests Integration fix for #121 landing after #181's `pin_tail_width`. Both arrived at the same idea — hold the retained raw xs window at a constant width instead of letting it breathe with `tokens % ratio` — for two different reasons, and both callers are real: * #181: CUDA-graph capture, where a moving allocation size is not slow but INVALID (a capture-time miss becomes an unstable graph memory node). Per-cache trigger, `pin_tail_width()`. * #121: serving, where the per-step reallocation is simply waste. Process-wide trigger, `ARC_V4_XS_PIN_WINDOW`. They agreed on the width and disagreed on everything else, so this keeps one width policy and both triggers: 1. `graph_tail_width()` and `window_capacity()` were the same expression (`span_groups * ratio + margin`) written twice. `graph_tail_width` now delegates. Two pins that computed "the pinned width" independently could drift, and capture and serving would then disagree about a number whose only value is that it does not move. 2. `retained_width` is gated on `pin_is_on()` = `self.pin_tail || xs_pin_window_enabled()`, so capture still gets its constant width when the env switch is off — which, for a capture-only run, it is. 3. #181's `pinned_base` form is dropped in favour of #121's. Both produce the same constant width; they differed in where the constancy came from. #181 moved `base` earlier so the logical span went constant; #121 leaves `base` at the retention point the type documents and widens the physical buffer. #121's is the one that composes, because it also moved `win_start` to the buffer's PHYSICAL start (`tokens - w_phys`) rather than `base`, which makes "buffer wider than the row promised" a representable state. `plan_xs_advance` still refuses to read below `base`, so the pin cannot change an answer. 4. Three `clone_in_cache_invariant_tests` assert the geometry of the UNPINNED window — exact tail widths `(4, 132)`, and that column 0 holds token `base`. Neither survives a pinned buffer, and #121 makes the pin the DEFAULT. They now run under `unpinned(..)`, the same `pin_test_override` shape #121 uses for its own affected tests. They are not testing a dead path: `ARC_V4_XS_PIN_WINDOW=0` is a supported mode and is the control arm of #121's own A/B. NOTE FOR REVIEW, not a defect: #121 flips the serving default — the window is PINNED unless `ARC_V4_XS_PIN_WINDOW` is set to 0. Its "+23.2% at uniform B=32" was measured on #121's own branch and is NOT re-measured here; this rebase changes which code that number describes. `cargo test -p mistralrs-core --lib`: 701 passed, 0 failed. Warning count back to master's 23 (the `xs_rolling` import is test-only and now lives in the test module). Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcKV): the xs window pin is OPT-IN, with the experiment that flips it named #121 shipped `ARC_V4_XS_PIN_WINDOW` default-ON, with `=0` as the control arm. Retracting the default and keeping the change. Why: the pin changes what the serving path RETAINS — at V4's HCA geometry the window is held at 144 columns where the resizing policy keeps as few as 4 at some residues — and the +23.2% at uniform B=32 that motivated defaulting it on was measured on #121's own pre-rebase branch, against a tree that no longer exists. That is an unmeasured default. TCFRAG was an unmeasured default too: it carried "UNVERIFIED ON HARDWARE — NEVER RUN" in its own header, held 63 GB and permanently broke a layer through a poisoned `OnceLock` (#209). So: unverified means default-off, and "unverified" means unmeasured, not new. The correctness argument #121 made for the pin is SOUND and is kept in place — it proves the pin changes no answer, which is necessary and not sufficient. The 40-step bit-identity A/B in `deepseek4` still guards it; its doc comment no longer claims to justify a default. ## Off is a temporary state with an owner, not the finish line Arc's larger problem is not unbuilt work — it is finished, correct, tested work left switched off. So this does NOT ship as a parked flag. The flip condition is written three places, one of them enforced: * `xs_pin_window_enabled_from`'s doc comment carries it in full, under a "FLIP CONDITION" heading, at the gate itself. * `capability_reachability.rs` — the file titled "the switched-off guard", which is where someone goes looking for exactly this — carries it as a registry entry, so the gate cannot go dark without CI going red. * `arcspec_perseq_ladder.sh`'s `ON_PINNED` arm IS the experiment. The experiment: one binary, uniform B=32, same prompt and seed, `ON` (flag unset) against `ON_PINNED` (`ARC_V4_XS_PIN_WINDOW=1`). Pass = ON_PINNED faster on aggregate tok/s with identical generated tokens. On pass the default flips to ON in the same change that records the number. ## The harness inverted with the default, and that is easy to get silently wrong The pin used to be the default, so the A/B's TREATMENT was the default arm and its CONTROL carried the flag. Now it is the other way round. Three things had to move together, and any one left behind would have produced a clean-looking and meaningless number: 1. The arm carries `=1` and is renamed `ON_UNPINNED` -> `ON_PINNED`. 2. The report's ratio pairing is swapped, so it stays treatment/control. Renaming the arm without this would have inverted every printed ratio. 3. The engagement guard's expectations are swapped: `ON` must log RESIZING, `ON_PINNED` must log PINNED, and the leak check now looks for a treatment that silently ran the control. Polarity is `== Some("1")`, split into the pure `xs_pin_window_enabled_from` and pinned by tests: unset is OFF, and `0`/`false`/`off`/`true`/`on`/`2`/`" 1"` are all OFF. #212 converted 23 `ARC_*` flags for this reason — `var_os(..).is_some()` made `ARC_FOO=0` mean ON, which turns an A/B control into a second treatment. A third test pins the thing most likely to be broken by this retraction: CUDA-graph capture must still get a constant width with the env flag off. It does — capture asks per-cache via `pin_tail_width()` and `pin_is_on` honours that trigger independently. Without that, capture would silently return to a per-step-varying allocation size, which is not a slow graph but an invalid one. `cargo test -p mistralrs-core -p mistralrs-quant -p mistralrs-vision`: green. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…to 0, opt-in (#217) * perf(ArcInfer/ArcGraph): device decode loop — N steps, N cuGraphLaunch, zero host syncs CUDA-graph capture already cut launch APIs per token 3,961 -> 3.8 and bought only ~8%. The reason is not that op count stopped mattering: replay() launches the graph and then blocks on cudaStreamSynchronize (graph.rs:362), the caller samples on the host, and the host-side argmax D2H (sampler.rs:1479) re-synchronizes anyway. Capture removed the cost of *issuing* the work and left the serialized host round trip between steps completely intact. Removing the sync alone buys nothing, so this removes the sync and the host sample together, behind an opt-in. Per step, all on one stream: cuGraphLaunch -> CudaSampler (on device, token to an i32 device buffer, never copied down) -> arc_graph_step_commit (scatters the token into the PINNED U32 buffer the captured graph reads, advances the device position, publishes into a pinned+mapped host ring). Nothing blocks. Fault detection replaces the sync rather than dropping it, with three signals that need no blocking call: non-blocking cudaStreamQuery (which NVIDIA documents as also returning errors from previous asynchronous launches), an idle-but-short ring (every step retired yet fewer tokens landed — a fault with no timeout heuristic in it), and a sticky device fault word the commit kernel latches before it writes anything. A spin backstop covers a kernel that hangs without ever erroring. The driver is expressed against a DeviceOps trait that deliberately has no synchronize method, so it is structurally incapable of blocking the stream, and it compiles without CUDA — 20 unit tests run against a scripted fake GPU on any host. Two mutations were checked: dropping the device-fault read and skipping one launch per burst each fail exactly the tests that name them. Aliasing: the logits tensor aliases storage the next launch overwrites. Safe here only because every access is stream-ordered and none is a host read; the ring the host drains is a separate pinned allocation no launch touches. The invariant that lets the aliasing tensor be returned at all is enforced at runtime, not just documented. Default OFF (ARC_GRAPH_DEVICE_LOOP). Greedy-only: the device sampler runs Splitmix64 against the host's Isaac64, so any stochastic mode would draw a different, equally valid token and silently break seeded reproducibility. Every refusal falls back to the existing path; one failure latches it off for the process. Kept as a pure insertion, away from replay(), to merge cleanly with PR #205. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcGraph): device loop stands aside for ARC_GRAPH_VERIFY_ALWAYS The device-loop branch sits earlier in the chain than the runner.replay(bs) arm that PR #205 extended with ARC_GRAPH_VERIFY_ALWAYS. Without this it would take every step and that switch would silently verify nothing — a clean diagnostic that never ran. Also cross-references #205's BOOBY TRAP note from replay_handles, since the same aliasing applies to the handles it returns. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * docs(ArcGraph): record the burst's scheduler-admission limitation A burst admits no new request for its whole length (default 4 decode steps). Harmless at batch 1 — the only shape admit() accepts — and not harmless above it, which is why admit() refuses batch_size != 1 outright rather than treating it as a tuning question. Stated in the module docs so it survives past the return it was first noted in. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcGraph): the device-loop sampling hook takes &mut Sequence The CUDA lane's whole purpose, on its first run against this branch: three jobs, one error, in the one block a macOS `cargo check` structurally cannot reach. error[E0596]: cannot borrow `*seq` as mutable, as it is behind a `&` reference --> mistralrs-core/src/pipeline/sampling.rs:557:12 557 | && seq.sampler().is_greedy_trivial() `Sequence::sampler()` takes `&mut self` (sequence.rs:885) even though it only clones an Arc, so the `&Sequence` parameter could never call it. `sample_sequence` already holds `&mut Sequence` and was reborrowing it down to `&`, which is why the call site compiled and the body did not. This is rustc's own suggested fix. Nothing here mutates the sequence, and the `&&` chain still borrow-checks: `sampler()` returns an owned `Arc<Sampler>`, so the mutable borrow ends before `seq.recognizer` is read. Verified by compiling the cuda body on a host with no CUDA — cfg removed, the three `arc_cuda_graph` calls stubbed — which puts the real `Sequence` through borrowck and passes. The `#[cfg(not(feature = "cuda"))]` twin moves with it so the two signatures cannot drift. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…els it proves were switched off (#196) * fix(ArcLab/ArcInfer): prefill measured 0.000 tok/s and aborted at teardown Three defects, all on the path that makes prefill measurable at all. 1. `pp N` reported `0.000±0.000` for every prompt benchmark. A prefill-only request (`max_len = 1`) finishes *during* the prompt step: `pipeline/sampling.rs` calls `seq.update_time_info()`, builds `group.get_usage()` and dispatches `Response::Done` from inside `step()`. But the default (non-paged) scheduler arm stamped `prompt_timestamp` and `total_prompt_time` only AFTER `step()` returned. So at the moment the usage was built `prompt_timestamp` was still `None`, `update_time_info` skipped `group.total_prompt_time`, and `get_usage` took its `== 0` branch and returned `avg_prompt_tok_per_sec: 0.0`. The PagedAttention arm already stamps before `step()`, with a comment saying exactly why. The fix was simply never applied to the arm that serves -- PagedAttention is banned here (it shadows the graph arm and measures zero tokens). Stamp on the default arm too. 2. The harness printed that zero into a results table and exited 0. A zero is not a measurement. Every row now has to carry evidence the engine processed the tokens the row claims to be about -- non-zero prompt tokens, the count actually asked for, a non-zero timed prefill, and a finite positive rate. A row that cannot exits 2 (environment failure, never 1) and prints no table. Adds a second, engine-independent instrument: wall-clock seconds and wall-clock prompt tok/s, so the engine's own numbers are only believed when an outside clock agrees. Also stops reusing one request id for every concurrent copy and repetition: `next_request_id()` was called once and the request cloned, so `concurrency * repetitions` sequences were in flight sharing one engine handle. 3. The binary aborted at teardown AFTER printing valid results. `Request::Terminate` only asks; nothing waited. `engine_handler` had no `.join()` anywhere in the workspace, so `main` returned while the engine thread was still releasing device memory and CUDA objects, and libc `exit()` ran CUDA's atexit handler concurrently with it -- a textbook `corrupted double-linked list` / SIGSEGV. `Drop for MistralRs` now joins each engine thread with a bounded 30s timeout, and says so loudly if it has to detach instead. Second half of the same race: `DedicatedDecodePath::new` binds the CUDA primary context and warns that skipping it gives "a SEGV in libcuda MOVAPS ... no error code, just a fault". Its `Drop` then made the same class of raw calls (`cuGraphExecDestroy`, many `cudaFree`) with no bind at all, made legal on a foreign thread by `unsafe impl Send`/`Sync`. It now binds exactly as `new()` does. This matters beyond tidiness: the abort libelled good runs. One chain already discarded an expensive valid measurement because its harness gated on rc==0. Parent system: ArcLab (harness) / ArcInfer (engine timing, teardown). * test(ArcLab): print the grouped-GEMM launch counters beside the results `grouped.rs` already states the rule: "a harness that reports a per-variant timing MUST show this counter advancing on the variant it claims to have measured, and NOT advancing on the others, or the number describes some other kernel." The counters existed and were public. Nothing in the workspace ever read them, so every grouped-vs-GEMV and variant-vs-variant comparison this harness could produce was unproven by construction. The bench now prints `grouped_launch_counts()` and the selected variant next to every results table, and says so explicitly when the total is zero -- which is the common case, because the grouped GEMM declines to run below its tile-fill boundary (~683 tokens at V4's top-6-of-256 routing) and the measurement is then the per-pair gather GEMV, not the kernel the reader will assume. Parent system: ArcLab. * perf(ArcQuant): ship the tuned grouped-GEMM kernel as the default (+19.5% prefill) The trellis grouped GEMM has three variants compiled into every binary. The tuned ones were measured long ago at +39.4% (v1) and +41.6% (v2) per m-tile, bit-identical in output, and then left unreachable: `grouped_variant()` fell back to `QTIP_GROUPED_VARIANT_BASELINE` whenever `ARC_QTIP_GROUPED_VARIANT` was unset, and the only callers of `set_grouped_variant` in the entire workspace are inside `examples/qtip_grouped_curve.rs`. So every production prefill since the variants landed has run variant 0, and the faster kernels existed only for a benchmark nobody ran in serving. ✅ MEASURED END-TO-END, H200, DeepSeek-V4-Flash qtip2b, pp2048 b=1, --no-paged-attn, on `arc-prefill`: variant 0 (was default): 561.7 prompt tok/s | 1.780 ms/tok | TTFT 3.675 s variant 2 (now default): 671.3 prompt tok/s | 1.490 ms/tok | TTFT 3.080 s => +19.5% prefill throughput, -16.2% TTFT The arm is proved from the runtime launch counters, not from the build: [129,0,0] on the control against [0,0,129] on the variant -- 43 layers x 3 expert matrices, and zero launches on the arms not under test. An engine-independent wall clock agrees (557.2 -> 664.9 tok/s, +19.3%). The end-to-end gain is smaller than the kernel gain because the expert gather is ~56% of a 2048-token prefill step; backing +41.6% on that share out of a measured +19.5% overall is consistent. Also: an unrecognised env value used to fall through to the baseline. With the baseline now the slow arm, a typo would silently cost ~20% of prefill, so unknown values warn and keep the default, and `baseline`/`0` is now an explicit arm so the A/B knob still selects the control. Parent system: ArcQuant / QTIP. --------- Co-authored-by: Nirupam Bhowmick <support@runcrate.ai>
… +69% (114.76 -> 193.46 tok/s) (#195) * perf(ArcInfer/ArcMoE): warp-butterfly Sinkhorn on top of the ladder — 1.36x Rebased onto release/openrouter-ready. The earlier version of this change was measured against the PRE-LADDER kernel and read 4.19x; that baseline is gone. Against the templated kernel the ladder actually shipped, the honest number is 1.36x. Recording both so the retracted figure is not quoted again. MEASURED, H200 / nvcc 12.4 / sm_90, n=1 hc=4 iters=20, 86 calls/token: pre-ladder (runtime hc, 192 B stack frame) 30.578 us/call 2.6297 ms/tok templated + shared/barriers (master today) 9.771 us/call 0.8403 ms/tok warp butterfly (this commit) 7.201 us/call 0.6193 ms/tok => 1.357x incremental, -0.221 ms/token. Bit-identical to the templated arm. The templated kernel got the row/tree buffers into registers but still stages the matrix through shared memory and pays TWO __syncthreads() per iteration -- 40 barriers and ~80 shared round trips per call -- purely to move each column to the thread that divides by it. hc <= 16 means every thread is a lane of ONE warp, so the column reduction can be a butterfly and the matrix never has to leave registers. Bit-identity by construction: candle's level s is `buf[t] = __fadd_rn(buf[t], buf[t + s])` for t < s. The butterfly gives lane t `__fadd_rn(buf[t], buf[t ^ s])`, which IS that expression for t < s and its commuted form for t >= s; IEEE-754 addition is commutative (only associativity fails), so lane 0 ends holding precisely candle's result. Lanes with row >= HC feed 0.0f, exactly candle's `shr[tid] = 0` identity padding. ARC_NO_SINKHORN_WARP=1 restores the templated kernel. The arms are separable in a profile by kernel name (`arc_sinkhorn::warp_kernel<N>` vs the anonymous-namespace `sinkhorn_normalize_f32_kernel<N>`), so which arm ran is provable from nsys alone and a stale object file cannot fake it. No Rust changes; the C ABI is unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * instrument(ArcQuant/QTIP): count expert-weight bytes read per step Answers whether the MoE path amortises expert weights across a batch. Flat bytes/step with batch => amortisation works and the aggregate deficit is elsewhere; linear => every token drags its own copy of the expert weights. METHOD: kernel-issued bytes, computed exactly from each launch's own geometry. Not a profiler sample and not DRAM traffic. `ncu` cannot attach on the rented H200 ("Failed to prepare kernel for profiling / Unknown Error on device 0", verified on this box, not inherited), there is no DCGM, and `nvidia-smi dmon` reports a utilisation percentage rather than bytes -- so no hardware DRAM counter is available and the accounting is done from the dispatch instead. The gather GEMV (qtip_gather_gemv.cu) is indexed `pair = blockIdx.y` with `row-block = blockIdx.x`, and each block loads its own expert's packed rows. Summed over blockIdx.x one pair reads one WHOLE expert, and nothing is shared between pairs, so a launch reads n_pairs * n_rows * packed_per_row bytes. The grouped GEMM stages a woken expert once per GROUPED_TILE_M pairs, so it is accounted as ceil(pairs_e / TILE_M) copies. Both arms report in the same units so they are directly comparable. Because these are loads ISSUED they upper-bound DRAM traffic: repeated reads of one expert inside a short window can be served by L2. That is the point -- 99-100% SM occupancy against a 3% memory controller is what re-reading a redundant working set looks like, and the redundancy factor is the quantity that must reach 1.0 for batching to amortise, whichever cache serves it. ARC_MOE_BYTE_PROBE=1 to enable; off by default. One relaxed atomic add per expert-GEMM launch when on -- no device sync and no D2H read, so it is capture-safe and does not perturb the dispatch it measures. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * test(ArcQuant/QTIP): pin the measured expert-byte curve MEASURED on H200, V4-Flash qtip2b, ARC_MOE_BYTE_PROBE=1. One expert is 2 MiB; a step is 129 gather invocations (43 layers x 3). Expert bytes per step: n_tokens GEMV arm grouped arm redundancy 1 1.5 GiB 1.5 GiB 1.00x 8 12.1 GiB 12.1 GiB 1.00x 64 96.8 GiB 64.5 GiB 1.50x 512 774.0 GiB 64.5 GiB 12.00x The GEMV arm is exactly linear in the batch -- 512x the users reads 512x the expert bytes, so nothing is shared. The grouped arm is FLAT at the whole expert working set once every expert is woken. Two independent anchors agree: - a real 8-user decode through the server reads 12.09 GiB/step, identical to the prefill-derived n_tokens=8 figure, so the sweep measured the same quantity that batched serving pays; - the amortised floor of 64.5 GiB/step is 95.1% of the 67.82 GiB of qtip2b shards on disk, i.e. it is the whole expert working set read once. This is a characterization test, not an endorsement: it encodes the defect so that fixing the dispatch fails here first and the numbers get updated. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * test(ArcQuant/QTIP): pin the corrected dispatch, with the verified numbers The shipped default (no env override) now selects the flat-bytes arm above the capture-safety floor. Verified end-to-end on H200: B old tile-fill gate shipped default delta 8 57.18 60.28 +5.4% 64 106.26 156.61 +47.4% 256 114.76 193.46 +68.6% Kernel identity is from the profile, not the gate condition: at n_tokens=256 the default runs qtip2b_grouped_gemm_kernel n=129 -- exactly one step's worth of gathers (43 layers x 3) -- while n_tokens=1 decode steps stay on qtip2b_gemv_tuned_kernel below the floor, which is the intended split. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
… reachable (#130) The headline is an ABI mismatch that explains a failure already on our record. `ffi.rs` declared `cudaGraphAddNode` with FIVE arguments and the first two transposed. CUDA 13.1 `cuda.h:21829` declares SIX, and the first parameter is an OUT pointer: CUresult cuGraphAddNode(CUgraphNode *phGraphNode, CUgraph hGraph, const CUgraphNode *dependencies, const CUgraphEdgeData *dependencyData, size_t numDependencies, CUgraphNodeParams *nodeParams); So the driver received a `CUgraph` handle where it expected `CUgraphNode *` and wrote the new node handle THROUGH it, while `numDependencies` received a pointer and `nodeParams` an uninitialised register. That is memory corruption from a declaration, and it matches the previously unexplained "host heap corruption" CUDA-graph capture failure we had recorded with no cause. Four more defects, each of which alone stops the autonomous loop: * Conditional handles were created with `flags = 0`. `cuda.h:21919` applies `defaultLaunchValue` only when `CU_GRAPH_COND_ASSIGN_DEFAULT` is set, so the condition read 0 at launch and a WHILE body executed ZERO times while every CUDA call returned success. Measured both ways on an H200. * The WHILE body was captured into a throwaway graph and then DESTROYED, leaving the conditional body empty. Bodies are now populated with `cuStreamBeginCaptureToGraph` (`cuda.h:1976` names it as the supported way), and a node-count assert refuses a body that recorded nothing. * `CUDA_CONDITIONAL_NODE_PARAMS` was missing its 5th field, `ctx`. * `DecodeState` allocated I64 while every kernel signature is `int32_t*`, so host and device disagreed about where element `i` lived; and `reset()` reallocated the buffers, silently invalidating the pointers a captured graph had baked in. `tensor_device_ptr` had no I32 arm, so `CudaSampler::sample` returned `unsupported dtype I32` on EVERY call -- the correct top-k/top-p sampler had never once executed on a GPU. Its test suite is a CPU simulator, which by construction cannot observe that. Sampling now runs on device and survives replay. The sampler previously wired into the graph body took `rng_offset` as a BY-VALUE kernel argument, which a captured graph bakes: measured 1 distinct token over 64 replays, versus 64 distinct for a device-resident `rng_state`. It also was not nucleus sampling -- it walked the vocabulary in TOKEN-ID order. The replacement is a hybrid, because measurement showed neither component wins everywhere (H200, vocab 129280, host+GPU verified exclusive before and after; support width -> legacy / hybrid us): 1 -> 134.9/166.6 8 -> 503.0/624.2 64 -> 3127.8/4178.5 512 -> 24320.7/4185.0 4096 -> 205125.2/4187.3 12928 -> 664200.0/4188.5 Exact enumeration while the nucleus is small, threshold bisection past a fixed budget, chosen block-uniformly ON DEVICE because a captured graph cannot ask the host which branch to take. Costs +24% on the peaked distributions real models produce -- 32 us, 0.05% of a 66.68 ms V4 decode step -- and removes a 664 ms cliff that was ~10x an entire decode step for a single token. No cap and no narrowing: exceeding the budget changes which ALGORITHM selects the nucleus, never which tokens it contains. Verified against the enumerating sampler with branch counters proving which path each case exercised: peaked TV=0.0000 A=20000 B=0 scrambled order TV=0.0000 A=20000 B=0 (15.6% per-draw agreement, so the orders genuinely diverged) diffuse TV=0.0000 A=0 B=3000 tie boundary TV=0.0000 A=20000 B=0 narrowed (ctrl) TV=0.5027 -> correctly DISAGREE Known and deliberately left: on TIED diffuse supports the fallback keeps 116 tokens where enumeration keeps 111 (TV=0.1103) -- it keeps every token sharing a boundary bit pattern, erring WIDER, never narrower. And the budget is derived from bisection's pass count (32) when it should follow the FALLBACK's (~70), which costs a 0.75x window at support ~64; correcting it needs a more diffuse fixture so the fallback stays covered. Not claimed: no V4 end-to-end number. Autonomous decode is unreachable on V4 by construction -- no PagedAttention means `cache_config` is None and the runner is never built (`normal.rs:1907`, pinned by `normal_loaders.rs:5687`) -- and `cuda.h:1971` bars alloc/free nodes inside a conditional body against 11,436 allocations per token. Co-authored-by: Claude (ArcGraph chain) <noreply@anthropic.com>
…es/step -> 301 (#198) * perf(ArcAttention): batch V4 RoPE over the cohort — 99,072 launches/step -> 301 `DeepSeekV2RotaryEmbedding::{forward,forward_inverse_tail}` gated the batched path on `seqlen_offsets.len() == 1` — the LENGTH of the offset vector, where the thing that matters is the DISTINCTNESS of its values. `len()` is the batch size, so the batched path was unreachable at every batch size above one. The offsets are uniform anyway. `scheduler::default_scheduler`'s `select_running_bucket` admits a forward pass only when cache lengths are exactly equal, and the one producer of a ragged dense batch is gated behind `ARC_MTP_PER_SEQ_KV`, which defaults off (`deepseek4::ragged_row_q0`). So the loop was performing B bit-identical recomputations and concatenating them back. Measured, B=256 pure decode, steady state (nsys, 204.80 tok/s): kernel share launches/step per (layer, seq) ucopy_bf16 10.2% 22,794 2 (2*43*256 = 22,016) copy2d_bf16 5.0% 33,610 3 (3*43*256 = 33,024) rope_i_bf16 3.9% 33,054 3 (3*43*256 = 33,024) uneg_bf16 0.8% 11,008 1 (1*43*256 = 11,008, exact) 19.9% of step time, and the only cost in the profile that gets WORSE as batch grows. `copy2d` is the 256-argument `Tensor::cat` that exists solely to undo the split the same function just made; it goes to zero. `uneg` is 11,008 launches to negate a [1,32] slice of a CONSTANT table, so `neg_sin` is now materialised once at construction — not a kernel to fuse, a kernel to delete. Cohort form is 43 layers * 7 = 301 launches (129 of them rope_i), which lands RoPE at the same order as fp8_matmul_tiled (301) and qtip2b_grouped_gemm (129) instead of dwarfing them. PREDICTION, not a result: this is a launch-count fix, the bytes moved are unchanged, and how much of the 19.9% actually returns must be measured on hardware. Ragged offsets still take the loop verbatim, so enabling ARC_MTP_PER_SEQ_KV behaves exactly as today. No ragged index_select path is added — nothing ships dark. Bit-identical by construction: with 2-D [T, D/2] cos/sin the kernel's stride_b is 0 (candle-nn/src/rotary_emb.rs), so every batch row reads the same cos/sin row and each output element is an independent two-multiply-one-add. Seven tests hold the pre-fix loop as a verbatim oracle and assert raw IEEE-754 bit equality (not a tolerance) for uniform AND ragged offsets, at seq_len 1 and 2, with a one-ULP negative control proving the comparator can fail. Engagement is counted (`rope_cohort_stats::counts`): a fast path that silently declined would pass every equality assertion trivially. The counter already earned itself — it caught a missing lock in one of these tests. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(ArcAttention): un-deaden the fused qk_norm_rope kernel above batch 1 `fused_qk_gate` declined with "per-sequence position offsets" whenever `seqlen_offsets.len() != 1` — the same length-vs-values confusion as the RoPE dispatch in the parent commit, and with the same consequence: the fused one-launch kernel was dead at EVERY batch size above one, i.e. in exactly the serving regime it was written for. Its own doc comment advertises "replacing 16 candle launches per layer with 1"; at B=256 it replaced nothing. The kernel is already batch-correct and needed no change. `qk_norm_rope.cu` flattens (b, t) into blockIdx.y, recovers `b = bt / seq_len`, and reads the table at `pos_offset + t` — a position independent of `b`, which is exactly right when the rows share an offset. Its grid is (n_heads + 1, batch * seq_len) and the wrapper sizes buffers as `batch * ...`. Nothing assumed batch == 1. Gate now takes the uniform offset via the shared `DeepSeekV2RotaryEmbedding::uniform_offset` helper, so both call sites cannot drift apart. Genuinely ragged rows still decline and still take the eager chain. Split from the parent commit so it can be reverted independently: this one swaps in a CUDA kernel that has never run above batch 1, and unlike the parent it is not bit-identical-by-construction — it is bit-identical by the kernel's own design, which `ARC_QK_VERIFY=1` checks against the eager chain at every layer. Verify before trusting; `ARC_QK_FUSED=0` restores the eager path from the same binary. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcAttention): one cohort predicate, one counter pair, per-thread test counts Integration fix for #198 landing after #204. Both PRs made the same correction -- `seqlen_offsets.len() == 1` tests the length of the vector where the property needed is the distinctness of its values -- and each shipped its own spelling of the predicate and its own engagement counters. `uniform_seqlen_offset`'s doc comment on the integration branch asked for exactly this collapse ("there must not be two spellings of the same predicate"); this is that collapse. 1. `pub mod rope_cohort_stats` was declared TWICE in one file (layers.rs:400 from #204, :1568 from #198). That is a hard name collision -- the file did not compile. #204's is the superset (it has `record_cohort` / `record_per_sequence` as well as `counts`), so #198's is the one that goes. 2. `DeepSeekV2RotaryEmbedding::uniform_offset` had a byte-identical body to `uniform_seqlen_offset`. Removed; `forward`, `forward_inverse_tail` and `deepseek4::fused_qk_gate` now call the free function, so the two cannot drift apart. The reasoning #198's doc comment carried and #204's did not -- that uniformity is *enforced* by `select_running_bucket` rather than merely observed, and the measured B=256 cost -- moves onto the survivor. 3. Collapsing the counters made them shared, and shared process-global counters cannot be delta-asserted from a parallel test binary. `deepseek_rope_cohort_ tests` held a private mutex, `rope_cohort_tests` held none (until now it was the only writer), and three tests duly failed with `(2, 1)` where they wanted `(1, 0)`. A shared mutex was the first fix and it is also wrong: it obliges every present and future test that touches ANY rotary to know about a lock in another module. Two such tests already exist and did not -- `deepseek4::tests::{standard_layer_rope_is_unscaled_and_compressed_is_not, mtp_block_takes_the_standard_unscaled_rope_table}` both drive `DeepSeekV2RotaryEmbedding::forward` -- and they race the cohort tests. That was caught by merging #198 and #121 together and running the full suite: `--test-threads=1` passes 707, the parallel run fails one. So the counters now also keep a `#[cfg(test)]` per-thread mirror, and the tests assert on `local_counts()`. Each test's delta is its own by construction, no cooperation from other modules is required, and the serving instrument -- the global `counts()` -- is byte-for-byte unchanged. Also drops #198's copy of `uniform_offset_tests_values_not_length`: with one predicate it asserted nothing the surviving copy does not, against the same six cases. `cargo test -p mistralrs-core --lib`: 699 passed, 0 failed; 707 passed on the merged #198+#181+#130+#121 tree. The 99,072 -> 301 launch figure remains #198's own branch measurement and is NOT re-measured here. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcAttention): the batched fused qk_norm_rope path is OPT-IN The parent commit un-deadens the fused Q/K kernel above batch 1 by fixing a length-vs-values gate. That is the right fix to the gate and the wrong default: it makes a CUDA kernel LIVE on a path it has never run on. This repo has paid for that twice in one week. * TCFRAG shipped default-ON carrying "UNVERIFIED ON HARDWARE — NEVER RUN" in its own header. It held 63 GB and permanently broke a layer through a poisoned `OnceLock`. #209 retired it to opt-in. * The fused-512 attention path ran for four days silently dropping the attention mask, at 12% agreement with the reference, because nobody ran the path it had quietly become live on. So the doctrine, and this commit: **unverified means default-off, and "unverified" means unmeasured, not new.** `ARC_QK_FUSED_COHORT` gates the batched path, default OFF. With it unset, `batch > 1` declines and takes the eager chain exactly as master does today — so the parent commit ships its gate fix with no behaviour change, and the cohort path is available to be measured rather than discovered in production. `batch == 1` is untouched: that path is already live and already exercised. Read by VALUE, `== Some("1")`, split into the pure `fused_cohort_enabled_from` so the polarity is testable without mutating the environment (racy across `cargo test`'s threads, `unsafe` since the 2024 edition, and meaningless against a `OnceLock` that latches the first read). Two tests pin it, and they run with no features and no GPU on every PR: unset is OFF, and `0`/`false`/ `off`/`true`/`on`/`2`/`"1 "` are all OFF. That second one is not pedantry — #212 converted 23 `ARC_*` flags this week because `var_os(..).is_some()` made `ARC_FOO=0` mean ON, which silently cancels any A/B whose control sets zero. To turn it on, once `ARC_QK_VERIFY=1` has bit-compared it against the eager chain on hardware and it is measured faster: `ARC_QK_FUSED_COHORT=1`. When that measurement exists, this default flips and the doc comment goes with it. `cargo test -p mistralrs-core --lib`: 701 passed, 0 failed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(ArcAttention): gather the RAGGED cohort's RoPE — one rank-3 rope_i, opt-in ARC_ROPE_COHORT The uniform-offset fix (#198, and the shared predicate it now rides on) left one arm untouched: rows at genuinely DIFFERENT positions still take the per-sequence loop — ~9 ops x 43 layers x B per step on the measured B=256 profile (~99,000 launches, 19.9% of step time), of which the copy2d half only undoes a split the same function just made. Nothing structural required the loop: candle's rope_i accepts a rank-3 [B, T, D/2] cos/sin whenever B matches the input batch (rope_check_cs, candle-nn/src/rotary_emb.rs, pinned rev 89ab14e; the CUDA wrapper sets stride_b for rank-3, the CPU path indexes i + b_i*t*d/2). So the ragged arm becomes: gather the per-row table rows once — [B*T] u32 indices, host-built because the offsets ARE host data, memoised per instance so the H2D and the three index_selects run once per step per device (the model shares one rotary per device per table kind), CLAUDE.md pitfall #5 — then ONE rope_i per projection. ~5 launches per layer where the loop paid ~9 x B. DEFAULT OFF behind ARC_ROPE_COHORT, read by VALUE via env_flag_is_set (=0 means OFF; #212), latched once. Unset keeps the loop byte-for-byte. Registered in capability_reachability.rs beside the fused-cohort entry, same doctrine: unverified means default-off, and unverified means unmeasured. Both ragged arms are named methods driven directly by deepseek_ragged_cohort_tests — bit-pattern parity (not tolerance) for mixed offsets at B in {1,3}, T in {1,2}, ascending and descending, forward and inverse tail, plus a stale-memo test (the cache must rebuild when the cohort moves) and a pinned unset-default. engaged_count in cuda/qk_norm_rope.rs also gains its first reader: a test that proves the counters advance, closing the silent-success hole where an A/B could read a dead path as live. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcAttention): bounds-guard the cohort gather — both ragged arms must fail LOUD The loop arm's narrow(0, offset, seq_len) hard-errors past the table end; the gathered arm fed the same out-of-range index to index_select, which is not reliably bounds-checked on the CUDA backend and can read garbage SILENTLY. A failure-mode asymmetry on exactly the arm the hardware A/B will exercise is the D18 house fault, so the offsets — host data, the check is free — are now guarded host-side in cohort_tables before any index reaches the device, with the offending offset, seq_len and table size named in the error. out_of_bounds_offset_fails_loud_on_both_arms pins: the loop arm errors, the cohort arm errors on forward AND inverse tail, the message is the GUARD's (a backend error cannot pass the test by coincidence), and offset 31 + T=1 on a 32-row table still works, so the bound is not off by one. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Nirupam Bhowmick <support@runcrate.ai> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…gate (#219) * feat(ArcSched/ArcKV): lift the three scheduler caps below the B>=214 gate — serve default 256, device-side ragged mask, prefill floor default 4 Parent system: ArcInfer / ArcSched + ArcKV. CAP 1 — admission default. mistralrs-cli's max_seqs default goes 32 -> 256 (arg default, struct default, default_max_seqs(), and the tune config template). At B=32 the physics ceiling is ~3,740 tok/s, so the 14,000 tok/s aggregate gate (B >= 214) was unreachable on the default configuration. Memory implications documented on the flag; PagedAttention admission is block-budgeted, DefaultScheduler admission is count-only (VRAM-aware check for the dense path is named follow-up, not built). CAP 2 — the mask that would make ragged decode read as a regression. CausalMasker::make_left_padded_causal_mask was a scalar triple loop over B * t_q * (past + t_q) host f32s plus a full-mask H2D copy, per decode step while a cohort is ragged — 4.2 MB/token at B=256 ctx-4096 (CLAUDE.md pitfall #5 on the hot path). It is now built on the device from broadcast comparisons; the only per-step H2D is the [B] lead vector, and the index rows come from a per-device cached arange grown geometrically. The host loop is kept verbatim as make_left_padded_causal_mask_host, the test oracle, and the two are pinned element-for-element across mixed lengths, sliding windows, multi-row queries and the degenerate past=0. CAP 3 — the prompt-starvation floor defaults ON at 4. Behind a 47-sequence decode cohort a fresh prompt wins 0 of 24 steps with no floor and 1 in 5 with floor=4 (this file's own fixture). ARC_PREFILL_FLOOR_STEPS=0 is the kill-switch restoring the pre-floor selection key for key; the env mapping is factored into floor_from() and pinned by test on both sides. The V4 ragged-pair default is left OFF deliberately: on this branch the pair is --v4-ragged-decode / ARC_V4_XS_PER_SEQ (the xs cache) plus ARC_MTP_PER_SEQ_KV (the per-row query positions in dsv4_attention via ragged_row_q0). The second gate lives outside this change's file fence, and opening only the first would let shorter V4 rows attend compressed blocks they have not reached — silently. The box run flips the two together on a measured pass. New/updated capability registry entries: the device-built mask, the scheduler's ragged_decode_supported read, the request_xs_per_sequence CLI/config latch, and the floor entry renamed to carry its new default. New tests: a_formed_cohort_admits_a_newcomer_when_ragged_decode_is_on (cohorts could previously only shrink), the_device_mask_matches_the_host_oracle_element_for_element, the_device_mask_and_the_oracle_refuse_the_same_inputs, the_floor_defaults_on_and_zero_is_the_kill_switch. All mutation-checked; values recorded in the test docs and the PR body. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcSched): a 1-token prompt keyed into the pinned ragged decode bucket — prompt-ness joins the BucketKey; typos lane green The collision the PR body flagged as UNRESOLVED is real, and is now fixed and pinned by test. With ragged decode published, decode sequences key the sentinel length 0 — and a 1-token prompt's cache_bucket_len() is len() - 1 == 0, the same number. The probe test showed the prompt riding into the decode cohort through the merged bucket's `<= 1` fast path: scheduled in the same pass (prompt = [4] alongside completion = [0, 1, 2, 3]), bypassing select_running_bucket, the floor's bookkeeping and the one-bucket-per-step invariant. Fix: `is_prompt` becomes a fourth BucketKey component, so a bucket is homogeneous by construction. The coalescing filter requires matching prompt-ness (a decode bucket must never idle to catch up to a prompt bucket), and the floor's is_prompt_bucket now reads the key instead of whichever sequence sits first in the bucket — removing an order-dependence at the same time. Test: a_one_token_prompt_is_not_folded_into_the_pinned_decode_bucket. Green: completion = [0, 1, 2, 3], prompt = [], waiting = 1. Mutated (fourth component pinned false): red with prompt = [4]. All 14 scheduler tests and the full 702-test core lib suite pass. Typos lane: "unparseable" -> "unparsable" (floor doc), and word-level allowlist entries for arange/ARANGE in .typos.toml — candle's Tensor::arange naming appearing as a segment inside cached_arange_f32 / ARANGE_CACHE, which the existing exact-identifier entries do not cover. `typos --config .typos.toml .` is clean repo-wide. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
… the V4 union mask across layers (#221) * fix(ArcAttention): the Metal sinks arm still dropped the caller's mask CORRECTNESS, the other backend. #199 gated the CUDA fused arm on mask.is_none(), but the Metal arm kept engaging on head_dim alone and calling flash_attn_sinks_metal — which takes no mask parameter at all, so unlike the CUDA kernel it cannot even refuse one at the interface. GPT-OSS reaches this arm on every macOS build: head_dim 64 is in Metal's fused set and the model passes a mask it asserts is Some (gpt_oss.rs). MEASURED on a real Metal device (M-series, head_dim 128, F32): max|fused - reference(causal only)| = 0.000001 max|fused - reference(masked)| = 0.860029 max|ref(causal) - ref(masked)| = 0.860029 <- the mask's whole effect The maskless kernel tracks the causal-only reference to 1e-6 and misses the masked one by the mask's entire effect: dropped, not approximated. The fix is the same one line as CUDA's: an explicit mask pins the call to the unfused mask-honoring path, with a once-per-process warn naming the deliberate perf loss. New tests: * metal_sinks_mask_tests (feature = "metal", real device): the baseline above, plus the gate — masked dispatch vs masked reference, diff 0.0. Mutation-checked on hardware: with the guard reverted the gate test fails at max abs diff 0.8600293; restored, 0.00000000. * padding_mask_is_honored_and_dropping_it_changes_the_output (CPU, always runs): a padding-column mask — the class no fused kernel can re-derive — survives the dispatch end to end. Mutation-checked: dropping the mask at the fall-through arm fails it at max abs diff 0.9353 vs 0.0 required. Rider (perf, contained): sinks_attn_cpu_varlen rebuilt its per-sequence mask with Tensor::from_vec — a host-to-device upload per sequence per layer per step on the GPU fall-through (CLAUDE.md pitfall #5). It is a pure function of (q_len, kv_len, window, dtype, device), now memoized in a single slot; uniform batches build it once per shape. varlen_mask_memo_hits_and_never_ serves_stale pins both memo failure modes (mutation-checked: with the key comparison removed it fails on the (5, 8, 3) entry). Also corrects the file docs that called the unfused fall-through "CPU": it is device-agnostic candle math and is exactly what V4's head_dim-512 layers run on CUDA. The to_vec1 D2H syncs on cu_seqlens remain — hoisting the host copy into FlashParams is shared plumbing, reported not fixed. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * perf(ArcInfer/ArcAttention): memoize the V4 union mask across layers The union mask is a pure function of nine integers — (q0, t_q, raw_base, t_k, t_c, ratio, window, dtype, device) — yet dsv4_attention rebuilt it once per LAYER per step. On the prefill path that is three Tensor::arange host-to-device uploads plus their casts, compares, the U8 cat and the 0/-inf conversion, every layer (CLAUDE.md pitfall #5: each arange with a GPU device is a CPU-to-GPU sync). a7b0890/64fe1d678 already killed the DECODE-path aranges; this removes the per-layer rebuild wholesale on both paths. Mechanics: the validity build is now lazy (build_valid closure, character-for-character the old code on a miss) and the additive [1, 1, t_q, n_keys] mask is memoized thread-local on the integer key — 4 FIFO slots, since one step touches at most three keys (the Standard / CSA / HCA layer classes). Per step, each distinct mask is built once instead of ~once per layer. Excluded on purpose: * graph_positions (fixed-capacity graph decode): the key would be constant across steps, and the ops must stay inside the capture region rather than be hoisted into a warmup-owned tensor whose later eviction would leave a recorded graph reading freed memory. That build is already arange-free. * ragged cohorts (row_q0): the per-row leads are content but not key. Tests (CPU, validating code not perf — D14): * memo_hits_without_building_and_misses_on_a_changed_key: a hit returns the cached handle WITHOUT running the builder (handle identity is the only observable — a memo that always misses returns correct values while doing nothing); an advanced q0 must miss. * memoized_masks_are_bit_identical_to_fresh_builds: decode at ctx 35 then 36 (ratio 4, window 8) — two masks with the SAME shape and different content (threshold crosses 8 -> 9). Sequential same-thread runs must be bit-identical to fresh-thread (empty-memo) references. Mutation-checked: with the memo's key comparison replaced by "serve the newest entry", both tests fail — the identity test at TensorId(1) != TensorId(6), the end-to-end test with ctx=36's output bits diverging from the fresh reference (first lane 1051680768 vs 1051066368). No behavioral change: a miss runs the exact pre-change build, and the existing 38 dsv4_attention tests (bit-identity, ragged, graph-geometry, absorbed-decode) all pass unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * docs(ArcAttention): reword 'aranges' for the Typos lane (config untouched; #219 owns the allowlist) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…ICE ends the 66 MB/step host copy that walled off the GPU sampler; 3 fast-path aggravators (#222) * perf(ArcInfer/ArcSample): port the STEP_us TOTAL/fwd/sample/other split onto the non-paged step arm V4 actually runs The only TOTAL/fwd/sample/other instrument lived on the PagedAttention arm, which V4 can never reach (DeepSeekV4Loader::supports_paged_attention is false). The DefaultInstructions arm — the path V4 runs — had never once reported how a step divides between GPU and host. Log format byte-identical to the paged arm's so tooling parses both. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcInfer/ArcSample): three measured fast-path aggravators + the device-eligibility predicate (a) has_penalties(): a DRY config at its disabled default (multiplier 0.0, a proven no-op in apply_dry_penalty) no longer disqualifies the GPU fast path or is_raw_argmax. (b) sample_top_kp_min_p: the server defaults (top_k=-1, top_p=1.0, min_p=0.0) ran a full-vocab sort_unstable_by (129,280 elements on V4, per sequence per token) whose result nothing consumed — skip it. Seeded draw provably identical (thread-local sort counter + pinned token 506 in tests). (c) sample_multinomial: the global rng MutexGuard was held (temporary lifetime extension) across get_top_logprobs and tokenizer.decode — narrowed to the draw. Plus Sampler::device_sampling_terminates, the single-source mirror of sample()'s device dispatch that ARC_SAMPLE_ON_DEVICE partitioning consults, and the ignored $0 host-cost probe at V4 vocab (D14: host evidence only). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * test(ArcInfer/ArcSample): partition split tests — copy discipline + seeded token parity for ARC_SAMPLE_ON_DEVICE Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * chore(ArcInfer/ArcSample): u32::try_from for the gather row index (clippy-by-hand for mistralrs-core) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…ispatch engagement log, three-arm sweep harness, SASS probe (#218) * fix(ArcKernels): the cuBLASLt FP8 branch never flattened rank-3 to 2-D The native branch flattens `[B, T, hidden]` to `[B*T, hidden]` before its GEMM. The dequantize + cuBLASLt branch did not — it passed `x` through untouched. V4's activation is rank-3, and bias is `None` on every V4 linear, so `UnquantLinear::forward` took `w.broadcast_left(B)` into `cublaslt.batch_matmul` with `stride_b = 0`. At decode T=1, making it a **B-way batched GEMM with m=1**: a batched GEMV wearing a GEMM's name, and the one shape cuBLASLt has no advantage in. The documented 27x was measured at prefill, where `[1, 512, hidden]` collapses to batch=1 and the call is a real GEMM. This invalidated the threshold A/B rather than merely slowing it: lowering `ARC_FP8_CUBLAS_MIN_M` selected the batched-GEVM shape, so the sweep was comparing the tiled kernel against a degenerate call and would have supported "cuBLASLt loses at decode" — a conclusion about a shape the change was never meant to select. Flattening makes the two branches comparable, which is the premise the A/B rests on. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(ArcKernels): pow2 shift fast path for the tiled FP8 GEMM's per-element scale divisions get_scale in fp8_matmul_tiled performs TWO runtime signed integer divisions per weight element staged into shared memory (n / block_size_y, k / block_size_x). Both divisors are kernel arguments, so nvcc expands each into a ~15-20 instruction reciprocal-and-fixup sequence — paid once per element of every weight tile, ~33.9G times per B=256 decode step by the source-derived count. The shipped weight_block_size is [128, 128]: both powers of two. This adds a POW2_SCALE template arm: the host launcher tests pow2-ness once per launch, computes shift = log2(block_size) on the host, and the kernel indexes the scale grid with two SHFs instead. Non-power-of-two geometries keep the division path, unchanged. Bit-identical by construction: for the non-negative indices used here, n >> log2(d) == n / d for every input — same quotient, same scale word fetched, same arithmetic after it. Dispatching on POW2_SCALE can never change what an A/B leg measures, only how fast the scalar arm runs. The C ABI (launch_fp8_matmul_{f16,bf16}) is unchanged; no Rust edits needed. Compile-unverified locally (macOS has no nvcc); covered by the nvcc CI lane (.github/workflows/cuda_compile_check.yaml), which builds this TU for sm_80/sm_90 on every PR touching it. The bank-conflict half of this lane's brief (+4 -> +1 tile padding) was already landed by PR #215 and is an ancestor of this branch; the kernel's stride-33 comment documents the bank arithmetic. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * feat(ArcLab): ARC_LOG_FP8_DISPATCH — engagement lines proving WHICH FP8 kernel served a forward Every threshold sweep on this lane needs an engagement assertion: a server-log line proving which kernel actually ran, per leg. None existed — a leg could silently measure the wrong arm (the ARC_NO_DEDICATED_DECODE incident produced two void A/Bs exactly this way, both arms running the same code). ARC_LOG_FP8_DISPATCH=1 (read by VALUE via env_flag_is_set, latched once per process) makes each dispatch path print ONE line per process: [arc-fp8-dispatch] path=<gemv_wide|gemv_warp|wmma|tiled|dequant_cublaslt> first_shape=mM_nN_kK Sites: the native dispatch chain in fp8_blockwise_matmul_impl (both dtype arms, mirroring the if/else exactly) and the ARC_FP8_CUBLAS_MIN_M divert to dequantize+cuBLASLt in BlockwiseFP8Linear::forward. Cost when unset: one latched bool read. When set: one relaxed atomic swap per call — no locks, so enabling it does not perturb the rates a sweep records. Registered in mistralrs-core/tests/capability_reachability.rs (Status::Live); the registry run passes (11/11). The Rust is cfg(cuda) — compile-unverified locally on macOS; covered by the cuda-typecheck job of the nvcc CI lane. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * feat(ArcLab): three-arm FP8 threshold sweep harness — the measurement allowed to move ARC_FP8_CUBLAS_MIN_M arc_fp8_cublas_min_m's doc and the WMMA kernel's doc both forbid merging a threshold from a two-arm result: the only sweep ever taken predates both the WMMA GEMM (#200) and the rank-3 flatten, so it compared cuBLASLt-as-batched- GEMV against the scalar kernel only. This script is the three-arm re-run they demand, on ONE binary, env toggles only, fresh server per leg (the gates are OnceLock-latched, so re-exporting into a live server would silently measure the previous leg). Design: the threshold candidates {5, 8, 64, 256, 512} are swept as the OFFERED DECODE BATCH M (a grid that varies only the env at one fixed batch routes every MIN_M <= batch leg to the same kernel and cannot name a crossover in M), and each ladder point runs all three arms: (a) cublaslt: ARC_FP8_CUBLAS_MIN_M=5 (b) wmma: ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=0 (c) tiled: ARC_FP8_CUBLAS_MIN_M=1000000, ARC_NO_FP8_WMMA=1 A leg counts only if ALL of: 1. ENGAGEMENT — its expected [arc-fp8-dispatch] path= line present in the server log and both rivals ABSENT (a WMMA leg whose log shows tiled is eligibility silently failing: VOID, loudly); 2. FLOOR — summed usage.completion_tokens >= MIN_TOKENS_FLOOR (default 1000); the rate divides by tokens the server SAYS it generated, never by requested max_tokens; 3. CANARY — greedy fixed-prompt stream vs the baseline leg (tiled @ M=5); cross-arm divergence is expected numerics and is REPORTED with its first index; a degenerate stream (<5 tokens / no finish_reason) voids the leg. The summary table names, per M, cublaslt vs min(wmma, tiled) and the FIRST clean M where cuBLASLt wins — the value ARC_FP8_CUBLAS_MIN_M may move to. No clean crossover => the default stays 512. Script only — NOT run here (no GPU; D14). House lock/preflight/provenance discipline copied from arcspec_token_identity_b8.sh. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * ci(ArcGate): free SASS probe for the FP8 dense-GEMM lane — LDS width census + WMMA HMMA presence Two questions this settles with nvcc + cuobjdump alone, no GPU: 1. Whether the 4-way bank-conflict cost model for fp8_matmul_tiled describes the compiled artifact: the model is a claim about SCALAR 32-bit LDS in the inner product, and if ptxas had vectorized those shared loads to LDS.128 its arithmetic would be about instructions that do not exist. The probe prints a per-kernel LDS/LDS.64/LDS.128 census. (Note: the +1 pad that kills the conflict also forecloses LDS.128 on tile rows — scalar LDS is the EXPECTED reading on current source.) 2. Whether the never-executed tensor-core GEMM (blockwise_fp8_gemm_wmma.cu) actually contains tensor-core instructions: HARD FAIL if the TU has zero HMMA/HGMMA — the three-arm sweep would otherwise compare cuBLASLt against two scalar kernels while calling one of them WMMA. Teeth: exactly 4 fp8_matmul_tiled kernels asserted (2 dtypes x 2 POW2_SCALE instantiations — a dropped instantiation makes the census about a kernel not in the build), and every tiled kernel must show >=1 LDS or the extractor lost the function body. Wired into the nvcc CI lane as a step beside the QTIP spill gate, so it runs for sm_80 and sm_90 on every Rust/kernel PR; also runnable standalone: arc-tools/fp8_gemm_sass_check.sh sm_90 Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcGate): SASS probe — replace GNU-awk-unsupported \< \> ERE escapes with cuobjdump -fun extraction The probe's first CI run went red on itself, honestly: GNU awk treats \< \> (and warns on \.) as plain characters in ERE, so the awk body-splitter's counting regexes matched NOTHING, the per-function census returned zero LDS for all 4 fp8_matmul_tiled instantiations, and the fail-on-zero guard refused to trust its own zero — which is exactly what it exists for. This revision removes regex body-splitting entirely: the mangled kernel names come from the Function listing (awk '{print $NF}', no regex), and `cuobjdump -sass -fun <mangled>` extracts each function's SASS directly. All counting is grep -cE with POSIX classes and bracket expressions only — no \< \> anywhere. Also, per review: * The two failure modes are now DISTINGUISHED: FAIL[extractor] (no instruction lines came back — the census is unverified) vs FAIL[zero-LDS] (body extracted, genuinely no shared loads — the conflict model has no subject). Body presence is judged on /*<hex>*/ instruction lines, so the two cannot be confused. * The VERDICT[tiled] line prints ONLY when the per-function census is complete and non-empty; otherwise it prints NOT ESTABLISHED. A TU-level cross-check line (whole cubin, includes the GEMV/MoE kernels) is always printed, labeled informational. * HMMA counting moved off the \< \> GNUism onto the same POSIX pattern. Verified locally against a stub cuobjdump with a realistic SASS fixture: happy path (4 kernels, LDSM excluded, LDS/LDS.64/LDS.128 counted per function, scalar arithmetic correct, exit 0) plus three negative controls (broken -fun -> FAIL[extractor] + verdict NOT ESTABLISHED; LDS-free body -> FAIL[zero-LDS]; HMMA-stripped WMMA TU -> hard fail), all exit 1. The real cuobjdump run is CI's — if -fun ever returns nothing there, the probe fails loudly rather than lying. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…s no longer outlive their sequence (#220) * fix(ArcInfer/ArcGraph): parked burst tokens no longer outlive their sequence — one funnel, one leak test A device-loop burst runs with eos_token_id = -1, so it always parks its full burst; a sequence stopping mid-burst (stop string, max_tokens, cancel, error, panic, eviction) left its undrained tokens parked, and the NEXT sequence's first sample_sequence — including the one after its prefill — consumed them as its own. One user's tokens in another user's response. clear_pending_tokens documented the contract but its only production call sites were fallback paths. The fix is one call at the one funnel every completion path passes through: Sequence::set_state stands the device loop down on every transition out of the running set (Done, Error, FinishedAborted, FinishedIgnored, Waiting, Swapped). Clearing on preemption is deliberate fail-closed: the sequence re-prefills on re-admission, so its parked tokens are stale. To make the funnel testable on machines without CUDA — the cfg(cuda) darkness is exactly what let this ship — arc-cuda-graph is now a non-optional dependency of mistralrs-core (its host surface is plain Rust; CUDA plumbing stays behind its own cuda feature, still enabled by mistralrs-core/cuda) and device_loop_pre_sampled_token loses its cfg split. The new leak test drives the real sample_sequence path both ways: wired, B's first token is its own argmax and the queue is empty; with the funnel removed, the queue holds 3 of A's tokens and B's first sampled token is 22 — A's parked id. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcInfer/ArcGraph): owner-tag the pending queue — the leak becomes structurally impossible, not funnel-dependent The completion funnel cannot see a sequence the DefaultScheduler moves aside WITHOUT a state change (the bucketing waitlist says so in its own doc), and batch growth has the same shape. So the queue itself now refuses: tokens are parked tagged with the owning sequence's id (published by the sampling pre-hook with eligibility), take requires a matching id, and a foreign queue is dropped whole with the grep-able instrument line "ArcGraph: dropped N foreign parked tokens". The forward's pending>0 short-circuit gains the two guards its aliasing contract always needed: seq_len == 1 (a parked queue at a prefill step must not replace the prefill) and pending_owned_by_current(). If the short-circuit fired and the sample still cannot take (scheduler swapped sequences in between), the step fails loudly via the aliased-logits marker instead of host-sampling another sequence's stale graph output. DeviceDecodeLoop finally resets: every stand_down (including the funnel's) bumps a device-state generation; run() compares and zeroes ring, cursors and fault word before a newer-generation burst. Reset at engagement, not completion, because the funnel cannot reach the pipeline-owned loop and between bursts is the only point with no in-flight ring writes. ARC_GRAPH_DEVICE_LOOP is now registered in capability_reachability.rs (Symbol: device_loop_enabled), closing the refuted premise from the previous commit. Interleave test (A waitlisted, no set_state): wired, B samples its own argmax 9 and the foreign tail is dropped; with the owner comparison removed, B's first sampled token is 22 — A's parked id. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…lips — GEMV_WIDE ON, QK_FUSED_COHORT ON, CUBLAS_MIN_M 512→64 (stacked: B=256 575→758 tok/s, b=1 40.8→43.8) (#223) * perf(ArcKernels): ARC_FP8_GEMV_WIDE default ON — measured +7.5% b=1 (43.80 vs 40.76 tok/s) Session-10 hardware A/B, box arc-s10-ledger (H200), candidate 4d03b9e, leg 3: 3x512 produced tokens, seeded t=0.7 p=0.95, engagement proven by [arc-fp8-dispatch] path=gemv_wide; canary1 byte-identical, canary2 coherent and arithmetically correct (numerics within the stated f32 re-association bound). The ledger's ~1.6x prediction did NOT materialise end-to-end; the default stands on the measured +7.5%. Polarity flips from 'only "1" enables' to 'only "0" disables' — still value-read, never presence-read (#212). Test renamed accordingly (only_literal_zero_disables) with the assertion values swapped. Registry entry added (capability_reachability.rs, ArcKernels). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * perf(ArcAttention): ARC_QK_FUSED_COHORT default ON — engaged + bit-identical at B=256 (599.35 vs 575.12 tok/s) Session-10 hardware A/B, box arc-s10-ledger (H200), candidate 4d03b9e, leg 4: fused qk_norm_rope ENGAGED at B=256 (the default arm DECLINED and took the eager chain), seeded canaries byte-identical, 0/512 request errors, aggregate 599.35 vs 575.12 tok/s. The +4.2% is inside the ~±5% inter-leg variance measured the same session (leg 4b no-op delta), so the flip stands on engaged + bit-identical + not-slower, exactly the bar the capability entry set. Doc comment and registry entry updated as the old comment promised; decline message now names the kill switch (ARC_QK_FUSED_COHORT=0). Polarity stays value-read (#212). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * perf(ArcKernels/ArcQuant): ARC_FP8_CUBLAS_MIN_M 512 -> 64 — the three-arm sweep this constant's own doc demanded Session-10, box arc-s10-ledger (H200), candidate 4d03b9e, arc-tools/fp8_threshold_sweep.sh, LADDER {8,64,256}, one binary, flatten fix present, >=1000-token floor. Clean rows (tok/s): M=64: cublaslt 552.4 | wmma 328.2 | tiled 223.1 -> cuBLASLt wins M=256: cublaslt 715.9 | wmma 620.4 | tiled 285.8 -> cuBLASLt wins First clean M where dequant+cuBLASLt beats min(WMMA, tiled): 64. M=8 rows were VOID on the token floor, so 5..63 stays with the native arms — a void row makes no claim. fp8_gemv_warp still owns M<=4 (the M=1 floor row, cuBLASLt 0.72x, still stands). Rows quoted in the doc per its own rule: a threshold is a claim about every value it excludes. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * perf(ArcKV): ARC_V4_FP8_KV default ON — the U8 CUDA leg finally measured: engaged, byte-identical, not slower Session-10, box arc-s10-ledger (H200), candidate 4d03b9e. Leg 10: seeded canaries byte-identical to the BF16 arm, b=1 41.49 vs 40.76 tok/s. The lane has no engagement log line (silent-success class, noted for follow-up), so engagement was proven by nsys kernel capture instead: arc_kv_fp8_quantize_kernel x16,297 + arc_kv_fp8_dequantize_kernel x32,594 in one b=1 run. 4-gate stacked verification (this + the three earlier flips, one binary): B=256 744.7 tok/s (vs 757.5 three-gate, inside the ±5% band) with canaries IDENTICAL through the stack; b=1 44.68 tok/s — the session's best b=1. Unlike wave43-BU (same default, shipped unrun, broke every forward), this default stands on a measurement. ARC_V4_FP8_KV=0 restores BF16; the ARC_V4_CAPTURE_PROBE veto is unconditional and now tested for the default arm too (write_kv_inplace has no U8 variant). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
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.
Arc → OpenRouter-ready: the single integration PR
This is that PR.
release/openrouter-readyis the integration branch, and this is the only pull request open againstmaster. Every other open PR targets the integration branch and merges into it; when the branch is ready, this one merge decision ships all of it.How to work in this model
release/openrouter-ready, and PR back into it. Not offmaster.master. If you find yourself opening a PR againstmaster, that is the mistake.base-branchCI lane enforces exactly this: it acceptsmasterorrelease/openrouter-readyand refuses everything else (ci(ArcGate): let the base-branch lane permit the single integration branch — 7 PRs are red for addressing, not quality #216). It is an exact-match allowlist of two, deliberately — see the comment on the lane for why it must not grow.Master stays protected
Widening that lane did not weaken
master.release/openrouter-readyreachesmasterthrough exactly one pull request — this one — and this PR's base ismaster, so it takes themasterarm of the same check every other PR does.CI completeis still the single required status check in branch protection and still aggregatesbase-branchin itsneeds. Nothing entersmasterwithout a full run whose base ismaster.This branch was re-cut, because merging the old one would have caused a regression
The previous
release/openrouter-readyhad diverged frommasterby ~40 commits in master's favour. Merging it as it stood would have reverted TCFRAG (#203) and re-broken the flag-polarity work (#212).It has been force-updated to current
masterplus only the work it uniquely carried, so the diff below is genuinely additive.What this branch carries today
Everything on
master, plus five commits:docs: the single OpenRouter-readiness gatedocs/engineering/OPENROUTER_READY.md— 16 items across blocking / product / trust / commercial, each with measured provenancedocs: the gate is 14,000 tok/s aggregate, and the scope is Arc completetest(ArcLab): batch sweeparc-tools/batch_sweep.py— fires B concurrent streaming requests and reports per-request vs aggregate with achieved concurrency, so the two anchors can be read together, plus memory-controller busy % sampled during decodefix(ArcQuant): the MoE gate never selected the grouped GEMM in servingfix(ArcQuant): same 683-token gate on the LUT rungqtip2, which is the rung the measurements actually exercisedThe keystone: the grouped GEMM was unreachable in real serving
Both QTIP rungs gated the trellis grouped GEMM on tile occupancy — switch only once each woken expert draws
GROUPED_TILE_M(=16) pairs. At V4's top-6-of-256 routing that needs6n/256 >= 16, i.e. n ≥ 683 tokens, which decode never reaches. The clause pinned every decode step, and every batch up to 682, on the per-pair gather-GEMV arm.Tile fill was the wrong quantity. At n=16 the tile is 7.5% full and the grouped kernel still wins by 16%. The real mechanism is that the grouped kernel stages a woken expert's bytes once per m-tile instead of once per (token, expert) pair — the GEMV arm's reads are exactly linear in
n_tokens, 12.00× redundant at n=512, while the grouped arm holds flat at one pass over the expert set — and that it runs on tensor cores.DECODE_REGIME_MAX_TOKENSsurvives as a floor (b=1 measures 0.47×, and it is the RUN-161 capture-safety rule).ARC_QTIP_ONDEVICE_MOE_MAX_TOKENSstill pins the GEMV arm for A/B, andARC_NO_QTIP_GROUPED_MOEremains the kill switch.The decision moved into
gather_policy::prefer_gather_gemvspecifically so it can be unit-tested — the dispatch site is inside#[cfg(feature = "cuda")], where a non-CUDAcargo testwould otherwise report green over code it never compiled. The tests pin every measured batch size to the grouped arm and carry a control asserting the retired clause would have excluded all of them.Carried forward without regression
The gate retraction was authored (#197) before #212 landed, and its original diff would have reverted
ARC_NO_QTIP_ONDEVICE_MOEfrom #212's by-value parser back to a presence check. That was resolved in #212's favour: the cherry-pick keepscrate::env_flag_is_set(...), soARC_NO_QTIP_ONDEVICE_MOE=0still means off.Verified on the re-cut branch:
tcfrag2b.rsstill selects onvalue == Some("1").=0means off #212'senv_flag_is_setconversions are intact — the count is unchanged frommasterin every touched file, and the only env-flag line in the whole diff is an explanatory comment.mistralrs-quant/src/qtip/tcfrag2b.rsandmistralrs-quant/src/env_flag.rsare not touched at all by this branch.cargo check -p mistralrs-quantclean; all 12gather_policytests pass, includingthe_grouped_gemm_is_reachable_at_decode_batch_sizesand thethe_retired_tile_fill_clause_would_have_excluded_all_of_themcontrol.Branch facts
8ae209008— currentmaster, which already carries #212, #203/TCFRAG, #205, #213, #214, #215 and the CI rule change #216d665956abarc-tools/batch_sweep.py,docs/engineering/OPENROUTER_READY.md,qtip/bitshift.rs,qtip/gather_policy.rs,qtip/mod.rsRegression proofs on the re-cut branch (
git merge-base --is-ancestor <c> HEAD→ 0 for each):cc5487ad3— TCFRAG-2B (perf(ArcQuant/ArcKernels): TCFRAG-2B — put the qtip2b trellis GEMV on the tensor cores #203) ✅ present9c127a2b1— TCFRAG made opt-in (fix(ArcQuant): make ARC_QTIP_TCFRAG opt-in — an unverified kernel is default-on for b=1 decode #209) ✅ present6ffdac7ae— the 23ARC_*flag-polarity fixes (fix(ArcGate): read boolean ARC_* flags by value, so=0means off #212) ✅ presentmistralrs-quant/src/qtip/tcfrag2b.rsandmistralrs-quant/src/env_flag.rsare not in the touched-file list at all, so neither can have regressed.The rest of the queue
Every other open PR now targets this branch. Four of them were stacks whose bottom PR was their base; they are now siblings, so merge order is load-bearing and no longer enforced by the base ref — each carries a note saying whether it is the bottom or the top:
V4CachedKvariant, + honourqs/softcapping#98Two are red for real defects rather than for addressing, and are labelled
needs-work: #184 (genuine test failures) and #108 (a real fat-binary failure). Do not read a green base check on those as readiness.