perf(v4): pin the xs window (+23.2% at uniform B=32) + the ladder that prices per-sequence advance — 4.46x at spread B=8 [MEASURED] - #121
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 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
Owning the +572 ms/step rebuild overheadCorrection first, since it was relayed elsewhere: Where the cost actually is
let batch_k = Tensor::zeros(dims_k, ...)?; // [B, .., capacity, head_dim]
let batch_v = Tensor::zeros(dims_v, ...)?;
for (i, src) in srcs.iter().enumerate() {
batch_k.slice_set(&pad_slack(&src.k, ..)?, 0, i * one_k[0])?;
batch_v.slice_set(&pad_slack(&src.v, ..)?, 0, i * one_v[0])?;
}The allocation is sized by capacity, the work is sized by tokens Note what the zeroing buys: nothing, in the common case. Every row is This is the same lever the window pin already pulled. The pin made every The changeReuse the batched buffers across steps instead of reallocating, and write only Two constraints this must not break, both established by measurement
Acceptance, measured in isolation rather than in a sweepTwo cells, one server per arm, repeated against a single server — not a ladder.
Both arms must hold |
|
ArcGate: HELD — parked above #104, and the reason is now specific rather than open-ended. Full analysis on #104. Short form: #104 and #122 (landed as master
#104's Two things that are not true, both checked rather than assumed:
This is scope, not a verdict. Nothing here is ranked down or closeable; the blocker is now one A/B measurement — does #104's design buy anything #122's does not, priced against its per-step rebuild. Also note: this PR currently shows 1 check, the comment bot — zero CI lanes have ever run on it, because it predates the stacked-PR trigger fix that is now on master. Any push starts the full 16 lanes. Nothing deleted; branch untouched. |
|
Separability checked, and the answer is no — recording it so nobody re-derives it. I lifted #116 onto master cleanly (now #128) and tried to do the same for this PR's measured +23.2% window pin. It does not separate. Two independent blockers, both verified by cherry-picking into a scratch worktree, not by reading: 1. let ragged_len = ragged_bucket_len(&running, &self.admission);
2. The pin itself is not independent of #116 either. This is not a verdict on the work. The pin is the only measured throughput win in the parked set — What unparks it: either #104's cadence A/B resolves (see my writeup there), or its owner deliberately rebases the #116 → #121 sequence onto master, dropping Nothing rebased, force-pushed, or deleted here. |
4090dd6 to
4c7b376
Compare
|
Rebased onto 🔴 The rebase surfaced a real collision: 3 master tests failCause: The failures are the pin doing exactly what it documents — Not resolved here, because both possible fixes are the author's call, not a rebase decision:
Conflict resolved during the rebase
|
Rebased onto master and retargeted — but held, not landed: it defaults ON and that breaks three tests on master.This was one of five PRs based on another PR's branch ( Rebasing is what made the defect visible — that is the whole point of the audit. The defect
let on = !matches!(
std::env::var("ARC_V4_XS_PIN_WINDOW").as_deref(),
Ok("0") | Ok("false") | Ok("off") | Ok("no")
);Unset does not match, so The old stacked base made this structurally invisible: those tests were not on the branch's ancestor, and no lane ever ran. Why it is not just flipped here
What is neededEither (a) default it off and update the ladder harness to pass Standing constraint for this session is do not flip defaults, and unproven-on-GPU behaviour stays off — so (a) is the smaller and safer change unless you know otherwise. |
|
Triage verdict: REBASE, not close. Freshest of the stale tail (58 commits behind) and still unique. Evidence against current master: Conflict: 1 file, Not merging now, deliberately: the headline numbers here (+23.2% at uniform B=32, 4.46x at spread B=8) were measured on this branch's base, not on today's master, and there is no GPU to re-measure on. The code may well be right; the numbers are stale and must not be carried forward as still-true. Land it when a box exists and the ladder can be re-run. |
Retargeted at the integration branchBase changed: The queue is being restructured to the shape the owner asked for: one PR open against Two things had to land on
This PR was not closed and is not considered stale. An audit of the queue found the overwhelming majority of it to be real work that was never merged, not noise. What you need to do: rebase onto |
…onfound 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>
…ch 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>
… 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>
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>
…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>
…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>
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>
4c7b376 to
2dfed66
Compare
… 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>
Rebased onto
|
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>
2dfed66 to
27c3fbf
Compare
…ps 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>
⚡ SWITCHED OFF ON PURPOSE — and here is what turns it back onThe default is retracted, not the change. Off is a temporary state with an
Recorded in three places, one of them enforced: the gate's own doc comment under Why the default is retractedThe pin changes what serving retains — 144 columns at V4's HCA geometry, The correctness argument this PR made for the pin is sound and kept: the The harness inverted with the default — three things had to move togetherThe pin used to be the default, so the treatment was the default arm and the
Polarity is A third test pins what this retraction was most likely to break: CUDA-graph |
… 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>
…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>
Stacked on #116. Two things: the throughput ladder that prices #116's fix, and
the window pin that came out of measuring it.
The headline
Per-sequence KV advance on realistic ragged-length traffic:
run_bucketms/bstep29.1317.2-> 43.61453-> 5001At B=8 with unequal prompt lengths, batching was negative before this stack
(11.17 tok/s against B=1's 20.82). It is now 2.4x better than B=1. That is the
case real serving is made of, and it is the case every uniform-prompt benchmark
we hold was blind to.
The other half, because a ratio alone is unreadable
tok/steprises where aggregate falls: uniform B=32 istok/step+6.0% andaggregate -20.4%, because
ms_per_batch_stepwent 715 -> 1287. Reportingone without the other is how a mechanism gets called a win or a regression when
it is neither.
The regression is located, and it is not per-sequence advance being
expensive. At identical batch width — uniform B=32,
running_bucket_size32.0 in both arms, nothing about the batch shape differing — the mechanism
costs +572 ms/step of pure overhead. And it is superlinear: +28 ms at B=8,
+572 ms at B=32 — 4x the batch, 20x the cost. (The spread B=128 pair is VOID — see below — so the earlier claim that
superlinearity turns its 2.5x-wider batch into a 0.81x result is withdrawn.
The superlinearity itself rests on the uniform B=8 and B=32 pairs, which are
clean in both arms.)
The cause is
issues_cache_in = !no_kv_cache && (cohort_changed || granted).The
granteddisjunct forcesCacheInstruction::Inevery step rather thanon cohort change, so the whole batched cache — 43 K/V slots plus 41
XsRollingslots — is rebuilt perpetually.
What
grantedis protecting is real and should not simply be removed: underper-sequence advance the rows' lengths genuinely diverge every step, so the
front-alignment moves, and
clone_in_cacheis the only producer of theper-row
lead_padthatresolve_row_cache_lensconsumes — without itper_seq_advanceis false on every step and the mechanism goes silently inert.The waste is not that the rebuild has no work; it is that the rebuild is
O(B x capacity)per slot when the geometry delta isO(B x tokens-per-step).Making that precise is the next change, and it is worth roughly 3x at spread
B=128 rather than the 28.3 ms/step the profiler priced elsewhere.
The window pin, priced
ONvsON_UNPINNEDdiffer in one flag, and the control was asserted to havegenuinely taken the other branch (
RESIZINGin its log,PINNEDin the others— read from the runtime log, never the ELF, since both strings are compiled in
unconditionally).
ms/bstep2371 -> 1287)Neutral everywhere, +23.2% at uniform B=32, worst case 0.983x. Restoring an
invariant the type already documents is free at worst and a quarter of
throughput at best — and one buffer that stopped reallocating gave back
1084 ms/step, which is the same lever the
clone_in_cachework pulls again.What is in here
scheduler::bucket_telemetry—SCHED[agg]carryingbuckets_per_step,running_bucket_size,offered_per_step, emitted on the same log fence asMTP[agg]so a cell's scheduler window and MTP window are identical. Withoutit, "aggregate did not move" cannot be told from "the batch never existed".
It confirms the bucketing law a fifth time: predicted
B / distinct lengthsvs measured
running_bucket_size— spread 1/1.06, 4/4.27, 16/17.2; uniform8/8.0, 32/32.0.
ARC_V4_XS_PIN_WINDOW, default on) and its proof — see thecommit message for the bound
W <= span_groups * ratio + margin - 1andpinning_the_window_is_numerically_inert, which requires exact equality ofcompressed rows and both time bases across 40 steps and asserts the unpinned
widths actually varied so the equality is not vacuous.
arc-tools/arcspec_perseq_ladder.sh— 3 arms x 2 regimes x B in {1,8,32,128},one server per arm, counters differenced across each cell's own fence.
What the harness had to learn, on the way
The first OFF arm produced a full set of plausible cells that were all wrong,
and the fixes are the reason the numbers above can be read:
requests carry
max_tokens 4096and keep running, sospread B=1wasmeasured with 81.2 sequences/step offered to the scheduler. Cells were
measuring their predecessors. Now drained on the engine's own cumulative
counter before the next fence opens.
uniform B=128started 128requests, hit 0 errors and produced 0 tokens in 65 s — the whole cell was
prefill, reported as
0.0 tok/s. The ramp is now adaptive: it ends when thebatch is actually producing, and
--warmupis a cap that voids the cell ifhit.
UNREADABLEverdictin three states — pass, fail, and cannot-answer — where cannot-answer includes
a real aggregate with absent engine counters, which is exactly what
uniform B=128 OFFwas. The report printsVOID, not a number.Two OFF cells are VOID, and the ON arm is clean throughout
Every cell records
offered_per_stepagainst the B it asked for. Re-verdicted ata 10% tolerance (the in-run guard used 1.25x, which was far too slack):
1.000x. The pincomparison is therefore between sixteen cells that were each handed precisely
the batch they asked for.
OFF.spread.k128was offered 151.95 sequences/step for B=128 (1.187x) —residual drain from the previous cell, because a shattered OFF batch drains far
more slowly than the counter-based detector assumed. VOID, and with it the
0.81x regression claim.
OFF.uniform.k128has a real aggregate and no engine counters. VOID —cannot-answer, and the 1.491x it implied does not stand.
OFF.spread.k32reads 1.045x — inside tolerance, noted.What this leaves untouched: the +572 ms/step at identical batch width is
uniform B=32, where BOTH arms read
offered 32.0andrun_bucket 32.0atexactly
1.000x. The pin's +23.2% is ON vs ON_UNPINNED with the sameper-sequence flags, both exactly
1.000x. Anything under ~10% — uniform B=8's0.954x, and the pin's B=1/B=8/B=128 cells — should be held until re-measured in
isolation, since a cross-arm null does not cover within-sweep order effects.
Engagement, asserted rather than assumed:
per_seq_stepsON=1459, OFF=0;per-sequence KV advance is ONandRagged batch admission is ONpresent in ONand absent in OFF; zero
xs rolling cacheerrors in every arm; the run wasexclusive on the card end to end and would have been discarded, not salvaged,
had a foreign compute app appeared.
Measured on H200
arc-graph-probe, ref4090dd62f, every arm's server loggingthat revision.