Eight cited defects that were reaching nobody: 4 wrong answers (default-on), 3 capacity, 1 dead loader - #204
Merged
Merged
Conversation
…de 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
…dow 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
…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
…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
… 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
…ng 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
…adding
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
…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
…troduced
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
Code Metrics Report━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Language Files Lines Code Comments Blanks ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ C Header 5 305 210 52 43 CSS 2 1181 1036 34 111 CUDA 78 27867 19381 5568 2918 Dockerfile 1 39 22 8 9 JavaScript 16 3546 2676 482 388 Jinja2 7 694 656 5 33 JSON 74 4600 4597 0 3 Makefile 1 6 5 0 1 Metal Shading Lan| 33 12224 9431 1142 1651 PowerShell 1 300 227 30 43 Python 147 15371 12677 824 1870 Shell 42 10258 6849 2734 675 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 204 45330 0 35183 10147 |- BASH 72 1655 1203 331 121 |- C 3 17 17 0 0 |- CUDA 2 84 56 16 12 |- JSON 18 708 708 0 0 |- PowerShell 1 1 1 0 0 |- Python 23 1008 787 113 108 |- Rust 66 2051 1716 77 258 |- TOML 6 207 164 0 43 |- YAML 5 41 36 5 0 (Total) 51102 4688 35725 10689 ───────────────────────────────────────────────────────────────────────────────── Rust 678 335044 288629 17172 29243 |- Markdown 497 30361 471 26183 3707 (Total) 365405 289100 43355 32950 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ Total 1337 502989 357396 92632 52961 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
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
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
… 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.
Eight defects that were each small, cited, and reaching nobody. One commit per fix, so any one reverts alone. Branched off
master@3c22cb5b5.No GPU was used. Every claim below is a source-level derivation or a CPU test. No speedup is asserted anywhere.
Base
master, because3c22cb5b5isorigin/masterand this gives an exact 9-commit diff. If this should ride inside #194's train instead,git rebase --onto origin/release/openrouter-ready 3c22cb5b5is clean — none of the files touched here differ betweenmasterandrelease/openrouter-ready.#198 (
perf/rope-cohort-uniform-offsets-194) is open and unmerged. It fixes the sameseqlen_offsets.len() == 1defect forDeepSeekV2RotaryEmbeddingand introducesuniform_offset+rope_cohort_stats.Those two DeepSeekV2 sites (
layers.rs:1631,:1674) are deliberately untouched here — they belong to #198. The seven sites this PR fixes are the other rotary implementations. The predicate added here,uniform_seqlen_offset, has a body identical to #198'suniform_offsetand its doc comment says so and asks for the two to be collapsed when the branches meet; the counter module keeps #198's names. Merge #198 first, or collapse the duplicate predicate on merge.What landed
gemma.rs:459read the literal-4096 serde default formax_seq_lenwhere its two neighbouring lines readcfg.max_position_embeddings— halving every Gemma's context, and the context reported to API clientsgemma2.rsgave its global-attention layers a 4096-entry ring buffer vianew_sliding; past 4096 tokens they see only the last 4096 keys, in ring order, against a non-windowed maskphi3_5_moe.rsusedbroadcast_minimumwhere HF'sscores.abs().clamp(min=t)ismaximum— the router's gate multiplier was wrong on every token, plus a divide-by-zero →infexposure HF does not havequantized_phi2.rs— two defects: the loader demanded an RMS-eps key phi2 GGUFs never carry (so it rejected every real file), which masked an FFN with no activation at all, deleted by1269bd8abin Aug 2024quantized_phi3.rs/quantized_starcoder2.rsread last-row logits from the padding on ragged batches;context_lenswas passed to every other GGUF arch and skipped for these twolayers.rs— the RoPE cohort defect at seven more sites, plus theqh == khgate that switched the fused CUDA kernel off for every GQA modelqwen3_next.rs— ato_vec1()device sync inside the layer loop on a loop-invariant tensor: 36 syncs/token → 1dsv4_attention.rs— V4 materialised its whole KV cache twice for one tensor, at two coupled pointsVerified premises that were wrong, and skipped
deepseek4.rs:2076wo_aBF16 cache — "this one is a one-liner": it is not, and there is no one-liner.self.wo_ais still live at three other places::2103(the non-grouped o_proj path),:4149/:4925/:4988(ISQ re-quantisation),:5132(UQFF serialisation). The FP8 original cannot be dropped. The BF16 cache is the deliberate price of a measured RUN-161 saving (~69 ms/token, 27% of decode). No change made.phi3_5_moe"Arc clamps againstmax_logitsrather than HF'smax − s": false. HF clamps against the max too (factor = scores.abs().clamp(min=mask_logits_threshold), andmask_logits_thresholdis the max). Only the min-vs-max inversion is real, and that is what was fixed.quantized_phi2"may be masked — check reachability": it was masked, and by a second bug worth fixing. Both are fixed here.layers.rssite list was partly stale.:1631/:1674are perf(ArcAttention): V4 RoPE was launched per-sequence — 99,072 launches/step -> 301 #198's DeepSeekV2 sites, not "seven more"; the actual seven are:659,:693,:1887,:2656,:2691,:2837,:2875.dsv4_attention"v_cat… materialised 4× where 1 suffices": the duplication is real but happens at two coupled points, not one. Thenarrow(..).contiguous()at:773runs before thecatand destroys the aliasing, so an identity test placed only at thecatalways says "not aliased". Fixing only thecatwould have been a no-op. The specific "4×/10.7 GiB" figure was not reproducible from source and is not claimed.Not shipped, with the reason
qwen3_next.rs:888num_layers: cfg.num_hidden_layers= 48 vs 12 real full-attention layers — the 4× paged-KV over-allocation is real, and this is not a small fix.num_layers()is overloaded three ways and only one of them wants 12:normal_loaders.rs:157—l.min(num_layers)for device placement; 12 would collapse layers 12–47 onto one device.normal.rs:1492—check_dense_layer_inventory(&tags, num_layers)over all 48 ISQ layer tags.cache_engine.rs:251/:288andpaged_attention/mod.rs— the paged KV sizing, the only caller that wants 12.The correct fix needs a separate
ModelConfigLike::num_kv_cache_layers()and a global→attention-slot map, becausecache_engine.rs:288sliceslayer_devicesby global index (attention layers are 3, 7, …, 47, which sit on different devices under a multi-GPU map) andqwen3_next.rs'sforwardindexeskv_cache[layer_idx]globally at:959.ModelConfigMetadatacannot carry the override without a new field across its 150 construction sites, andNormalModel::config()returns it concretely from ~45 models. That is a structural change to the paged allocator, in upstream-owned files, that cannot be validated without a GPU — shipping it unverified risks out-of-bounds slot aliasing on a real run. Design recorded here for a follow-up.Verification
cargo test -p mistralrs-core -p mistralrs-quant -p mistralrs-vision --lib— 668 / 328 / 2 passed, 0 failedcargo check -p mistralrs-core --lib— 24 warnings, identical to3c22cb5b5--no-deps -p arc-bench -p arc-engine -p arc-cuda-graph -p arc-cli -p mistralrs-quant --tests --examples -- -D warnings) — greencargo fmt --allnot used. Every touched file's own-file rustfmt deviation count was measured against3c22cb5b5and is unchanged or lower.23 new CPU tests. They assert on raw IEEE-754 bit patterns where the claim is "no arithmetic changed", and several carry explicit vacuity guards —
minimum_and_maximum_actually_disagree_herefailed on the first input tried and was only made to pass by finding an input inside the band where the two formulas actually diverge.Engagement is asserted, not assumed:
rope_cohort_stats::counts()andfused_v_stats::counts()exist so a fast path cannot sit dark behind a green test, and the RoPE and aliased-V tests fail if the counter does not advance.What is UNVERIFIED ON HARDWARE
Everything. Specifically, these are reachable only on CUDA and were read from kernel source rather than run:
qh == khremoval in bothRotaryEmbedding::forwardandGptOssRotaryEmbedding::forward—kernels/rotary/rotary.cutakesnum_headsandnum_kv_headsas separate arguments and walks them as two independent loops, so GQA is natively supported;cohort_rotary_table's tiling for the CUDA cos/sin cache — its shape and bit-equality against thecatit replaces are CPU-tested, but the kernel launch is not;is_cuda()device test on the CUDA rotary arm.🤖 Generated with Claude Code
https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7