fix(ArcGraph): correct cuGraphAddNode ABI; make GPU-autonomous decode reachable - #130
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 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ |
|
Triage verdict: REBASE, not close — LIVE and valuable. Both signals confirmed on current master, with one important correction. Confirmed unique against master:
Correction to the premise this was triaged under. The blocking So the honest statement is: this PR is genuinely live and unique at the Conflict: 1 file, Not merging now: the new |
Retargeted at the integration branchBase changed: The queue is being restructured to the shape the owner asked for: one PR open against Two things had to land on
This PR was not closed and is not considered stale. An audit of the queue found the overwhelming majority of it to be real work that was never merged, not noise. What you need to do: rebase onto |
0440dfb to
b4c4546
Compare
… test counts Integration fix for #198 landing after #204. Both PRs made the same correction -- `seqlen_offsets.len() == 1` tests the length of the vector where the property needed is the distinctness of its values -- and each shipped its own spelling of the predicate and its own engagement counters. `uniform_seqlen_offset`'s doc comment on the integration branch asked for exactly this collapse ("there must not be two spellings of the same predicate"); this is that collapse. 1. `pub mod rope_cohort_stats` was declared TWICE in one file (layers.rs:400 from #204, :1568 from #198). That is a hard name collision -- the file did not compile. #204's is the superset (it has `record_cohort` / `record_per_sequence` as well as `counts`), so #198's is the one that goes. 2. `DeepSeekV2RotaryEmbedding::uniform_offset` had a byte-identical body to `uniform_seqlen_offset`. Removed; `forward`, `forward_inverse_tail` and `deepseek4::fused_qk_gate` now call the free function, so the two cannot drift apart. The reasoning #198's doc comment carried and #204's did not -- that uniformity is *enforced* by `select_running_bucket` rather than merely observed, and the measured B=256 cost -- moves onto the survivor. 3. Collapsing the counters made them shared, and shared process-global counters cannot be delta-asserted from a parallel test binary. `deepseek_rope_cohort_ tests` held a private mutex, `rope_cohort_tests` held none (until now it was the only writer), and three tests duly failed with `(2, 1)` where they wanted `(1, 0)`. A shared mutex was the first fix and it is also wrong: it obliges every present and future test that touches ANY rotary to know about a lock in another module. Two such tests already exist and did not -- `deepseek4::tests::{standard_layer_rope_is_unscaled_and_compressed_is_not, mtp_block_takes_the_standard_unscaled_rope_table}` both drive `DeepSeekV2RotaryEmbedding::forward` -- and they race the cohort tests. That was caught by merging #198 and #121 together and running the full suite: `--test-threads=1` passes 707, the parallel run fails one. So the counters now also keep a `#[cfg(test)]` per-thread mirror, and the tests assert on `local_counts()`. Each test's delta is its own by construction, no cooperation from other modules is required, and the serving instrument -- the global `counts()` -- is byte-for-byte unchanged. Also drops #198's copy of `uniform_offset_tests_values_not_length`: with one predicate it asserted nothing the surviving copy does not, against the same six cases. `cargo test -p mistralrs-core --lib`: 699 passed, 0 failed; 707 passed on the merged #198+#181+#130+#121 tree. The 99,072 -> 301 launch figure remains #198's own branch measurement and is NOT re-measured here. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Rebased onto
|
ABI independently verified against a machine-generated binding
pub struct CUDA_CONDITIONAL_NODE_PARAMS {
pub handle: CUgraphConditionalHandle,
pub type_: CUgraphConditionalNodeType,
pub size: ::core::ffi::c_uint,
pub phGraph_out: *mut CUgraph,
pub ctx: CUcontext,
}Field-for-field identical to this PR's corrected Two more, same source:
(Note the NVIDIA doc page lists the struct's fields alphabetically, not in |
04d4162 to
5ed2b8e
Compare
… reachable
The headline is an ABI mismatch that explains a failure already on our record.
`ffi.rs` declared `cudaGraphAddNode` with FIVE arguments and the first two
transposed. CUDA 13.1 `cuda.h:21829` declares SIX, and the first parameter is
an OUT pointer:
CUresult cuGraphAddNode(CUgraphNode *phGraphNode, CUgraph hGraph,
const CUgraphNode *dependencies,
const CUgraphEdgeData *dependencyData,
size_t numDependencies,
CUgraphNodeParams *nodeParams);
So the driver received a `CUgraph` handle where it expected `CUgraphNode *`
and wrote the new node handle THROUGH it, while `numDependencies` received a
pointer and `nodeParams` an uninitialised register. That is memory corruption
from a declaration, and it matches the previously unexplained "host heap
corruption" CUDA-graph capture failure we had recorded with no cause.
Four more defects, each of which alone stops the autonomous loop:
* Conditional handles were created with `flags = 0`. `cuda.h:21919` applies
`defaultLaunchValue` only when `CU_GRAPH_COND_ASSIGN_DEFAULT` is set, so the
condition read 0 at launch and a WHILE body executed ZERO times while every
CUDA call returned success. Measured both ways on an H200.
* The WHILE body was captured into a throwaway graph and then DESTROYED,
leaving the conditional body empty. Bodies are now populated with
`cuStreamBeginCaptureToGraph` (`cuda.h:1976` names it as the supported way),
and a node-count assert refuses a body that recorded nothing.
* `CUDA_CONDITIONAL_NODE_PARAMS` was missing its 5th field, `ctx`.
* `DecodeState` allocated I64 while every kernel signature is `int32_t*`, so
host and device disagreed about where element `i` lived; and `reset()`
reallocated the buffers, silently invalidating the pointers a captured graph
had baked in.
`tensor_device_ptr` had no I32 arm, so `CudaSampler::sample` returned
`unsupported dtype I32` on EVERY call -- the correct top-k/top-p sampler had
never once executed on a GPU. Its test suite is a CPU simulator, which by
construction cannot observe that.
Sampling now runs on device and survives replay. The sampler previously wired
into the graph body took `rng_offset` as a BY-VALUE kernel argument, which a
captured graph bakes: measured 1 distinct token over 64 replays, versus 64
distinct for a device-resident `rng_state`. It also was not nucleus sampling --
it walked the vocabulary in TOKEN-ID order.
The replacement is a hybrid, because measurement showed neither component wins
everywhere (H200, vocab 129280, host+GPU verified exclusive before and after;
support width -> legacy / hybrid us):
1 -> 134.9/166.6 8 -> 503.0/624.2 64 -> 3127.8/4178.5
512 -> 24320.7/4185.0 4096 -> 205125.2/4187.3 12928 -> 664200.0/4188.5
Exact enumeration while the nucleus is small, threshold bisection past a fixed
budget, chosen block-uniformly ON DEVICE because a captured graph cannot ask
the host which branch to take. Costs +24% on the peaked distributions real
models produce -- 32 us, 0.05% of a 66.68 ms V4 decode step -- and removes a
664 ms cliff that was ~10x an entire decode step for a single token.
No cap and no narrowing: exceeding the budget changes which ALGORITHM selects
the nucleus, never which tokens it contains. Verified against the enumerating
sampler with branch counters proving which path each case exercised:
peaked TV=0.0000 A=20000 B=0
scrambled order TV=0.0000 A=20000 B=0 (15.6% per-draw agreement, so the
orders genuinely diverged)
diffuse TV=0.0000 A=0 B=3000
tie boundary TV=0.0000 A=20000 B=0
narrowed (ctrl) TV=0.5027 -> correctly DISAGREE
Known and deliberately left: on TIED diffuse supports the fallback keeps 116
tokens where enumeration keeps 111 (TV=0.1103) -- it keeps every token sharing
a boundary bit pattern, erring WIDER, never narrower. And the budget is
derived from bisection's pass count (32) when it should follow the FALLBACK's
(~70), which costs a 0.75x window at support ~64; correcting it needs a more
diffuse fixture so the fallback stays covered.
Not claimed: no V4 end-to-end number. Autonomous decode is unreachable on V4
by construction -- no PagedAttention means `cache_config` is None and the
runner is never built (`normal.rs:1907`, pinned by `normal_loaders.rs:5687`) --
and `cuda.h:1971` bars alloc/free nodes inside a conditional body against
11,436 allocations per token.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
5ed2b8e to
b372a43
Compare
…ix/arcgraph-cugraphaddnode-abi
… test counts Integration fix for #198 landing after #204. Both PRs made the same correction -- `seqlen_offsets.len() == 1` tests the length of the vector where the property needed is the distinctness of its values -- and each shipped its own spelling of the predicate and its own engagement counters. `uniform_seqlen_offset`'s doc comment on the integration branch asked for exactly this collapse ("there must not be two spellings of the same predicate"); this is that collapse. 1. `pub mod rope_cohort_stats` was declared TWICE in one file (layers.rs:400 from #204, :1568 from #198). That is a hard name collision -- the file did not compile. #204's is the superset (it has `record_cohort` / `record_per_sequence` as well as `counts`), so #198's is the one that goes. 2. `DeepSeekV2RotaryEmbedding::uniform_offset` had a byte-identical body to `uniform_seqlen_offset`. Removed; `forward`, `forward_inverse_tail` and `deepseek4::fused_qk_gate` now call the free function, so the two cannot drift apart. The reasoning #198's doc comment carried and #204's did not -- that uniformity is *enforced* by `select_running_bucket` rather than merely observed, and the measured B=256 cost -- moves onto the survivor. 3. Collapsing the counters made them shared, and shared process-global counters cannot be delta-asserted from a parallel test binary. `deepseek_rope_cohort_ tests` held a private mutex, `rope_cohort_tests` held none (until now it was the only writer), and three tests duly failed with `(2, 1)` where they wanted `(1, 0)`. A shared mutex was the first fix and it is also wrong: it obliges every present and future test that touches ANY rotary to know about a lock in another module. Two such tests already exist and did not -- `deepseek4::tests::{standard_layer_rope_is_unscaled_and_compressed_is_not, mtp_block_takes_the_standard_unscaled_rope_table}` both drive `DeepSeekV2RotaryEmbedding::forward` -- and they race the cohort tests. That was caught by merging #198 and #121 together and running the full suite: `--test-threads=1` passes 707, the parallel run fails one. So the counters now also keep a `#[cfg(test)]` per-thread mirror, and the tests assert on `local_counts()`. Each test's delta is its own by construction, no cooperation from other modules is required, and the serving instrument -- the global `counts()` -- is byte-for-byte unchanged. Also drops #198's copy of `uniform_offset_tests_values_not_length`: with one predicate it asserted nothing the surviving copy does not, against the same six cases. `cargo test -p mistralrs-core --lib`: 699 passed, 0 failed; 707 passed on the merged #198+#181+#130+#121 tree. The 99,072 -> 301 launch figure remains #198's own branch measurement and is NOT re-measured here. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…es/step -> 301 (#198) * perf(ArcAttention): batch V4 RoPE over the cohort — 99,072 launches/step -> 301 `DeepSeekV2RotaryEmbedding::{forward,forward_inverse_tail}` gated the batched path on `seqlen_offsets.len() == 1` — the LENGTH of the offset vector, where the thing that matters is the DISTINCTNESS of its values. `len()` is the batch size, so the batched path was unreachable at every batch size above one. The offsets are uniform anyway. `scheduler::default_scheduler`'s `select_running_bucket` admits a forward pass only when cache lengths are exactly equal, and the one producer of a ragged dense batch is gated behind `ARC_MTP_PER_SEQ_KV`, which defaults off (`deepseek4::ragged_row_q0`). So the loop was performing B bit-identical recomputations and concatenating them back. Measured, B=256 pure decode, steady state (nsys, 204.80 tok/s): kernel share launches/step per (layer, seq) ucopy_bf16 10.2% 22,794 2 (2*43*256 = 22,016) copy2d_bf16 5.0% 33,610 3 (3*43*256 = 33,024) rope_i_bf16 3.9% 33,054 3 (3*43*256 = 33,024) uneg_bf16 0.8% 11,008 1 (1*43*256 = 11,008, exact) 19.9% of step time, and the only cost in the profile that gets WORSE as batch grows. `copy2d` is the 256-argument `Tensor::cat` that exists solely to undo the split the same function just made; it goes to zero. `uneg` is 11,008 launches to negate a [1,32] slice of a CONSTANT table, so `neg_sin` is now materialised once at construction — not a kernel to fuse, a kernel to delete. Cohort form is 43 layers * 7 = 301 launches (129 of them rope_i), which lands RoPE at the same order as fp8_matmul_tiled (301) and qtip2b_grouped_gemm (129) instead of dwarfing them. PREDICTION, not a result: this is a launch-count fix, the bytes moved are unchanged, and how much of the 19.9% actually returns must be measured on hardware. Ragged offsets still take the loop verbatim, so enabling ARC_MTP_PER_SEQ_KV behaves exactly as today. No ragged index_select path is added — nothing ships dark. Bit-identical by construction: with 2-D [T, D/2] cos/sin the kernel's stride_b is 0 (candle-nn/src/rotary_emb.rs), so every batch row reads the same cos/sin row and each output element is an independent two-multiply-one-add. Seven tests hold the pre-fix loop as a verbatim oracle and assert raw IEEE-754 bit equality (not a tolerance) for uniform AND ragged offsets, at seq_len 1 and 2, with a one-ULP negative control proving the comparator can fail. Engagement is counted (`rope_cohort_stats::counts`): a fast path that silently declined would pass every equality assertion trivially. The counter already earned itself — it caught a missing lock in one of these tests. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(ArcAttention): un-deaden the fused qk_norm_rope kernel above batch 1 `fused_qk_gate` declined with "per-sequence position offsets" whenever `seqlen_offsets.len() != 1` — the same length-vs-values confusion as the RoPE dispatch in the parent commit, and with the same consequence: the fused one-launch kernel was dead at EVERY batch size above one, i.e. in exactly the serving regime it was written for. Its own doc comment advertises "replacing 16 candle launches per layer with 1"; at B=256 it replaced nothing. The kernel is already batch-correct and needed no change. `qk_norm_rope.cu` flattens (b, t) into blockIdx.y, recovers `b = bt / seq_len`, and reads the table at `pos_offset + t` — a position independent of `b`, which is exactly right when the rows share an offset. Its grid is (n_heads + 1, batch * seq_len) and the wrapper sizes buffers as `batch * ...`. Nothing assumed batch == 1. Gate now takes the uniform offset via the shared `DeepSeekV2RotaryEmbedding::uniform_offset` helper, so both call sites cannot drift apart. Genuinely ragged rows still decline and still take the eager chain. Split from the parent commit so it can be reverted independently: this one swaps in a CUDA kernel that has never run above batch 1, and unlike the parent it is not bit-identical-by-construction — it is bit-identical by the kernel's own design, which `ARC_QK_VERIFY=1` checks against the eager chain at every layer. Verify before trusting; `ARC_QK_FUSED=0` restores the eager path from the same binary. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcAttention): one cohort predicate, one counter pair, per-thread test counts Integration fix for #198 landing after #204. Both PRs made the same correction -- `seqlen_offsets.len() == 1` tests the length of the vector where the property needed is the distinctness of its values -- and each shipped its own spelling of the predicate and its own engagement counters. `uniform_seqlen_offset`'s doc comment on the integration branch asked for exactly this collapse ("there must not be two spellings of the same predicate"); this is that collapse. 1. `pub mod rope_cohort_stats` was declared TWICE in one file (layers.rs:400 from #204, :1568 from #198). That is a hard name collision -- the file did not compile. #204's is the superset (it has `record_cohort` / `record_per_sequence` as well as `counts`), so #198's is the one that goes. 2. `DeepSeekV2RotaryEmbedding::uniform_offset` had a byte-identical body to `uniform_seqlen_offset`. Removed; `forward`, `forward_inverse_tail` and `deepseek4::fused_qk_gate` now call the free function, so the two cannot drift apart. The reasoning #198's doc comment carried and #204's did not -- that uniformity is *enforced* by `select_running_bucket` rather than merely observed, and the measured B=256 cost -- moves onto the survivor. 3. Collapsing the counters made them shared, and shared process-global counters cannot be delta-asserted from a parallel test binary. `deepseek_rope_cohort_ tests` held a private mutex, `rope_cohort_tests` held none (until now it was the only writer), and three tests duly failed with `(2, 1)` where they wanted `(1, 0)`. A shared mutex was the first fix and it is also wrong: it obliges every present and future test that touches ANY rotary to know about a lock in another module. Two such tests already exist and did not -- `deepseek4::tests::{standard_layer_rope_is_unscaled_and_compressed_is_not, mtp_block_takes_the_standard_unscaled_rope_table}` both drive `DeepSeekV2RotaryEmbedding::forward` -- and they race the cohort tests. That was caught by merging #198 and #121 together and running the full suite: `--test-threads=1` passes 707, the parallel run fails one. So the counters now also keep a `#[cfg(test)]` per-thread mirror, and the tests assert on `local_counts()`. Each test's delta is its own by construction, no cooperation from other modules is required, and the serving instrument -- the global `counts()` -- is byte-for-byte unchanged. Also drops #198's copy of `uniform_offset_tests_values_not_length`: with one predicate it asserted nothing the surviving copy does not, against the same six cases. `cargo test -p mistralrs-core --lib`: 699 passed, 0 failed; 707 passed on the merged #198+#181+#130+#121 tree. The 99,072 -> 301 launch figure remains #198's own branch measurement and is NOT re-measured here. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix(ArcAttention): the batched fused qk_norm_rope path is OPT-IN The parent commit un-deadens the fused Q/K kernel above batch 1 by fixing a length-vs-values gate. That is the right fix to the gate and the wrong default: it makes a CUDA kernel LIVE on a path it has never run on. This repo has paid for that twice in one week. * TCFRAG shipped default-ON carrying "UNVERIFIED ON HARDWARE — NEVER RUN" in its own header. It held 63 GB and permanently broke a layer through a poisoned `OnceLock`. #209 retired it to opt-in. * The fused-512 attention path ran for four days silently dropping the attention mask, at 12% agreement with the reference, because nobody ran the path it had quietly become live on. So the doctrine, and this commit: **unverified means default-off, and "unverified" means unmeasured, not new.** `ARC_QK_FUSED_COHORT` gates the batched path, default OFF. With it unset, `batch > 1` declines and takes the eager chain exactly as master does today — so the parent commit ships its gate fix with no behaviour change, and the cohort path is available to be measured rather than discovered in production. `batch == 1` is untouched: that path is already live and already exercised. Read by VALUE, `== Some("1")`, split into the pure `fused_cohort_enabled_from` so the polarity is testable without mutating the environment (racy across `cargo test`'s threads, `unsafe` since the 2024 edition, and meaningless against a `OnceLock` that latches the first read). Two tests pin it, and they run with no features and no GPU on every PR: unset is OFF, and `0`/`false`/ `off`/`true`/`on`/`2`/`"1 "` are all OFF. That second one is not pedantry — #212 converted 23 `ARC_*` flags this week because `var_os(..).is_some()` made `ARC_FOO=0` mean ON, which silently cancels any A/B whose control sets zero. To turn it on, once `ARC_QK_VERIFY=1` has bit-compared it against the eager chain on hardware and it is measured faster: `ARC_QK_FUSED_COHORT=1`. When that measurement exists, this default flips and the doc comment goes with it. `cargo test -p mistralrs-core --lib`: 701 passed, 0 failed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * perf(ArcAttention): gather the RAGGED cohort's RoPE — one rank-3 rope_i, opt-in ARC_ROPE_COHORT The uniform-offset fix (#198, and the shared predicate it now rides on) left one arm untouched: rows at genuinely DIFFERENT positions still take the per-sequence loop — ~9 ops x 43 layers x B per step on the measured B=256 profile (~99,000 launches, 19.9% of step time), of which the copy2d half only undoes a split the same function just made. Nothing structural required the loop: candle's rope_i accepts a rank-3 [B, T, D/2] cos/sin whenever B matches the input batch (rope_check_cs, candle-nn/src/rotary_emb.rs, pinned rev 89ab14e; the CUDA wrapper sets stride_b for rank-3, the CPU path indexes i + b_i*t*d/2). So the ragged arm becomes: gather the per-row table rows once — [B*T] u32 indices, host-built because the offsets ARE host data, memoised per instance so the H2D and the three index_selects run once per step per device (the model shares one rotary per device per table kind), CLAUDE.md pitfall #5 — then ONE rope_i per projection. ~5 launches per layer where the loop paid ~9 x B. DEFAULT OFF behind ARC_ROPE_COHORT, read by VALUE via env_flag_is_set (=0 means OFF; #212), latched once. Unset keeps the loop byte-for-byte. Registered in capability_reachability.rs beside the fused-cohort entry, same doctrine: unverified means default-off, and unverified means unmeasured. Both ragged arms are named methods driven directly by deepseek_ragged_cohort_tests — bit-pattern parity (not tolerance) for mixed offsets at B in {1,3}, T in {1,2}, ascending and descending, forward and inverse tail, plus a stale-memo test (the cache must rebuild when the cohort moves) and a pinned unset-default. engaged_count in cuda/qk_norm_rope.rs also gains its first reader: a test that proves the counters advance, closing the silent-success hole where an A/B could read a dead path as live. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf * fix(ArcAttention): bounds-guard the cohort gather — both ragged arms must fail LOUD The loop arm's narrow(0, offset, seq_len) hard-errors past the table end; the gathered arm fed the same out-of-range index to index_select, which is not reliably bounds-checked on the CUDA backend and can read garbage SILENTLY. A failure-mode asymmetry on exactly the arm the hardware A/B will exercise is the D18 house fault, so the offsets — host data, the check is free — are now guarded host-side in cohort_tables before any index reaches the device, with the offending offset, seq_len and table size named in the error. out_of_bounds_offset_fails_loud_on_both_arms pins: the loop arm errors, the cohort arm errors on forward AND inverse tail, the message is the GUARD's (a backend error cannot pass the test by coincidence), and offset 31 + T=1 on a 32-row table still works, so the bound is not off by one. Parent system: ArcInfer / ArcAttention. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01WK8VgocBUrr5djjE1ZfCNf --------- Co-authored-by: Nirupam Bhowmick <support@runcrate.ai> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
The ABI bug, first — it closes an open mystery
arc-cuda-graph/src/ffi.rsdeclaredcudaGraphAddNodewith five arguments and the first two transposed. CUDA 13.1cuda.h:21829declares six, and the first parameter is an out-pointer:The driver therefore received a
CUgraphhandle where it expectedCUgraphNode *and wrote the new node handle through it, whilenumDependenciesgot a pointer andnodeParamsan uninitialised register.This gives a cause to the "host heap corruption" CUDA-graph capture failure already on our record, which had been logged with no explanation.
Why graph replay had never once launched
Four further defects, each individually fatal to the autonomous loop:
flags = 0cuda.h:21919appliesdefaultLaunchValueonly withCU_GRAPH_COND_ASSIGN_DEFAULT; the condition read 0 and the WHILE body ran zero times while every CUDA call returned successcuStreamBeginCaptureToGraph(cuda.h:1976names it as the supported way) with a node-count assertCUDA_CONDITIONAL_NODE_PARAMSmissing its 5th fieldctxcuda.h:1958DecodeStateallocated I64 againstint32_t*kernels;reset()reallocated buffersAnd
tensor_device_ptrhad no I32 arm, soCudaSampler::samplereturnedunsupported dtype I32on every call — the correct top-k/top-p sampler had never executed on a GPU. Its tests are a CPU simulator, which by construction cannot observe that.Measured on an H200 (CUDA 13.1, sm_90)
GPU-autonomous decode works. WHILE conditional node, one launch for N steps: 0.0156 host calls/step at N=128 vs 2.0 for launch-and-sync-per-step. Device-side loop overhead is constant in body size (+2.90/+2.39/+1.98 µs at 1/64/256 kernels) — 0.004% of a 66.68 ms V4 step. The control that makes it mean something: one instantiated graph run with device-side
target=1thentarget=Nreturned counter 1 then N, so the device chose the trip count.Replay safety. The old in-graph sampler took
rng_offsetas a by-value kernel argument, which a captured graph bakes: 1 distinct token over 64 replays, vs 64 distinct with a device-residentrng_state. It also was not nucleus sampling — it walked the vocabulary in token-id order.Sampler cost (host + GPU verified exclusive before and after):
Neither component wins everywhere, so this ships a hybrid: exact enumeration while the nucleus is small, threshold bisection past a fixed budget, chosen block-uniformly on device because a captured graph cannot ask the host which branch to take. Costs +24% on peaked distributions (32 µs, 0.05% of a decode step) and removes a 664 ms cliff that was ~10× an entire decode step for one token.
Distribution-verified against the enumerating sampler, with branch counters proving which path each case exercised —
TV=0.0000on peaked, scrambled-order (15.6% per-draw agreement, so the orders genuinely diverged), diffuse (A=0 B=3000), and tie-boundary cases; the deliberately-narrowed control correctly disagreed atTV=0.5027.Known, and deliberately left
TV=0.1103) — it keeps every token sharing a boundary bit pattern, erring wider, never narrower.cache_configisNoneand the runner is never built (normal.rs:1907, pinned bynormal_loaders.rs:5687) — andcuda.h:1971bars alloc/free nodes inside a conditional body against 11,436 allocations per token.🤖 Generated with Claude Code