perf(ArcAttention): V4 RoPE was launched per-sequence — 99,072 launches/step -> 301 - #198
Merged
heydryft merged 6 commits intoAug 21, 2026
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 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
This was referenced Aug 21, 2026
Merged
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
…lt-on), 3 capacity, 1 dead loader (#204) * fix(ArcModels): Gemma reported a 4096 context because it read the serde default `Model::new` set `max_seq_len: default_max_position_embeddings()` -- the literal `4096` fallback used when a config file is silent -- in a struct literal whose two adjacent lines read `cfg.max_position_embeddings` correctly. Every Gemma checkpoint ships an explicit `max_position_embeddings` (8192 for gemma-2b and gemma-7b), so the usable context of every Gemma was halved. `NormalModel::max_seq_len` is not cosmetic: the engine consults it for prompt rejection, prompt truncation, `is_done`, CUDA-graph block sizing, and the context length reported to API clients. All five were being told 4096 while the KV cache next to them was sized for 8192. Gemma was the only model in `models/` that did this -- deepseek2/3/4, gemma2, glm4*, gpt_oss, llama, mistral, mixtral, phi2, phi3 and the rest all read `cfg.max_position_embeddings` in the same position. UNVERIFIED ON HARDWARE: no GPU was used. The change is a one-line substitution of a config-derived value for a constant; there is no arithmetic to measure. Parent system: ArcModels. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcInfer/ArcKV): Gemma-2's global layers were given a sliding-window ring `Model::new` built its cache with `NormalCache::new_sliding(cfg.num_hidden_layers, _, Some(cfg.sliding_window))`, which fills `num_hidden_layers` slots with the SAME `KvCache::Rotating` prototype. Gemma-2 alternates SWA and global attention -- `Attention::new` sets `use_sliding_window: layer_idx.is_multiple_of(2)`, "Order is SWA, global, SWA" -- so the odd, global-attention layers were handed a `sliding_window`-entry ring buffer as well. Below `sliding_window` tokens that is only wasteful. At and beyond it the ring wraps, so a global layer sees the last `sliding_window` keys in RING order rather than positional order, while `forward` pairs it with the non-windowed `attention_mask`. The mask cannot compensate -- it has no way to know the rows were rotated -- so the model degrades silently rather than erroring. gemma-2-9b's `max_position_embeddings` is 8192 against a 4096 window, so the broken regime is inside the supported context, not past it. `NormalCache::from_types` already exists for exactly this and `qwen3.rs` already uses it; gemma2 predates it. The per-layer kinds are now built by `cache_types`, whose predicate is pinned against `Attention::new`'s by test. Tests (CPU, no GPU required): - `cache_types_match_attention_layers` walks all 42 layers of a gemma-2-9b shaped config and asserts SWA layers get a window and global layers get a full-length cache, with the right window and max_seq_len values. - `global_layers_get_a_full_length_cache_not_a_ring` asserts on the cache actually constructed: 21 of 42 slots rotating, not 42. Before this change all 42 were rotating, so that count is what fails. Parent system: ArcInfer / ArcKV. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcMoE): Phi-3.5-MoE's router clamped the wrong way -- min where HF has max HF PhiMoE's `sparsemixer` computes its routing mask as factor = scores.abs().clamp(min=mask_logits_threshold) mask = ((mask_logits_threshold - scores) / factor) > (2 * jitter_eps) `torch.clamp(min=x)` is a LOWER BOUND -- elementwise `max(input, x)`. Both call sites in this port spelled it `broadcast_minimum`, i.e. `min(|s|, t)`, the exact opposite. Verified against `transformers/src/transformers/models/phimoe/modeling_phimoe.py`, where the top-1 and top-2 blocks are identical in this respect. Consequence: `min <= max`, and the numerator `t - s` is non-negative for the top-1 threshold, so the ratio was systematically too LARGE. The `> 2 * jitter_eps` mask therefore over-fired, sending extra experts to `-inf` before the softmax and changing the gate multiplier on every token of every forward. The model still produces plausible text, which is why this survived. There is also a structural difference, not just a numeric one. With `clamp(min=t)` and `t > 0` the denominator is floored at `t`, so a score of exactly zero divides by `t`. Under `minimum` the denominator became `min(0, t) == 0`, producing `inf` and an unconditional mask -- an exposure HF does not have. The duplicated expression is now one function, `sparsemixer_jitter_mask`, so the two call sites cannot drift apart again. Tests (CPU, no GPU required), against a scalar reference written straight from the HF formula so it cannot drift with the code it checks: - `matches_hf_clamp_min_semantics` - `minimum_and_maximum_actually_disagree_here` -- the vacuity guard. It fails if the chosen input cannot separate the bug from the fix, which it did on the first input tried; the mask only notices inside the band where the two ratios straddle `2 * jitter_eps` (`s in [0.98, 1/1.02)` for `t = 1`, `eps = 0.01`). - `zero_score_does_not_divide_by_zero` -- pins the `inf` exposure above. Parent system: ArcInfer / ArcMoE. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * perf(ArcModels): Qwen3-Next synced the device 36 times per token for one vector `Model::forward`'s layer loop called `indices.to_vec1()?` on the recurrent state indices inside each linear-attention arm. `to_vec1` is a device-to-host copy, i.e. a full stall of the compute stream, and `indices` is bound once ABOVE the loop and never written to -- the loop hands it to `pool.gather_*` and `pool.scatter_*`, which read it and mutate the pool, never the index tensor. Qwen3-Next has `full_attention_interval == 4`, so 36 of its 48 layers are linear attention: 36 identical syncs per forward where the data supports one. Hoisted to a single read before the loop. The two `bail!`s that lived next to it move with it. The emptiness check is now gated on the same `has_linear_layers` predicate the existing "indices are required" bail uses, so a model with no linear-attention layers sees no new failure mode. The per-layer offset-divergence check stays in the loop -- `pool` is per-layer, so `first_offset` is genuinely not loop-invariant. UNVERIFIED ON HARDWARE: no GPU was used, and a host sync costs nothing measurable on a CPU device. The claim is 36 device syncs become 1, which is a count of `to_vec1` call sites, not a timing. Parent system: ArcModels. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * perf(ArcAttention): RoPE tested the LENGTH of seqlen_offsets at seven more sites Every rotary implementation in `layers.rs` gated its batched path on `seqlen_offsets.len() == 1` -- the number of ENTRIES -- when the property it needs is whether those entries are EQUAL. `len()` is the batch size, so the batched path was unreachable at every batch size above one, while the values were uniform anyway: the scheduler admits a forward pass only when the participating cache lengths are equal. The per-sequence loop was therefore performing B bit-identical recomputations of the same rotation and concatenating them back together. This is the same defect PR #198 fixed for `DeepSeekV2RotaryEmbedding`. Payoff there is measured, not projected: the `rope` node alone spent 728.6 ms/step of HOST time against 163.5 ms/step of device time at B=256 on an H200, and the fix moved a step from 1082.3 ms to 794.0 ms (1.36x). Same shape, five more model families. NOT A CLAIM ABOUT THESE SITES -- no GPU was used here, and no speedup is asserted for them. Seven sites, all `candle_nn::rotary_emb::{rope, rope_i}` or the fused CUDA kernel: - `PhiRotaryEmbedding::forward`, both arms (partial-rotary and full) - `Phi4MMRotaryEmbedding::forward` - `RotaryEmbedding::forward`, CUDA and candle arms - `GptOssRotaryEmbedding::forward`, CUDA and candle arms `DeepSeekV2RotaryEmbedding`'s two sites are deliberately untouched -- they belong to PR #198, which is still open against `release/openrouter-ready`. The predicate here is `uniform_seqlen_offset`, a free function with a body identical to that PR's `DeepSeekV2RotaryEmbedding::uniform_offset`; the doc comment says so and asks for them to be collapsed when the branches meet. The `rope_cohort_stats` counters keep that PR's names. Second, independent defect in the same hunks: `RotaryEmbedding::forward` gated the fused CUDA kernel on `qh == kh` and `GptOssRotaryEmbedding::forward` on `qh == k.dim(1)?`. That is false for every GQA model -- qwen2, qwen3, mistral, gemma2, phi3, and GPT-OSS itself (64 query heads, 8 KV heads) -- so the kernel was dead for the only architectures that reach it. The equality was never a kernel requirement: `mistralrs_quant::rotary::apply_rotary_inplace` passes `num_heads` and `num_kv_heads` as separate arguments and `kernels/rotary/rotary.cu` walks them as two independent loops (`nq = num_heads * rot_dim`, `nk = num_kv_heads * rot_dim`). The only shape it enforces is `(num_tokens, head_size)` matching between q and k, which GQA preserves. Read from the kernel source; UNVERIFIED ON HARDWARE. The CUDA arm also gains an `is_cuda()` device test. It previously checked only the compile-time feature, so a cuda-enabled binary running on CPU or Metal fell into `apply_rotary_inplace`'s "expects a cuda tensor" bail. Widening the head-count gate would have widened that exposure too. The CUDA cos/sin tables cannot simply narrow for a uniform cohort the way the candle paths can: the kernel indexes `cos_cache + token_idx * rot_dim` over the b-major flattened token axis and hard-bails unless the cache is exactly `(b * seq_len, rot_dim)`. `cohort_rotary_table` therefore tiles -- one broadcast plus one copy for a uniform cohort, and the original `b`-way `cat` for a ragged one, which is the only thing that can express it. Tests (CPU, no GPU required), asserting on raw IEEE-754 bit patterns rather than a tolerance, since the claim is that no arithmetic changed at all: - `uniform_offset_tests_values_not_length` -- `[7; 256]` is uniform. - `uniform_cohort_is_bit_identical_to_the_loop` and `..._for_gpt_j_style` -- both rope styles, against the pre-fix loop held as a verbatim oracle. Each asserts the `rope_cohort_stats` COHORT counter advanced, so the test cannot pass while the fast path sits dark. - `ragged_cohort_keeps_the_loop_and_stays_per_row` -- ragged cohorts must still loop, must still be per-row correct, and must genuinely differ between rows. - `cohort_rotary_table_matches_the_cat_it_replaces` -- shape AND bit equality against the `cat` for uniform, single-row and two ragged cases. This is the only CPU-reachable assertion on the CUDA table layout. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcModels): the phi2 GGUF loader was unreachable, and hid a missing GELU Two defects, one masking the other. 1. `PropsGGUF::try_from` listed `attention.layer_norm_rms_epsilon` in `has_required_keys`, so it demanded `phi2.attention.layer_norm_rms_epsilon`. Phi-2 is a LayerNorm architecture -- this very file builds a `candle_nn::LayerNorm` from that epsilon, complete with a bias -- and no phi2 GGUF carries an RMS-eps key. llama.cpp loads phi2's epsilon with `ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps)` and builds the graph with `LLM_NORM` (`src/models/phi2.cpp`); `gguf-py` spells that key `{arch}.attention.layer_norm_epsilon` (`gguf/constants.py`). So the loader rejected every real file. `quantized_starcoder2.rs`, the other LayerNorm-based GGUF loader in this tree, already reads the non-RMS key. The epsilon is now resolved by `layer_norm_eps`, which accepts either spelling and errors naming both. Nothing that loaded before stops loading. 2. `Mlp::forward` had no activation: `ffn_up` straight into `ffn_down`. Two stacked linear maps compose to a single linear map, so the whole FFN contributed no nonlinearity. `git show 1269bd8` ("Implement GPTQ quantization", Aug 2024) is the deletion: - xs.apply(&self.ffn_up)?.gelu()?.apply(&self.ffn_down) + MatMul.qmethod_matmul(&MatMul.qmethod_matmul(xs, &*self.ffn_up)?, ...) The sibling `quantized_phi3.rs`, rewritten in the same commit, kept its activation -- a transcription slip, not a decision. GELU, not SiLU, and specifically the tanh approximation: `models/phi2.rs` (the unquantized path) applies `cfg.hidden_act`, which is `gelu_new` in microsoft/phi-2's config.json, and llama.cpp builds phi2's FFN with `LLM_FFN_GELU`. Candle's `Tensor::gelu` is the tanh form; `gelu_erf` is not. Defect 2 alone would have been unobservable, since defect 1 meant the code never ran. Fixing 1 without 2 would have shipped a phi2 that loads and answers with a linear FFN. Tests (CPU, no GPU required): - `accepts_the_layer_norm_key_llama_cpp_writes` -- the file shape that used to be rejected outright. - `still_accepts_the_rms_spelling` -- the fix cannot regress anything that worked before. - `missing_epsilon_names_both_keys` - `gelu_is_not_a_linear_map` -- asserts GELU breaks additivity, i.e. that the activation is what stops the FFN collapsing to one matmul. - `gelu_is_the_tanh_approximation_not_erf` -- pins which GELU, since the two forms are separate functions in candle. UNVERIFIED ON HARDWARE: no GPU was used, and no phi2 GGUF was loaded end to end. The metadata-key claim is read from llama.cpp and gguf-py source; the activation claim is read from this repo's own git history and its unquantized sibling. Parent system: ArcModels. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * fix(ArcModels): phi3 and starcoder2 GGUF read their logits from the padding Both models ended `forward` with .i((.., seq_len - 1, ..)) `seq_len` is the batch MAXIMUM, not any particular sequence's length. On a ragged batch that takes the last row of the right-hand padding for every short sequence -- and the same wrong row for all of them -- so every sequence but the longest was sampled from a padded position. The repo already documents this as banned, at `pipeline/inputs_processor.rs`: 🔑 `new_len`, NOT `padded.len()`. `padded.len()` is the batch max, so on a ragged batch every short row would read its logits out of the right-hand PADDING — same wrong row for every sequence. `context_lens` carries the per-row `(start, len)` and `extract_logits` narrows row by row. It was simply never plumbed into these two `forward`s: `pipeline/gguf.rs` passed `context_lens` to Llama, Phi2, Qwen, Qwen3, Qwen3MoE and both X-LoRA arms, and skipped exactly these two. Extraction still happens BEFORE the output projection, so the "only the kept rows go through the vocab matmul" property of the original code is preserved. The result shape becomes `[b, len, vocab]`, matching every other GGUF model here. Uniform batches are unaffected: there every row's `start + len` is `seq_len`, so the change is the identity -- which is why this survived. Tests (CPU, no GPU required), in `pipeline/mod.rs` where `extract_logits` lives, on hidden states whose values encode `(batch, token)` so a wrong row is identifiable rather than merely unequal: - `ragged_batch_reads_each_sequences_own_last_row` - `last_row_indexing_reads_padding_on_a_ragged_batch` -- states the bug as a disagreement and asserts the two disagree, so it cannot pass vacuously. - `uniform_batch_is_unchanged` -- pins the identity case. UNVERIFIED ON HARDWARE: no GPU was used and no GGUF was loaded end to end. Note that the phi3 path is additionally gated by `quantized_phi3.rs`'s own metadata requirements, which were not audited here. Parent system: ArcModels. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * perf(ArcInfer/ArcAttention): V4 copied its whole KV cache twice, for one tensor Under V4's fused MQA the compressor stores ONE tensor and attention reads it as both K and V: three of `deepseek4.rs`'s four `dsv4_attention` call sites pass the same binding twice (`&k, &k` at :1836, `&k_full, &k_full` at :1885, `&k_cached, &k_cached` at :1977). `dsv4_attention` then materialised it twice, at two separate points: 1. the raw working-set narrowing -- `Some(v.narrow(2, rel_base, keep)?.contiguous()?)` next to the identical call for K; 2. the union build -- `Tensor::cat(&[v, comp], 2)?.contiguous()?` next to the identical call for K. Both allocate and copy the full retained key span, once per layer per step. The existing comment above (1) already names the cost -- "the `Tensor::cat` below, which copies the whole raw cache (twice, once for K and once for V) on every decode step" -- without noticing that when V is K the second copy is the same bytes as the first. The two are coupled, which is why this is not a one-line change at the `cat`: two independent `narrow(..).contiguous()` calls produce two distinct handles even when their input was one tensor, so an identity test placed after (1) always says "not aliased". The flag is computed once, before the narrowing, and threaded down. `v_aliases_k` is a HANDLE test, not a value test. Candle's `Tensor` is a refcounted handle carrying a `TensorId` that survives `clone`, so `k.id() == v.id()` holds exactly when the two arguments name one tensor with one layout -- the only condition under which the duplicated work is guaranteed bit-identical. Distinct tensors holding equal values decline and keep both copies, so this can only be conservative, never wrong. `deepseek4.rs:1821` passes genuinely distinct `&k_full, &v_full` and is unaffected. `fused_v_stats::counts()` reports `(aliased, materialized)`: skipping a copy leaves no trace in the output, so a harness that sees `aliased == 0` on a V4 decode must treat that as an environment failure rather than a result. Tests (CPU, no GPU required): - `alias_predicate_is_handle_identity` -- the decision, with no shared state: same binding and `clone` alias; `copy()` and a different tensor do not. - `aliased_v_is_bit_identical_and_engages` -- passing K twice must equal passing a distinct `copy()` of K, compared on raw IEEE-754 bit patterns, and must advance the counters (lower bounds; the counters are process-global and this file's other 60 tests drive them concurrently). - `distinct_v_is_not_aliased` -- negative control against "always reuse k_cat", which would silently compute attention over keys. UNVERIFIED ON HARDWARE: no GPU was used. The claim is that two of the copies of a tensor are provably redundant and are no longer made -- a count of allocations, not a measured time or a measured memory figure. No byte or millisecond number is asserted. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 * chore: rustfmt the code this wave added, and clear the warnings it introduced Scoped cleanup for the eight fixes in this branch. No behaviour change. - `layers.rs`, `gemma2.rs`: rustfmt applied to the lines this wave added. `cargo fmt --all` is NOT used -- it reformats ~90 upstream files, which the fork policy forbids. Both files were checked against `3c22cb5b5` first: `layers.rs` was rustfmt-clean at baseline and every deviation was in new code; `gemma2.rs` had 5 pre-existing own-file deviations and now has 4, because replacing the multi-line `NormalCache::new_sliding(..)` call happened to remove one. Every other file this branch touches was already unformatted upstream and its deviation count is unchanged. - `quantized_phi3.rs`, `quantized_starcoder2.rs`: `seq_len` and starcoder2's `IndexOp` import became unused when `.i((.., seq_len - 1, ..))` was replaced by `extract_logits`. - `lib.rs`: re-export `models::dsv4_attention::fused_v_stats`. `models` is a private module, so without this the engagement counters are dead code -- and shipping a counter nothing can read is the same failure the counter exists to prevent. Placed alongside the existing `pub use models::deepseek4::{..}` escape hatch, with a `//` comment rather than `///` so it does not split the sorted `use` group. Verified: `cargo test -p mistralrs-core --lib` 668 passed / 0 failed; `cargo check -p mistralrs-core --lib` at 24 warnings, the same count as `3c22cb5b5`; the scoped clippy lane (`--no-deps -p arc-bench -p arc-engine -p arc-cuda-graph -p arc-cli -p mistralrs-quant --tests --examples -- -D warnings`) green. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
heydryft
force-pushed
the
release/openrouter-ready
branch
from
August 21, 2026 15:36
17597e6 to
d665956
Compare
heydryft
added a commit
that referenced
this pull request
Aug 21, 2026
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. Three things were duplicated, and two of them did not merely duplicate, they broke: 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. With the counters shared, `deepseek_rope_cohort_tests`'s private mutex stopped serialising anything: `rope_cohort_tests` took no lock at all, because until now it was the only writer. Three tests failed with `(2, 1)` where they wanted `(1, 0)` -- the counter catching a missing lock for the second time in its life. The lock now lives beside the counters it protects, as `rope_cohort_stats::engagement_lock`, and both modules take that one. 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. 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
force-pushed
the
perf/rope-cohort-uniform-offsets-194
branch
from
August 21, 2026 16:21
654a6fe to
75779d3
Compare
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
force-pushed
the
perf/rope-cohort-uniform-offsets-194
branch
from
August 21, 2026 16:47
75779d3 to
f287c7c
Compare
Contributor
Author
Rebased onto
|
Contributor
Author
⚡ SWITCHED OFF ON PURPOSE — and here is what turns it back onRead this before assuming the PR is finished. Off is a temporary state with
Also recorded in-repo, enforced, at Why it is off rather than onThe kernel is batch-correct by construction and the dispatch bug that kept it
"Unverified" means unmeasured, not new. Polarity is
|
…tep -> 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>
…ch 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>
… 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>
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>
…_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
heydryft
force-pushed
the
perf/rope-cohort-uniform-offsets-194
branch
from
August 21, 2026 19:19
417ea25 to
192f8ff
Compare
…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
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 #194 (
release/openrouter-ready). Built and verified onorigin/release/openrouter-ready@17597e620. The same two commits are also onperf/rope-cohort-uniform-offsets, branched from real master3c22cb5b5, where they are equally green —mastercarries the identical defect at the identical lines (layers.rs:1631,:1674).The bug is one word
DeepSeekV2RotaryEmbedding::{forward,forward_inverse_tail}gated the batched path onseqlen_offsets.len() == 1— the length of the offset vector, where what matters is the distinctness of its values.len()is the batch size, so the batched path was unreachable at every batch size above one.And the values are uniform anyway.
scheduler::default_scheduler::select_running_bucketadmits a forward pass only when cache lengths are exactly equal; the one producer of a ragged dense batch is gated behindARC_MTP_PER_SEQ_KV, which defaults off (deepseek4::ragged_row_q0, whose own doc says "the offsets would be uniform"). So the loop was performing B bit-identical recomputations and concatenating them back together.What it cost
Real profile, B=256 pure decode, steady state (
nsys --delay 330 --duration 25, 204.80 tok/s):ucopy_bf16copy2d_bf16rope_i_bf16uneg_bf1619.9% of step time, and the only cost in the profile that gets worse as batch grows. Residuals are 0.1–3.5%, within window-edge quantisation (25 s at 204.80 tok/s over B=256 is ~20 steps).
Attributions, since one was previously wrong:
ucopy_bf16is candle's strided copy (cuda_backend/mod.rs, non-contiguous branch ofcopy_strided_src) — the two last-dim narrows made contiguous per sequence. It is not the mHC F32 round trip.copy2d_bf16iscat_contiguous, which issues one launch per argument even for dim 0 — the 256-argumentTensor::catthat exists solely to undo the split the same function just made. The k-pathcatmoves 128 bytes per launch.uneg_bf16is not a rotary-half op and should not be fused: it isself.sin.narrow(..)?.neg()?inside the loop, negating a[1,32]slice of a constant table 11,008 times per step.neg_sinis now materialised once at construction. Not a kernel to fuse — a kernel to delete.Prediction, to be confirmed or refuted on hardware
Cohort form is 43 layers × 7 = 301 launches, 129 of them
rope_i— landing RoPE at the same order asfp8_matmul_tiled(301) andqtip2b_grouped_gemm(129) instead of dwarfing them.This is scope, not a promise. It is a launch-count fix; the bytes moved are unchanged. The
copy2d5.0% anduneg0.8% are work removed outright. Theucopy10.2% +rope_i3.9% keep their bytes and shed 99.4% of their launches — at ~8 KB/launch that is launch-bound, but how much of the 19.9% actually returns must be measured. I have no GPU; this needs a box.Correctness
Bit-identical by construction, not by tolerance: with 2-D
[T, D/2]cos/sin the kernel'sstride_bis 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. Batching changes which launch computes an element, never the arithmetic.Seven tests hold the pre-fix loop as a verbatim oracle and assert raw IEEE-754 bit equality — for uniform and ragged offsets, at
seq_len1 and 2 — plus a one-ULP negative control proving the comparator can fail. Ragged offsets still take the loop verbatim, so enablingARC_MTP_PER_SEQ_KVbehaves exactly as today. No raggedindex_selectpath is added; nothing ships dark.Engagement is counted (
rope_cohort_stats::counts) because a fast path that silently declined would pass every equality assertion trivially. It already earned itself — it caught a missing lock in one of these tests.Second commit, separately revertable
fused_qk_gatedeclined with"per-sequence position offsets"for the same length-vs-values reason, so the fused one-launchqk_norm_ropekernel was dead at every batch above 1 — the regime it was written for. Its doc advertises "replacing 16 candle launches per layer with 1"; at B=256 it replaced nothing.The kernel needed no change:
qk_norm_rope.cuflattens(b,t)intoblockIdx.y, recoversb = bt / seq_len, and reads the table atpos_offset + t— independent ofb, exactly right for uniform rows. Grid is(n_heads+1, batch*seq_len); buffers are sizedbatch * …. Nothing assumedbatch == 1.Split out because, unlike the first commit, this one is not bit-identical-by-construction — it swaps in a CUDA kernel that has never run above batch 1. Verify with
ARC_QK_VERIFY=1(bit-compares against eager at every layer);ARC_QK_FUSED=0restores eager from the same binary.Verification run
cargo test -p mistralrs-core --lib→ 654 passed, 0 failed (647 baseline + 7 new). Scoped clippy lane → 0 errors.mistralrs-coreclippy warnings 314 = 314 baseline (no new).rustfmt --checkonlayers.rs→ 0 diffs;deepseek4.rs→ 28, exactly the pre-existing count (fork policy: no mass-reformat).Not verified: anything on a GPU.
Separate, larger blast radius
The identical
len() == 1pattern is atlayers.rs:704-709and:2886-2891on the genericRotaryEmbedding, so every non-V4 model on those paths has the same defect. Deliberately not touched here — different models, different test surface.Lane update (08-21): rebased to
55160369c, e6b0ccf audited, ragged arm gatheredRebase. All four commits rebased clean onto
release/openrouter-ready@55160369c(#181, #121, #217, #196, #195, #130 landed under the branch). No conflicts; the env-flag revert trap (std::env::var(..).is_ok()presence spellings) did not reappear — every gate in this diff reads by value (env_flag_is_set/fused_cohort_enabled_from).e6b0ccf (the dropped distinctness hunk) is fully subsumed — recovered as a verification, not a commit. Audited hunk by hunk against the new head: its
uniform_seqlen_offsetpredicate and generic-rotary tests are already on the integration branch (via the #196 lineage; baselayers.rs:385,:4096-4102); its two V4 dispatch hunks are carried by this branch's first and third commits in strictly stronger form (counters,neg_sinprecompute, bit-pattern oracles instead of max-abs-diff). A cherry-pick emptied after conflict resolution toward the stronger versions and was dropped, per the standing rule.New commit — the RAGGED arm, gathered (
ARC_ROPE_COHORT, default OFF). The uniform fix left rows at genuinely different positions on the per-sequence loop. candle'srope_iaccepts rank-3[B, T, D/2]cos/sin whenBmatches the input batch (rope_check_cs, candle-nnrotary_emb.rs:228-239at pinned rev89ab14e; CUDA wrapper setsstride_bat:126-130; CPU path indexesi + b_i*t*d/2), so the ragged arm becomes one gatheredrope_iper projection:index_selectper step per device, not per layer (pitfall feat(v4): load + wire the full MTP decoder block for speculative drafting #5).rope_i(+ up to 2contiguous), inverse tail = 1rope_i+ 1contiguous+ 1 NoPEcat⇒ ~5-7 launches/layer, ~230-300/step at 43 layers, batch-independent — versus the loop's ~9 ops × 43 × B ≈ 99,072/step at B=256. Same order as the uniform path's 301.env_flag_is_set,=0means OFF), registered incapability_reachability.rsbeside the fused-cohort entry, same doctrine: unverified means unmeasured, not new. Flag unset keeps the loop byte-for-byte; the uniform arm is untouched and stays unconditionally on.cuda::qk_norm_rope::engaged_countgains its first reader — a test proving both counters advance — closing the silent-success hole where an A/B could read a permanently-declining path as live. The fused path's engagement log lines (first engage + every 100k) were already in.Mutation checks (fix reverted → test fails → fix restored → passes):
ragged_cohort_forward_matches_the_loop_bitwise1063410531vs loop1049743782(offsets[0,4,9], T=1); 3 tests failcohort_gather_rebuilds_when_the_cohort_moves[1065353216 ×4](the offset-0 row, 1.0f) vs correct[3207025968, 1064028839, 1065339796, 1065353082]len() == 1(the decline-gate bug)uniform_offset_tests_values_not_lengthuniform_seqlen_offset(&[7; 256]): leftNonevs rightSome(7)cohort_forward_is_bit_identical_to_the_per_sequence_loop(0, 1)vs right(1, 0)— the batched arm no longer engages; 4 tests failcohort_tablesout_of_bounds_offset_fails_loud_on_both_armsindex-select invalid index 32 with dim size 32) — the pinned guard message is gone; on CUDA this same case is not reliably checked and reads garbage silentlyrope cohort gather out of bounds: offset 32 + seq_len 1 exceeds the 32-row cos/sin table; 722/722 passVerification at
227c34f8a:cargo check -p mistralrs-coreclean;cargo test -p mistralrs-core --lib→ 722 passed, 0 failed;capability_reachability→ 11 passed (including the newARC_ROPE_COHORTregistry entry); scoped clippy lane → exit 0, 0 errors. Still not verified: anything on a GPU — the flag stays OFF until the box A/B runs.🤖 Generated with Claude Code