Repository navigation
ArcGraph: three replay-correctness blockers fixed; capture now blocked on host heap corruption, not on V4 - #181
Merged
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
|
CI fixed ( What I fixed. The only red check was On the stale title. This PR says capture is 'now blocked on host heap corruption, not on V4'. That blocker has since been root-caused and fixed — dangling host pointers in 7,077 capture-time H2D copies, addressed in candle via I have not merged it because I have not audited its blast radius against the current capture path, and the ArcGraph area has live in-flight work on top of it ( |
heydryft
added a commit
that referenced
this pull request
Aug 20, 2026
Same root cause as #181, and the identical fix: `.typos.toml` already exempted the bare Rust identifier `arange` (candle's/numpy's `Tensor::arange`) under [type.rust.extend-identifiers], but that only matches WHOLE identifiers — the word inside `matches_the_arange_expression_it_replaced` still tripped the gate. Add the word-level exemption rather than renaming, because `arange` is a real API name, not a misspelling of 'arrange'. This branch and #181 both carry those test functions, so both hit it. Whichever lands first makes the other's copy redundant. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Contributor
Author
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 |
…d refuse to instantiate on a miss
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.
heydryft
force-pushed
the
arcgraph/replay-correctness
branch
from
August 21, 2026 16:27
6424f40 to
9c36af9
Compare
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
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>
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
… 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>
Contributor
Author
Rebased onto
|
| commit | outcome |
|---|---|
9795d15a6 geometry invariant · 09b3b82d5 BLOCKER 3 · 9d1d07c02 BLOCKER 2a · f49b991cf ring-slot remainder · 5c7f9677f measurement runner · a35bcec0d mark the five driver calls · 09469fee1 stop swallowing the warm-pass sync |
already on the integration branch (via #205, fdf0b6c19) — skipped by git rebase itself |
2743a4400, fb58fc63c (the two arange commits) |
dropped — same content as #190, already landed as c1d2b27f4 / 64fe1d678, and the positions_f32 helper was deliberately deleted by 6fcb87f47. See #190 for the full argument. |
6424f40f7 .typos.toml widening |
dropped — 6fcb87f47 chose renaming over widening on purpose, citing 745c871dd / e4eb59dfb, and no identifier in the tree now contains arange inside a longer snake_case name, so the entry is dead. |
492b7baa7 pre-warm the capture from a measured alloc profile |
kept — this is the residual. |
The candle pin would have reverted three fixes
This is worth reading even if nothing else here is.
492b7baa7 bumps the fork pin to 1fa534db for prewarm_alloc_cache,
reset_capture_misses and capture_misses. That was right when it was written.
It is wrong now: the integration branch has since moved to 89ab14ef, and the
two revs have diverged.
$ gh api repos/aeonmindai/candle/compare/1fa534db...89ab14ef
{"status":"diverged","ahead":3,"behind":1,"commits":[
"perf(cuda): bound the caching allocator — it retained every buffer and freed none",
"test(cuda): assert the cap's real contract, not an idealised one",
"fix(cuda): eviction must not run inside the capture window"]}
So pinning at 1fa534db would have silently reverted the bounded caching
allocator — which #213 (bound the alloc cache against real headroom) rests on
— plus a capture-window eviction fix. Green build, no conflict, no commit to
blame.
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. All four required pub fns read out of
candle-core/src/cuda_backend/device.rs at that rev before relying on it:
drain_alloc_cache_and_free, reset_capture_misses, capture_misses,
prewarm_alloc_cache.
Also
normal.rs's same-step-probe block is re-indented +4. That is the exact
counterpart of the de-indent 6fcb87f47 applied on the integration branch,
which it did because this commit's if !misses.is_empty() nesting level was
not yet present. It is now.
Verification
cargo check --workspacegreen on the new candle pin;cargo check -p arc-cuda-graph --all-targetsgreen.cargo test -p mistralrs-core --lib: 693 passed, 0 failed. 707 on the merged#198 + #181 + #130 + #121tree.mistralrs-corewarning count unchanged at 23.typosclean.- TCFRAG still opt-in (
tcfrag2b.rs:273,value == Some("1"));env_flag_is_setper-file counts identical to the integration branch.
The pre-warm behaviour itself is not verified — it needs a CUDA box, and none was used.
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
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>
Contributor
Author
📌 DO NOT DOWNGRADE THE CANDLE PIN —
|
| compared against | result |
|---|---|
89ab14ef (integration branch's pin) |
ahead 6, behind 0 |
1fa534db (this branch's original pin) |
ahead 8, behind 0 |
All four required entry points read directly out of
candle-core/src/cuda_backend/device.rs at that rev before relying on it:
drain_alloc_cache_and_free (:879), reset_capture_misses (:998),
capture_misses (:1004), prewarm_alloc_cache (:1054).
Rule for anyone rebasing this branch: the pin may move forward to a rev that
contains 9586979db. It may never move to one that does not. The Cargo.toml
comment states this inline, with the gh api compare command to re-verify it.
CI has confirmed the bump: cargo check (cuda, workspace), nvcc compile (sm_80) and nvcc compile (sm_90) all passed on this pin.
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…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>
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
… 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>
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…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>
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.
Stacked on #179 (fixed launch geometry) and
agent/graph-static-input(dangling-host-pointerarange), whose overlapping fixes are resolved here rather than picked between.What was wrong, and what this changes
Blocker 1 — RoPE resolved its position on the host.
DeepSeekV2RotaryEmbedding::forwarddidcos.narrow(0, seqlen_offsets[0], seq_len). Capture folds that offset into the recorded kernel's arguments, so every replay rotated at the capture step's position — the near-constantmax|Δ| ≈ 30seen in three earlier runs.cos_sin_forgathers from the device position buffer under graph mode and narrows otherwise.index_selectof[p]isnarrow(0, p, 1), so this changes when the offset is resolved, not what is computed. Applied toforward_inverse_tailtoo — it de-rotates the MLA output and would otherwise have frozen whileforwardmoved.Blocker 3 — the graph KV window was a prefix, not a ring.
append_graphread a fixed0..read_capacitywindow but wrote at the absolute position, so from tokenread_capacityon, the new row landed outside the window that is read.read_capacityissliding_window, so a ring of that many slots holds exactly the key set the raw branch is defined to attend. Permuting keys is invariant here: softmax over keys is permutation-invariant, K and V go to the same slot, and V4 rotates K before caching so each row carries its own absolute RoPE. The validity mask needed no change —slot <= positionis already ring-correct in both regimes.Blocker 2a — the retained raw tail was an unpinned width. This turned out to be the live defect, and it is now measured.
advance_uniformretains only what a future compressed row can consume, a function oftokens % ratio, givingratiodistinct widths on a period ofratio.pin_tail_widthholds it atspan_groups * ratio + margin, an upper bound on what the unpinned policy would keep.Measured
cuGraphLaunch→ 700ILLEGAL_ADDRESS327 680 B at hidden 4096 in BF16 is 8192 B/token = 40 tokens — the raw tail, one token wider than the 319 488 B (39-token) miss on the preceding warmup step. That is the ratio-128 (HCA) tail, which grows one token per step for 128 steps, which is why no practical warmup count ever covered it. (The ratio-4 tail has only 4 distinct widths — the same defect once recorded as "four unwarmed allocation sizes".)
What still blocks replay
Not V4, and not the graph geometry. With misses at zero the forward records cleanly and then dies in
cuGraphInstantiateWithFlagswith glibccorrupted size vs. prev_size— host heap corruption, not a CUDA error. The ABI is correct against the box'scuda-12.4header. Capture requiresARC_CANDLE_ALLOC_CACHE, whose deferred-free bookkeeping is the component already known to leak (132 allocs/step, zero frees) and is being repaired on a separate branch;ARC_ALLOC_CACHE_MAX_MBis not in this tree. Retest capture once that lands.Blocker 2b is architectural, not mechanical, and is documented at the call site.
cuGraphLaunchexecutes only recorded kernels — no host code runs — and the V4 compressor history lives in host-owned Rust state (XsRollingCache::tail, reassigned to a freshTensoreach step), so under replay it stops advancing. The raw KV half escapes this only because it writes through a recorded kernel into one fixed-address buffer. Separately, the compressor fires on 1-in-ratiosteps, so consecutive decode steps do not execute the same kernel set at all. Making the compressed branch replay-correct needstailto become a device ring advanced by a recorded kernel, plus either a conditional graph node or one graph per phase.Scope
Default OFF; every change is inside the graph-decode arm, gated behind
ARC_V4_CAPTURE_PROBE+ARC_CANDLE_ALLOC_CACHE.pin_taildefaults false, so the shipping path is untouched — all 16xs_rollingunit tests pass unchanged, and the capture-off leg above is clean.Known limitation, named rather than fixed: the ring assumes every write since position 0 used
position % capacity, which holds only while the prompt is ≤sliding_window. A longer prompt needs a one-time seeding pass at graph-mode entry.🤖 Generated with Claude Code
https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H