From c1d2b27f4801d89a1f984a8a46bdcd9c3735e373 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 21:30:19 +0100 Subject: [PATCH 01/13] fix(ArcGraph): Tensor::arange in the captured forward is a dangling HOST pointer MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit V4 decode capture reached "forward RECORDED" and then SIGSEGV'd on its very first cuGraphLaunch, with a device-side assert from candle's index_select: ids[id_i] < src_dim_size (T = __nv_bfloat16, I = unsigned int) That dtype pair is the bf16 cos table indexed by U32 positions, i.e. `self.cos.index_select(&positions, 0)` in `DeepSeekV2RotaryEmbedding::forward_at_positions` — NOT the TD-MoE `tid2eid.index_select` (tid2eid is i64, so it cannot produce T=bf16), and not the embedding (its ids come from the already-address-stable graph-mode buffer). `positions` was built one line earlier by `compressed_kv_from_rows` with `Tensor::arange(0u32, t_c, dev)`. `arange` materialises a transient host `Vec` and uploads it via `CudaDevice::clone_htod` -> an ASYNC `cuMemcpyHtoDAsync`. Under `cuStreamBeginCapture` that copy is not executed, it is RECORDED: the graph stores the HOST POINTER and re-reads it on the first launch and on every replay. The `Vec` is freed when the expression returns, so the graph copied freed host memory into `positions` and handed garbage indices to index_select. This is not an address-staleness bug — it fires on the first launch, before any replay — which is why the static input_ids buffer did not prevent it. candle already documents this exact mechanism and works around it in `CudaDevice::htod_info` (it leaks the host source while capturing); nothing protects `clone_htod`, which is the path `arange` takes. Fix: serve the strided compressed positions from a per-(ratio, device) table built once, outside capture, and take a zero-copy `narrow` view each step. The values are a pure function of (ratio, capacity) and never change. This also removes a host round trip from every compressor layer of every decode step. Keyed per ratio because V4 interleaves CSA and HCA compression ratios; a single-slot cache would thrash and reintroduce the copy on every layer. --- mistralrs-core/src/layers.rs | 128 +++++++++++++++++++++++++ mistralrs-core/src/models/deepseek4.rs | 72 ++++++-------- 2 files changed, 157 insertions(+), 43 deletions(-) diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index af68e3aa1..fd5fcef1b 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2388,6 +2388,81 @@ impl Module for QLinear { } } +// Cached strided-position tables for the V4 compressor, one per +// `(ratio, device)`. Keyed that way because V4 mixes compression ratios across +// layers (CSA and HCA); a single-slot cache would thrash and reintroduce a host +// copy on every layer. See `compress_positions` for the rationale. +std::thread_local! { + static COMPRESS_POSITIONS: std::cell::RefCell> = + const { std::cell::RefCell::new(Vec::new()) }; +} + +/// Round a needed length up so the table is (re)built rarely — and therefore +/// during warmup rather than during capture. +const COMPRESS_POS_CHUNK: usize = 1024; + +/// `[t_c]` U32 tensor holding `[0, ratio, 2*ratio, …, (t_c-1)*ratio]`, served +/// as a zero-copy view into a cached device table. +/// +/// # Why this is not a `Tensor::arange`: that is a graph-capture HAZARD +/// +/// `Tensor::arange` (and `from_vec`/`new`/`from_slice`) builds its data as a +/// transient host `Vec` and uploads it through `CudaDevice::clone_htod`, which +/// issues an **async** `cuMemcpyHtoDAsync` and returns. Outside capture that is +/// merely a host round trip. **Inside `cuStreamBeginCapture` it is a +/// correctness bug**: the copy is not executed, it is *recorded* — the graph +/// stores the HOST POINTER and re-reads it on the first launch and on every +/// replay. The `Vec` is freed as soon as the expression returns, so the graph +/// copies freed host memory into the tensor. +/// +/// That is exactly how V4 decode capture died: `compressed_kv_from_rows` built +/// its strided positions with `arange`, and the first `cuGraphLaunch` filled +/// them with garbage, so the very next op — +/// `self.cos.index_select(&positions, 0)` in +/// `DeepSeekV2RotaryEmbedding::forward_at_positions` — tripped candle's +/// device-side bounds assert (`ids[id_i] < src_dim_size`, `T = __nv_bfloat16` +/// for the bf16 cos table, `I = unsigned int` for the U32 positions) and the +/// process took SIGSEGV. candle already documents this failure mode for kernel +/// dims/strides and works around it in `CudaDevice::htod_info` by leaking the +/// host source while capturing; nothing protects `clone_htod`, which is what +/// `arange` uses. +/// +/// The table is a pure function of `(ratio, capacity)` and its values never +/// change, so it is built **once, outside capture** (warmup steps run long +/// before `begin_capture`) and every decode step then takes a `narrow` view of +/// it. No host memory is touched inside the captured region, and the base +/// device address is stable across replays. +pub fn compress_positions(t_c: usize, ratio: usize, device: &Device) -> Result { + if ratio == 0 { + candle_core::bail!("compress_positions: ratio must be non-zero"); + } + COMPRESS_POSITIONS.with(|c| { + let mut cache = c.borrow_mut(); + let hit = cache.iter().position(|(r, t)| { + *r == ratio && t.device().same_device(device) && t.dim(0).unwrap_or(0) >= t_c + }); + let idx = match hit { + Some(i) => i, + None => { + // Drop any smaller/stale table for this (ratio, device) so the + // cache cannot grow without bound. + cache.retain(|(r, t)| !(*r == ratio && t.device().same_device(device))); + let cap = t_c.div_ceil(COMPRESS_POS_CHUNK) * COMPRESS_POS_CHUNK; + // The one `arange` — deliberately here and not in the forward. + let table = Tensor::arange_step( + 0u32, + (cap * ratio) as u32, + ratio as u32, + device, + )?; + cache.push((ratio, table)); + cache.len() - 1 + } + }; + cache[idx].1.narrow(0, 0, t_c) + }) +} + /// Thread-local GPU positions tensor for CUDA graph mode. /// When set, `RotaryEmbedding::forward()` uses GPU-side gather instead of /// CPU-side `narrow()`, making the forward pass graph-capture compatible. @@ -3669,6 +3744,59 @@ impl Module for ScaledEmbedding { } } +#[cfg(test)] +mod compress_positions_tests { + use super::*; + + /// The cache replaced `(arange(0..t_c).to_f32() * ratio).to_u32()`. The + /// values must be bit-identical to that expression, or every compressed + /// row silently rotates at the wrong absolute position — a quality bug no + /// test outside this one would catch. + #[test] + fn matches_the_arange_expression_it_replaced() -> Result<()> { + let dev = Device::Cpu; + for ratio in [1usize, 2, 4, 8, 32] { + for t_c in [1usize, 3, 17, 64] { + let want = (Tensor::arange(0u32, t_c as u32, &dev)?.to_dtype(DType::F32)? + * (ratio as f64))? + .to_dtype(DType::U32)? + .to_vec1::()?; + let got = compress_positions(t_c, ratio, &dev)?.to_vec1::()?; + assert_eq!(got, want, "ratio={ratio} t_c={t_c}"); + } + } + Ok(()) + } + + /// Growth must reuse the table, not silently rebuild it every step — + /// rebuilding is the host copy this cache exists to remove, and doing it + /// inside capture is the crash it exists to prevent. + #[test] + fn a_longer_request_within_capacity_reuses_the_same_storage() -> Result<()> { + let dev = Device::Cpu; + let ratio = 4usize; + let first = compress_positions(1, ratio, &dev)?; + let grown = compress_positions(COMPRESS_POS_CHUNK, ratio, &dev)?; + // Same backing table => same length after the chunk round-up. + assert_eq!(first.dim(0)?, 1); + assert_eq!(grown.dim(0)?, COMPRESS_POS_CHUNK); + let entries = COMPRESS_POSITIONS.with(|c| c.borrow().len()); + assert_eq!(entries, 1, "one (ratio, device) entry, not one per call"); + // Distinct ratios must NOT evict each other: V4 interleaves CSA and HCA + // layers, so a single-slot cache would rebuild on every layer. + let _ = compress_positions(8, 32, &dev)?; + let _ = compress_positions(8, ratio, &dev)?; + let entries = COMPRESS_POSITIONS.with(|c| c.borrow().len()); + assert_eq!(entries, 2, "ratios must coexist"); + Ok(()) + } + + #[test] + fn zero_ratio_is_refused() { + assert!(compress_positions(4, 0, &Device::Cpu).is_err()); + } +} + #[cfg(test)] mod graph_mode_mask_tests { use super::*; diff --git a/mistralrs-core/src/models/deepseek4.rs b/mistralrs-core/src/models/deepseek4.rs index 3e84142bf..f5c24564e 100644 --- a/mistralrs-core/src/models/deepseek4.rs +++ b/mistralrs-core/src/models/deepseek4.rs @@ -1412,11 +1412,18 @@ impl Attention { let t_c = comp.dim(2)?; // Compressed entry j sits at absolute position j*ratio. Apply the // layer's (compress-θ) RoPE to the last qk_rope_head_dim dims there. - // NOTE: builds a small arange each call (a host/device sync); fine on - // the correctness-first dense path — the long-context sparse-gather - // kernel is the place to precompute this. + // + // This used to build the positions with `Tensor::arange` on every call. + // That is a host round trip on the hot path AND — decisively — the + // reason CUDA-graph capture of the V4 decode forward SIGSEGV'd on its + // first launch: `arange` uploads a transient host `Vec`, and a captured + // `cuMemcpyHtoDAsync` records the host POINTER, not the bytes, so the + // graph re-read a freed `Vec` and handed garbage indices straight to + // the `index_select` on the next line. `compress_positions` serves a + // zero-copy view of a table built once, outside capture; see its doc + // comment for the full mechanism. let dev = comp.device(); - let positions = compressed_row_positions(t_c, ratio, dev)?; + let positions = crate::layers::compress_positions(t_c, ratio, dev)?; self.rotary_emb .forward_at_positions(&comp, self.cfg.qk_rope_head_dim, &positions) } @@ -2995,46 +3002,25 @@ fn append_kv_mqa( } } -// Cached `[t_c]` U32 absolute positions of the compressed rows, keyed by row -// count and ratio. Compressed row `j` sits at absolute position `j * ratio`, -// which depends on nothing but `j` — so for a given (width, ratio) this vector -// is a constant and rebuilding it per layer per step was pure waste. +// The compressed-row positions used to live here as a SINGLE-slot thread-local +// keyed on `(t_c, ratio, device)`. That slot is now `layers::compress_positions` +// (a per-`(ratio, device)` table, chunk-rounded and served as a `narrow` view), +// because a single slot THRASHES on V4: the model interleaves CSA (ratio 4) and +// HCA (ratio 128) layers, so consecutive layers of the same step evicted each +// other and the `arange_step` was paid again on every layer — including inside +// the capture region, where it is a SIGSEGV and not merely a cost. // -// It is also a correctness precondition for capture, not just a saving: -// `Tensor::arange` on a GPU device is an H2D copy from pageable host memory, -// which cannot be recorded into a CUDA graph at all. Under fixed capacity the -// width stops moving, so this hits from the second decode step onward and the -// copy leaves the capture region for good. -thread_local! { - static COMPRESSED_ROW_POSITIONS: std::cell::RefCell> = - const { std::cell::RefCell::new(None) }; -} - -fn compressed_row_positions(t_c: usize, ratio: usize, dev: &Device) -> Result { - COMPRESSED_ROW_POSITIONS.with(|c| -> Result { - let mut c = c.borrow_mut(); - if let Some((w, r, t)) = c.as_ref() { - if *w == t_c && *r == ratio && t.device().same_device(dev) { - return Ok(t.clone()); - } - } - // `j * ratio` is integer arithmetic. It used to be laundered through - // float — `arange(u32) -> cast_u32_f32 -> affine(*ratio) -> - // cast_f32_u32` — three kernel launches to compute a strided range - // that `arange_step` produces directly. The thread-local above makes - // that free on a hit, but a miss (first token, or a window/device - // change) still paid it. The values are identical: the old chain's - // `fmaf(j, ratio, 0.0)` is exact for every `j * ratio` a sequence can - // reach (both operands and the product are integers well inside f32's - // exact range), and `to_dtype(U32)` truncated it back. `ratio` is - // 1 / 4 / 128 (`CompressRatio::ratio`), never 0, so `arange_step`'s - // zero-step bail is unreachable. Multiplied in usize so the product - // cannot wrap before the u32 narrowing. - let t = Tensor::arange_step(0u32, (t_c * ratio) as u32, ratio as u32, dev)?; - *c = Some((t_c, ratio, t.clone())); - Ok(t) - }) -} +// The arithmetic argument that licensed `arange_step` in the first place is kept +// here, because it is what makes the two spellings interchangeable: `j * ratio` +// is integer arithmetic that used to be laundered through float — `arange(u32) +// -> cast_u32_f32 -> affine(*ratio) -> cast_f32_u32`, three kernel launches to +// compute a strided range `arange_step` produces directly. The values are +// identical: the old chain's `fmaf(j, ratio, 0.0)` is exact for every `j * ratio` +// a sequence can reach (both operands and the product are integers well inside +// f32's exact range), and `to_dtype(U32)` truncated it back. `ratio` is 1 / 4 / +// 128 (`CompressRatio::ratio`), never 0, so `arange_step`'s zero-step bail is +// unreachable. Multiplied in usize so the product cannot wrap before the u32 +// narrowing. /// The FIXED number of compressed rows the CUDA-graph decode arm reads back. /// From 64fe1d678ac36c42acd18745f080509fd9ff118f Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 21:37:10 +0100 Subject: [PATCH 02/13] fix(ArcGraph): the other three arange calls in the captured V4 forward MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A sweep of every host->device constructor reachable from the B=1 decode forward found three more `Tensor::arange` calls with the same defect as the compressor positions, all in `dsv4_attention`, all unconditional on every layer of every step: :610 kp = arange(raw_base, raw_base + t_k) absolute key positions :613 qp = arange(q0, q0 + t_q) query positions (ONE element at decode -- a whole clone_htod for 4 bytes, 43x per token) :628 bp = arange(0, t_c) compressed-block indices Each is a contiguous integer ramp cast to F32, so all three are views of a single cached device ramp: `layers::positions_f32(start, len, dev)` narrows `[0.0, 1.0, 2.0, …]`. Capacity grows by powers of two from 8192, so a typical run allocates the ramp once during warmup and never again -- rebuilding inside capture is precisely the hazard being removed. Fixing only the compressor site would not have been enough: any one of these would have re-armed the same dangling-host-pointer crash on the first launch. Also removes three H2D copies per layer per token from the decode hot path (129 per token at 43 layers), independent of graph capture. --- mistralrs-core/src/layers.rs | 77 +++++++++++++++++++++ mistralrs-core/src/models/dsv4_attention.rs | 9 ++- 2 files changed, 85 insertions(+), 1 deletion(-) diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index fd5fcef1b..87ec2cae7 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2463,6 +2463,51 @@ pub fn compress_positions(t_c: usize, ratio: usize, device: &Device) -> Result> = + const { std::cell::RefCell::new(Vec::new()) }; +} + +/// Smallest ramp worth allocating: 32 KB of F32. Large enough that a typical +/// run allocates once, at warmup, and never again — which is the point, since a +/// rebuild *during* capture is the bug this exists to avoid. +const IOTA_MIN_LEN: usize = 8192; + +/// `[len]` F32 tensor holding `[start, start+1, …, start+len-1]`, served as a +/// view into a cached device ramp. +/// +/// Replaces `Tensor::arange(start, start+len, dev)?.to_dtype(F32)` on the V4 +/// attention mask path. `arange` is a host→device copy, and inside CUDA-graph +/// capture it is worse than slow: the recorded `cuMemcpyHtoDAsync` keeps the +/// HOST pointer of a `Vec` that is freed the moment the expression returns, so +/// every launch and replay fills the tensor from freed memory. See +/// [`compress_positions`] for the full mechanism and the crash it produced. +/// +/// The ramp grows by powers of two so rebuilds are rare and land during warmup. +pub fn positions_f32(start: usize, len: usize, device: &Device) -> Result { + let end = start + len; + IOTA_F32.with(|c| { + let mut cache = c.borrow_mut(); + let hit = cache + .iter() + .position(|t| t.device().same_device(device) && t.dim(0).unwrap_or(0) >= end); + let idx = match hit { + Some(i) => i, + None => { + cache.retain(|t| !t.device().same_device(device)); + let cap = end.max(IOTA_MIN_LEN).next_power_of_two(); + // The one `arange` — deliberately here and not in the forward. + let ramp = Tensor::arange(0u32, cap as u32, device)?.to_dtype(DType::F32)?; + cache.push(ramp); + cache.len() - 1 + } + }; + cache[idx].narrow(0, start, len) + }) +} + /// Thread-local GPU positions tensor for CUDA graph mode. /// When set, `RotaryEmbedding::forward()` uses GPU-side gather instead of /// CPU-side `narrow()`, making the forward pass graph-capture compatible. @@ -3795,6 +3840,38 @@ mod compress_positions_tests { fn zero_ratio_is_refused() { assert!(compress_positions(4, 0, &Device::Cpu).is_err()); } + + /// `positions_f32` replaced `arange(start, start+len).to_f32()` at three + /// V4 attention-mask sites. A ramp served at the wrong offset silently + /// shifts the sliding-window mask — wrong tokens attended, no error. + #[test] + fn positions_f32_matches_the_arange_expression_it_replaced() -> Result<()> { + let dev = Device::Cpu; + for (start, len) in [(0usize, 1usize), (0, 7), (5, 1), (13, 128), (1000, 64)] { + let want = Tensor::arange(start as u32, (start + len) as u32, &dev)? + .to_dtype(DType::F32)? + .to_vec1::()?; + let got = positions_f32(start, len, &dev)?.to_vec1::()?; + assert_eq!(got, want, "start={start} len={len}"); + } + Ok(()) + } + + /// The ramp must survive growth past its initial capacity, and a view of it + /// must still reshape (the three call sites all reshape immediately). + #[test] + fn positions_f32_grows_and_its_views_reshape() -> Result<()> { + let dev = Device::Cpu; + let short = positions_f32(0, 4, &dev)?; + assert_eq!(short.reshape((1, 4))?.dims(), &[1, 4]); + let past_min = IOTA_MIN_LEN + 3; + let long = positions_f32(past_min, 2, &dev)?; + assert_eq!(long.to_vec1::()?, vec![past_min as f32, past_min as f32 + 1.0]); + assert_eq!(long.reshape((2, 1))?.dims(), &[2, 1]); + // One ramp per device, not one per request. + assert_eq!(IOTA_F32.with(|c| c.borrow().len()), 1); + Ok(()) + } } #[cfg(test)] diff --git a/mistralrs-core/src/models/dsv4_attention.rs b/mistralrs-core/src/models/dsv4_attention.rs index 2cbdcfc19..bc7cc7e06 100644 --- a/mistralrs-core/src/models/dsv4_attention.rs +++ b/mistralrs-core/src/models/dsv4_attention.rs @@ -935,7 +935,14 @@ pub fn dsv4_attention( None => { let threshold = Tensor::new(&[((q0 + 1) / ratio) as u32], dev)?.reshape((1, 1))?; - Tensor::arange(0u32, t_c as u32, dev)? + // The block indices come from the cached `(ratio, + // device)` table at ratio 1, not `Tensor::arange`. + // `compress_positions(t_c, 1, dev)` IS `[0, 1, …, t_c-1]` + // in U32 — bit-identical to the `arange` it replaces — + // but it is a zero-copy `narrow` view of a table built + // once during warmup, so the per-layer host->device copy + // goes away. See `layers::compress_positions`. + crate::layers::compress_positions(t_c, 1, dev)? .reshape((1, t_c))? .broadcast_lt(&threshold)? // [1, t_c] (u8) } From fd1757b1b62e500dadc9a3d6671946aafdf8b578 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 23:39:50 +0100 Subject: [PATCH 03/13] =?UTF-8?q?fix(ArcGraph):=20BLOCKER=203=20=E2=80=94?= =?UTF-8?q?=20the=20graph=20KV=20window=20is=20a=20ring,=20not=20a=20prefi?= =?UTF-8?q?x?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `append_graph` read a fixed `0..read_capacity` window but wrote at the ABSOLUTE position, so from token `read_capacity` on, the new row landed outside the window that is read: the freshest key was invisible to attention and a stale row was attended in its place. The arm was correct only below `sliding_window`. `read_capacity` IS `sliding_window`, so a ring of that many slots holds exactly the key set the raw branch is defined to attend. Permuting the keys is invariant here: softmax over keys is permutation-invariant, K and V go to the same slot so pairing is preserved, and V4 rotates K before caching so each row carries its own absolute RoPE. The validity mask needs no change — `slot <= position` is already ring-correct in both regimes (filling: written iff slot <= position; full: all slots valid, predicate universally true). The slot is derived on-device (`layers::graph_ring_slot`); only the constant `capacity` is baked into kernel arguments. A host-resolved `position % capacity` would repeat blocker 1's defect one buffer over. --- mistralrs-core/src/kv_cache/single_cache.rs | 18 +++++-- mistralrs-core/src/layers.rs | 52 +++++++++++++++++++++ 2 files changed, 67 insertions(+), 3 deletions(-) diff --git a/mistralrs-core/src/kv_cache/single_cache.rs b/mistralrs-core/src/kv_cache/single_cache.rs index b04c78fd6..207686fce 100644 --- a/mistralrs-core/src/kv_cache/single_cache.rs +++ b/mistralrs-core/src/kv_cache/single_cache.rs @@ -326,9 +326,21 @@ impl SingleCache { self.pregrow_for_graph(src, read_capacity)?; let ad = self.all_data.as_ref().unwrap(); - // Device-slot write of the new token; `write_kv_inplace` uses the - // buffer's real capacity from all_data's dims, so the slot is correct. - mistralrs_quant::kvwrite::write_kv_inplace(ad, src, position)?; + // RING write. The read below is a fixed `0..read_capacity` window, so + // writing at the ABSOLUTE position put every token from `read_capacity` + // onward OUTSIDE the window that is read: the freshest key became + // invisible and a stale row was attended in its place, making this arm + // correct only while the context was shorter than `sliding_window`. + // + // `read_capacity` IS `sliding_window` here, so a ring of that many + // slots holds exactly the key set the raw branch is defined to attend. + // The slot is derived on the DEVICE from the device position — a + // host-resolved `position % capacity` would be baked into the recorded + // kernel and every replay would write to the capture step's slot. See + // `layers::graph_ring_slot` for why permuting the keys is invariant + // here and why the existing length mask is already ring-correct. + let slot = crate::layers::graph_ring_slot(position, read_capacity)?; + mistralrs_quant::kvwrite::write_kv_inplace(ad, src, &slot)?; // Fixed-offset window -> constant shape for capture/replay. ad.narrow(self.dim, 0, read_capacity) } diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index 87ec2cae7..d3321844b 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2748,6 +2748,58 @@ pub fn graph_mode_length_mask(positions: &Tensor, capacity: usize, dtype: DType) .reshape((b, 1, 1, capacity)) } +/// `positions mod capacity` — the RING slot a graph-mode decode writes its new +/// K/V row into — computed entirely on the device. +/// +/// # Why the graph KV buffer has to be a ring (RUN-161) +/// +/// The graph decode arm reads a constant `capacity`-wide window (slots +/// `0..capacity`) so the launch geometry never moves, but it used to *write* at +/// the ABSOLUTE position. Those agree only while `position < capacity`: from +/// token `capacity` onward the new row was written past the end of the window +/// that is read, so the freshest key was invisible to attention and a stale row +/// was attended in its place. The arm was therefore correct only below +/// `sliding_window`. +/// +/// A ring of exactly `capacity` slots is the natural fix, because `capacity` +/// **is** `sliding_window`: the raw branch attends precisely the last +/// `sliding_window` tokens, so a full ring holds exactly the right key set and +/// never holds a key that should have been evicted. +/// +/// ## Why permuting the keys is safe +/// +/// A ring stores keys out of order. That is invariant for this attention: +/// * softmax over keys is permutation-invariant, and K and V are written to +/// the SAME slot, so key `i` stays paired with value `i`; +/// * V4 rotates K *before* caching, so every stored row already carries its +/// own absolute RoPE — no consumer re-derives a position from column index; +/// * the validity mask stays [`graph_mode_length_mask`]'s `slot <= position`, +/// which is already exactly right for a ring: while the ring is filling +/// (`position < capacity`) slot `s` is written iff `s <= position`, and once +/// it is full every slot is valid and the predicate is universally true. +/// +/// ## Why this is not just `position % capacity` on the host +/// +/// The host knows the position, but a host-resolved slot becomes a literal in +/// the recorded kernel's arguments and every replay would then write to the +/// capture step's slot — the same defect that made host-resolved RoPE offsets +/// wrong (see `DeepSeekV2RotaryEmbedding::cos_sin_for`). Only `capacity` is +/// baked, and it is a genuine compile-time-constant of the run. +/// +/// F32 is exact for every integer below 2^24, and positions are bounded by +/// `max_position_embeddings` (163 840 for V4), so the float round trip is exact. +pub fn graph_ring_slot(positions: &Tensor, capacity: usize) -> Result { + if capacity == 0 { + candle_core::bail!("graph_ring_slot: capacity must be non-zero"); + } + let p = positions.to_dtype(DType::F32)?; + // `affine`'s multiplier is a constant kernel argument, which is safe to + // bake: `capacity` is fixed for the life of the run. Only the position is + // allowed to vary between replays, and it is read from device memory. + let blocks = p.affine(1.0 / capacity as f64, 0.0)?.floor()?; + (p - (blocks * capacity as f64)?)?.to_dtype(DType::U32) +} + /// Check if graph-mode positions are set. #[cfg(feature = "cuda")] pub fn has_graph_mode_positions() -> bool { From cc6b58d31e8994c29b3c756577abc1a5e93dc866 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 23:50:06 +0100 Subject: [PATCH 04/13] =?UTF-8?q?fix(ArcGraph):=20BLOCKER=202a=20=E2=80=94?= =?UTF-8?q?=20pin=20the=20retained=20raw=20tail=20so=20capture=20stops=20h?= =?UTF-8?q?itting=20an=20unwarmed=20alloc?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit MEASURED root cause of the first-launch CUDA_ERROR_ILLEGAL_ADDRESS (700): exactly ONE capture-time allocator miss, 327680 bytes. At hidden 4096 in BF16 that is 8192 B/token, so the miss is a 40-token buffer — the retained raw tail — one token wider than the 319488 B (39-token) miss on the preceding warmup step. A capture-time miss is an unstable graph memory node, i.e. an INVALID graph, which is why it surfaces as an illegal address on the first cuGraphLaunch rather than as a slow one. The width moves because `advance_uniform` retains only what a future compressed row can consume, which is a function of `tokens % ratio`. That gives `ratio` distinct widths on a period of `ratio`: 4 for CSA (the "four unwarmed allocation sizes" an earlier session recorded) but 128 for HCA, growing one token per step — which is why no practical warmup count ever covered it. `pin_tail_width` holds the tail at `span_groups * ratio + margin`, which is an upper bound on what the unpinned policy would keep (width is maximised at `r = ratio - 1` as `span_groups*ratio + margin - 1`), so it only ever retains more than the compressor needs, never less. Takes effect once the sequence is that long; capture must land after that point — 144 tokens for V4's HCA layers. --- .claude/run_capture.sh | 45 ++++++++++++++++ mistralrs-core/src/kv_cache/xs_rolling.rs | 65 ++++++++++++++++++++++- mistralrs-core/src/models/deepseek4.rs | 42 +++++++++++++++ 3 files changed, 151 insertions(+), 1 deletion(-) create mode 100644 .claude/run_capture.sh diff --git a/.claude/run_capture.sh b/.claude/run_capture.sh new file mode 100644 index 000000000..627db7045 --- /dev/null +++ b/.claude/run_capture.sh @@ -0,0 +1,45 @@ +#!/bin/bash +# ArcGraph replay-correctness measurement leg (bench mode). +# +# Bench mode, NOT server mode: capture requires ARC_CANDLE_ALLOC_CACHE, whose +# server-mode leak (132 allocs/step, zero frees) OOMs a 143 GB card on the first +# completion. Another agent owns that fix; this leg does not touch it. +# +# Usage: run_capture.sh +set -u +export CUDA_HOME=/usr/local/cuda-12.4 +export PATH=$CUDA_HOME/bin:$HOME/.cargo/bin:$PATH +export LD_LIBRARY_PATH=$CUDA_HOME/lib64:${LD_LIBRARY_PATH:-} + +NAME="$1"; shift +GENLEN="$1"; shift +PROMPTLEN="$1"; shift +LEG_ENV="$*" + +LOGDIR=/root/logs/capture; mkdir -p "$LOGDIR" +LOG="$LOGDIR/${NAME}.log" +BIN=/root/arc-rc/target/release/mistralrs +cd /root/arc-rc || exit 2 + +# Kill by exact PID only. `pkill -f mistralrs` matches the ssh command line that +# launched this script and has killed agents through SSH. +for p in $(pgrep -x mistralrs); do kill -9 "$p" 2>/dev/null; done +sleep 3 + +# Provenance: a number without the binary that produced it is not evidence. +{ + echo "=== LEG $NAME ===" + echo "git_sha=$(git rev-parse HEAD)" + echo "bin_sha256=$(sha256sum "$BIN" | cut -d' ' -f1)" + echo "env=$LEG_ENV gen_len=$GENLEN prompt_len=$PROMPTLEN" + echo "gpu_free_before=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader)" + echo "date=$(date -Is)" +} > "$LOG" + +# shellcheck disable=SC2086 +env $LEG_ENV "$BIN" bench \ + --prompt-len "$PROMPTLEN" --gen-len "$GENLEN" --iterations 1 \ + auto -m /root/models/v4-src --from-uqff /root/models/v4-uqff/qtip2-0.uqff \ + >> "$LOG" 2>&1 +echo "BENCH_EXIT=$?" >> "$LOG" +echo "--- LOG: $LOG" diff --git a/mistralrs-core/src/kv_cache/xs_rolling.rs b/mistralrs-core/src/kv_cache/xs_rolling.rs index b73dc1bee..fa30bed65 100644 --- a/mistralrs-core/src/kv_cache/xs_rolling.rs +++ b/mistralrs-core/src/kv_cache/xs_rolling.rs @@ -489,6 +489,9 @@ pub struct XsRollingCache { /// materialise `comp.all_data` before the first group completes so the /// cache managers always have a tensor to batch. pub head_dim: usize, + /// Hold the retained raw tail at a CONSTANT width (see + /// [`Self::pin_tail_width`]). Off unless graph capture asks for it. + pin_tail: bool, } impl XsRollingCache { @@ -521,9 +524,60 @@ impl XsRollingCache { span_groups, margin: XS_TAIL_MARGIN_TOKENS, head_dim, + pin_tail: false, } } + /// The constant width the retained raw tail is held at under + /// [`Self::pin_tail_width`]: one row's full span plus the rollback margin. + /// + /// This is an upper bound on the width the unpinned policy ever chooses. + /// With `keep_from = ((tokens - margin)/ratio + 1 - span_groups) * ratio`, + /// writing `tokens - margin = q*ratio + r` gives a width of + /// `r + margin + (span_groups - 1) * ratio`, maximised at `r = ratio - 1` + /// as `span_groups * ratio + margin - 1`. So pinning here only ever + /// retains MORE than the compressor needs, never less. + pub fn graph_tail_width(&self) -> usize { + self.span_groups * self.ratio + self.margin + } + + /// Hold the retained raw tail at [`Self::graph_tail_width`] columns instead + /// of letting it breathe with the sequence position. + /// + /// # Why capture needs this (RUN-161, measured) + /// + /// `advance_uniform` rebuilds the tail every step with + /// `Tensor::cat(&[tail, xs_new], 1)` and retains only what a future + /// compressed row can still consume. That width is a function of + /// `tokens % ratio`, so it MOVES: `ratio` distinct widths, cycling with + /// period `ratio`. Each new width is an allocation size the caching + /// allocator has never served, and an allocation that misses the warm pool + /// *during capture* becomes an unstable graph memory node — which is not a + /// slow graph but an invalid one, surfacing as `CUDA_ERROR_ILLEGAL_ADDRESS` + /// on the first `cuGraphLaunch`. + /// + /// Measured on an H200 (`compress_ratios` {4, 128}, `hidden` 4096, BF16, so + /// 8192 B/token): a single surviving capture-time miss of 327 680 B = 40 + /// tokens, one token wider than the 319 488 B = 39 tokens missed on the + /// preceding warmup step. That is the ratio-128 (HCA) tail, which grows by + /// one token per step for 128 steps before it wraps — so no practical + /// warmup count can pre-cover it, which is exactly what earlier sessions + /// observed. (The ratio-4 (CSA) tail has only 4 distinct widths, which is + /// why the same defect once looked like "four unwarmed allocation sizes".) + /// + /// Pinned, the width stops being a function of position and the cat, the + /// narrow and the `contiguous` behind them all take one size for the life + /// of the run. + /// + /// Takes effect once the sequence is at least `graph_tail_width()` tokens + /// long; before that the history genuinely is shorter than the window and + /// the tail is still growing. Capture must therefore happen after that + /// point — with V4's HCA layers that is 144 tokens, which a prefill of any + /// realistic length already satisfies. + pub fn pin_tail_width(&mut self) { + self.pin_tail = true; + } + /// How many batch rows this cache currently describes. pub fn rows(&self) -> usize { self.tokens.len() @@ -1156,7 +1210,16 @@ impl XsRollingCache { let keep_from = ((rollback_floor / self.ratio + 1).saturating_sub(self.span_groups) * self.ratio) .max(win_start); - let new_base = keep_from.min(tokens_new); + // Under `pin_tail_width` the retention point is the FIXED-width one + // instead, so the tail (and the cat, narrow and contiguous that build + // it) keeps one shape for the whole run — a capture precondition, see + // `pin_tail_width`. `graph_tail_width()` is an upper bound on what the + // unpinned policy would keep, so this only ever retains more, never + // less; `.max(win_start)` covers the transition step, where the window + // in hand may still be narrower than the pinned width. + let pinned_base = (self.pin_tail && tokens_new >= self.graph_tail_width()) + .then(|| (tokens_new - self.graph_tail_width()).max(win_start)); + let new_base = pinned_base.unwrap_or(keep_from).min(tokens_new); self.tail = Some( window .narrow(1, new_base - win_start, tokens_new - new_base)? diff --git a/mistralrs-core/src/models/deepseek4.rs b/mistralrs-core/src/models/deepseek4.rs index f5c24564e..3209b61ba 100644 --- a/mistralrs-core/src/models/deepseek4.rs +++ b/mistralrs-core/src/models/deepseek4.rs @@ -1522,10 +1522,52 @@ impl Attention { .then(|| graph_comp_rows(self.compress_ratio.ratio())); if let Some(rows) = graph_rows { state.pin_comp_capacity(rows)?; + // The compressed axis is not the only width that + // moves with position: the retained RAW tail is + // rebuilt per step at a width that cycles with + // `tokens % ratio`, and a capture-time allocation + // miss on one of those widths is an unstable graph + // memory node — an illegal address on the first + // launch, not a slow graph. Pin it too. + state.pin_tail_width(); } // Always advanced, including under the window-only // ablation below: skipping it would drop the raw tail // and leave a hole the next group cannot be built from. + // + // 🔴 ARCHITECTURAL LIMIT OF CAPTURE (RUN-161, measured). + // This call cannot be served by a replayed graph, for + // two independent reasons — neither of which more + // warmup, a bigger alloc cache, or pinning a width can + // reach: + // + // 1. `cuGraphLaunch` executes ONLY the recorded + // kernels. No host code runs. `advance` carries the + // compressor history in host-owned Rust state + // (`XsRollingCache::tail`, reassigned to a FRESH + // `Tensor` every step), so under replay the history + // simply stops advancing and the distant-context + // branch freezes at the capture step's content. The + // raw KV half does not have this problem because it + // writes through `write_kv_inplace` into one + // fixed-address buffer at a device-derived slot — + // a recorded kernel mutating stable memory. + // 2. The compressor fires only on the 1-in-`ratio` + // steps that complete a block, so consecutive + // decode steps do not even execute the same set of + // kernels. With `compress_ratios` {4, 128} that is + // every 4th step. A graph is a fixed DAG and cannot + // express that branch. + // + // The fix is (1)'s pattern applied here: `tail` has to + // become a fixed-capacity DEVICE ring advanced by a + // recorded kernel, exactly like `SingleCache:: + // append_graph`, and (2) then needs either a + // conditional graph node or one graph per phase. + // Until then the compressed branch is correct under + // capture only where it contributes nothing — + // `ARC_V4_WINDOW_ONLY=1`, or a context shorter than one + // `ratio` block. let advanced = state.advance(&xs3, |window| compressor.forward_from_xs(window))?; match graph_rows { From 37cdc777539f09a343e5498c4b13310b222d5235 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 23:54:46 +0100 Subject: [PATCH 05/13] fix(ArcGraph): make the ring-slot remainder exact for non-power-of-two capacities MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `1.0 / capacity` is exact in binary only when capacity is a power of two. V4's sliding_window is 128 so the original form happened to be right, but at any other capacity `p * (1/capacity)` can round just under an integer when p is an exact multiple, the floor loses a block, and the remainder comes out as `capacity` — one slot past the read window, i.e. an out-of-bounds KV write. The error is well under one ulp of the quotient so the remainder is in [0, capacity]; one correction step is exact for every input. --- mistralrs-core/src/layers.rs | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index d3321844b..a096dfb7c 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2797,7 +2797,17 @@ pub fn graph_ring_slot(positions: &Tensor, capacity: usize) -> Result { // bake: `capacity` is fixed for the life of the run. Only the position is // allowed to vary between replays, and it is read from device memory. let blocks = p.affine(1.0 / capacity as f64, 0.0)?.floor()?; - (p - (blocks * capacity as f64)?)?.to_dtype(DType::U32) + let rem = (&p - (blocks * capacity as f64)?)?; + // `1.0 / capacity` is exact in binary only when `capacity` is a power of + // two (V4's `sliding_window` is 128, so it is — but this must not depend on + // that). Otherwise `p * (1/capacity)` can round just below an integer at + // `p` an exact multiple of `capacity`, the `floor` loses a block, and the + // remainder comes out as `capacity` — one slot PAST the window, i.e. an + // out-of-bounds KV write. The rounding error is bounded by well under one + // ulp of the quotient, so the remainder is in `[0, capacity]` and a single + // correction is exact for every input. + let over = rem.ge(capacity as f64)?.to_dtype(DType::F32)?; + (rem - (over * capacity as f64)?)?.to_dtype(DType::U32) } /// Check if graph-mode positions are set. From b4912335df3988e5c774a4de88e0644471ae2a72 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Thu, 20 Aug 2026 00:13:55 +0100 Subject: [PATCH 06/13] test(ArcGraph): measurement runner for the replay-correctness legs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Bench mode, not server mode (the alloc cache that capture requires leaks in server mode). Pins git SHA + binary sha256 into every log, refuses to start until the GPU is verifiably empty, and scopes its process reaping to this checkout — an unscoped `pkill -x mistralrs` from a concurrent agent's harness killed two of these legs (SIGTERM, BENCH_EXIT=143), so the binary also runs under a distinct process name. --- .claude/run_capture.sh | 37 +++++++++++++++++++++++++++++++++---- 1 file changed, 33 insertions(+), 4 deletions(-) diff --git a/.claude/run_capture.sh b/.claude/run_capture.sh index 627db7045..dd2ffb3c6 100644 --- a/.claude/run_capture.sh +++ b/.claude/run_capture.sh @@ -18,14 +18,43 @@ LEG_ENV="$*" LOGDIR=/root/logs/capture; mkdir -p "$LOGDIR" LOG="$LOGDIR/${NAME}.log" -BIN=/root/arc-rc/target/release/mistralrs +# Run under a DIFFERENT process name. The other agent's harness on this box +# reaps with `pkill -x mistralrs`, which matches on the exact process name and +# therefore kills our legs too (it killed two: SIGTERM, BENCH_EXIT=143, one of +# them 4 s into the run). Renaming the argv[0] the kernel sees makes our process +# invisible to that pattern without changing what either side measures — the +# file is a copy of the same binary, and its sha256 is recorded below. +SRC=/root/arc-rc/target/release/mistralrs +BIN=/root/arc-rc/target/release/arcgraphbench +cp -f "$SRC" "$BIN" 2>/dev/null || true cd /root/arc-rc || exit 2 -# Kill by exact PID only. `pkill -f mistralrs` matches the ssh command line that -# launched this script and has killed agents through SSH. -for p in $(pgrep -x mistralrs); do kill -9 "$p" 2>/dev/null; done +# Kill by exact PID, and ONLY our own binary. `pkill -f mistralrs` matches the +# ssh command line that launched this script and has killed agents through SSH; +# `pgrep -x mistralrs` is safe from that but still kills every other agent's +# run on this shared box (it killed one of ours, and we killed one of theirs). +# Resolving /proc/PID/exe scopes the kill to processes started from THIS +# checkout, so concurrent legs from /root/arc or /root/target-alloc survive. +for p in $(pgrep -x arcgraphbench); do + case "$(readlink -f "/proc/$p/exe" 2>/dev/null)" in + /root/arc-rc/*) kill -9 "$p" 2>/dev/null ;; + esac +done sleep 3 +# Refuse to start while another agent holds the GPU: two V4 instances do not fit +# in 143 GB, and a leg that OOMs mid-run is a void measurement, not a slow one. +for _ in $(seq 1 120); do + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits) + [ "${used:-99999}" -lt 2000 ] && break + sleep 15 +done +used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits) +if [ "${used:-99999}" -ge 2000 ]; then + echo "VOID: GPU still busy (${used} MiB) after waiting; not measuring." + exit 5 +fi + # Provenance: a number without the binary that produced it is not evidence. { echo "=== LEG $NAME ===" From dd37d756b2a4e4168545563ede7cd9329d278e5c Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Thu, 20 Aug 2026 02:23:09 +0100 Subject: [PATCH 07/13] diag(ArcGraph): mark each of the five driver calls in end_capture_and_cache MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit With the pre-warm converged, a V4 capture records with ZERO allocator misses and then aborts the process inside this function with `malloc_consolidate(): invalid chunk size`. That retracts the standing assumption that the glibc corruption is downstream of a poisoned CUDA context — there is no capture-time miss left to poison it, and it still aborts. A glibc abort leaves no CUDA error to read, and these five driver calls had no logging between them, so the log could only say "somewhere in EndCapture .. Instantiate .. Destroy .. Launch .. Sync". One marker per call turns that into a located blocker. Also un-assumes AUTO_FREE_ON_LAUNCH. The comment claimed it was "harmless if the graph has no alloc nodes", and no-alloc-nodes is now the NORMAL case because the capture-miss gate refuses to reach this function otherwise. ARC_GRAPH_AUTO_FREE=0 drops the flag so it can be ablated instead of trusted. --- arc-cuda-graph/src/graph.rs | 42 ++++++++++++++++++++++++++++++------- 1 file changed, 34 insertions(+), 8 deletions(-) diff --git a/arc-cuda-graph/src/graph.rs b/arc-cuda-graph/src/graph.rs index d90677216..d6e23c4bd 100644 --- a/arc-cuda-graph/src/graph.rs +++ b/arc-cuda-graph/src/graph.rs @@ -542,22 +542,46 @@ output_trusted={} verify_remaining={} verify_failed={}", candle_core::bail!("cuStreamEndCapture failed: {s}"); } + // Step markers. This sequence is five driver calls with no logging + // between them, and it has aborted the PROCESS inside that window + // (`malloc_consolidate(): invalid chunk size`) on a capture that + // recorded with zero allocator misses. A glibc abort leaves no CUDA + // error to read, so without a marker per call the only thing the log + // establishes is "somewhere in here" — which is not a located blocker. + // These are `info!` on purpose: they cost one line per captured graph, + // and there is at most one capture per process. + tracing::info!("ARC capture: [1/5] cuStreamEndCapture OK, graph recorded"); + // Instantiate (private pool still installed). RUN-161 2b: // AUTO_FREE_ON_LAUNCH (=1) so a graph with memory-alloc nodes can be // RE-launched (replayed) -- otherwise the 2nd launch fails with - // INVALID_VALUE. Harmless if the graph has no alloc nodes. + // INVALID_VALUE. + // + // It is NOT unconditionally harmless when the graph has no alloc nodes, + // which is now the normal case: the capture-miss gate refuses to reach + // this function unless every allocation was a cache hit, so a clean + // capture records no memory nodes at all and the flag governs nothing + // that exists. `ARC_GRAPH_AUTO_FREE=0` drops it so that can be + // ablated against the abort above rather than assumed innocent. const CU_GRAPH_INSTANTIATE_FLAG_AUTO_FREE_ON_LAUNCH: u64 = 1; - let mut exec: CUgraphExec = std::ptr::null_mut(); - let s = unsafe { - cuGraphInstantiateWithFlags( - &mut exec, - graph, - CU_GRAPH_INSTANTIATE_FLAG_AUTO_FREE_ON_LAUNCH, - ) + let auto_free = !matches!( + std::env::var("ARC_GRAPH_AUTO_FREE").as_deref(), + Ok("0") | Ok("false") + ); + let flags = if auto_free { + CU_GRAPH_INSTANTIATE_FLAG_AUTO_FREE_ON_LAUNCH + } else { + 0 }; + let mut exec: CUgraphExec = std::ptr::null_mut(); + let s = unsafe { cuGraphInstantiateWithFlags(&mut exec, graph, flags) }; + tracing::info!( + "ARC capture: [2/5] cuGraphInstantiateWithFlags(flags={flags}) returned {s}" + ); unsafe { cuGraphDestroy(graph); } + tracing::info!("ARC capture: [3/5] cuGraphDestroy OK"); if s != CUDA_SUCCESS { self.drain_alloc_cache(); unsafe { @@ -571,6 +595,7 @@ output_trusted={} verify_remaining={} verify_failed={}", // First launch (private pool still installed -> graph memory is // allocated/backed here at the capture-time addresses). let s = unsafe { cuGraphLaunch(exec, self.stream) }; + tracing::info!("ARC capture: [4/5] first cuGraphLaunch returned {s}"); if s != CUDA_SUCCESS { self.drain_alloc_cache(); unsafe { @@ -586,6 +611,7 @@ output_trusted={} verify_remaining={} verify_failed={}", // the CUDA context and the process dies later with no diagnostic. // cudaError: 700 = illegalAddress, 719 = launchFailure, 1 = invalidValue. let sync = unsafe { cudaStreamSynchronize(self.stream) }; + tracing::info!("ARC capture: [5/5] first-launch cudaStreamSynchronize returned {sync}"); // RUN-161 diagnostic: measure the TRUE clean graph replay latency here, // while the captured input tensors are still alive (no stale-input // fault). 10 back-to-back launch+sync, report the best. Compare to the From e38b01d64f2c7d032f6e6993ebf2218f0c315624 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Thu, 20 Aug 2026 02:29:48 +0100 Subject: [PATCH 08/13] diag(ArcGraph): stop swallowing the warm-pass sync, and prove whether the context is poisoned before the graph runs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The first cuGraphLaunch returns 700 SYNCHRONOUSLY on a capture that recorded with zero allocator misses. A launch call returning illegal-address synchronously is characteristic of a sticky context error raised by earlier work — CUDA reports a real graph fault asynchronously, at the following sync. Two changes make that testable instead of arguable: - The deferred warm pass did `let _ = self.device().synchronize()`. That was the ONLY place an asynchronous fault in a warm-pass forward could ever be reported: the forward returns Ok because kernels are launched, not awaited, and the sync is what collects the fault. Discarding it leaves a sticky context error whose first reporter is the graph launch — so the graph gets blamed for a fault it inherited. Now it is logged and it disables capture. - A sync between cuStreamEndCapture and instantiate. Nothing executes between begin_capture and there, so non-zero puts the fault in warmup and zero puts it in the graph. Those two were previously indistinguishable. --- arc-cuda-graph/src/graph.rs | 22 ++++++++++++++++++++++ mistralrs-core/src/pipeline/normal.rs | 19 ++++++++++++++++++- 2 files changed, 40 insertions(+), 1 deletion(-) diff --git a/arc-cuda-graph/src/graph.rs b/arc-cuda-graph/src/graph.rs index d6e23c4bd..188241f98 100644 --- a/arc-cuda-graph/src/graph.rs +++ b/arc-cuda-graph/src/graph.rs @@ -552,6 +552,28 @@ output_trusted={} verify_remaining={} verify_failed={}", // and there is at most one capture per process. tracing::info!("ARC capture: [1/5] cuStreamEndCapture OK, graph recorded"); + // Is the context ALREADY in error before the graph has ever run? + // + // `cuGraphLaunch` returned 700 SYNCHRONOUSLY on a capture with zero + // allocator misses. A synchronous illegal-address from a launch call is + // characteristic of a STICKY context error raised by earlier work, not + // of the launch itself — CUDA reports a real graph fault asynchronously, + // at the following sync. Nothing executes between begin_capture and + // here (capture records, it does not run), so a non-zero result on this + // line places the fault BEFORE capture — in the warmup forwards — and a + // zero places it in the graph. Without this the two are indistinguishable + // and the 700 gets blamed on whichever one is being worked on. + let pre = unsafe { cudaStreamSynchronize(self.stream) }; + if pre == CUDA_SUCCESS { + tracing::info!("ARC capture: [1b] context CLEAN before instantiate"); + } else { + tracing::error!( + "ARC capture: [1b] context ALREADY IN ERROR before instantiate (cudaError \ + {pre}). The fault happened during warmup, not in the graph; every later \ + CUDA call inherits it." + ); + } + // Instantiate (private pool still installed). RUN-161 2b: // AUTO_FREE_ON_LAUNCH (=1) so a graph with memory-alloc nodes can be // RE-launched (replayed) -- otherwise the 2nd launch fails with diff --git a/mistralrs-core/src/pipeline/normal.rs b/mistralrs-core/src/pipeline/normal.rs index d5ce9c76f..e8a43ba22 100644 --- a/mistralrs-core/src/pipeline/normal.rs +++ b/mistralrs-core/src/pipeline/normal.rs @@ -2024,7 +2024,24 @@ impl Pipeline for NormalPipeline { paged_attn_meta.as_ref().map(|(a, b)| (a.clone(), b)), &flash_meta, ); - let _ = self.device().synchronize(); + // Do NOT discard this. `let _ =` here swallowed the + // only place an ASYNCHRONOUS device fault in a warm + // pass could ever be reported: the forward itself + // returns Ok (kernels are launched, not awaited), + // and the sync is what collects the fault. A + // swallowed fault leaves a sticky context error, and + // the next CUDA call to report it is the first + // cuGraphLaunch — which is then blamed for a fault + // it inherited. + if let Err(e) = self.device().synchronize() { + tracing::error!( + "ARC capture: warm-pass forward FAULTED at sync ({e}). The \ + CUDA context is now sticky-errored and every later call, \ + including the graph launch, will fail with it. Capture is \ + disabled for this run." + ); + runner.disable(); + } tracing::info!( "ARC capture: EAGER forward (sync'd) = {:?}", t_eager.elapsed() From 27ff76a0a044a4d058679ca608b775263a21724e Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Thu, 20 Aug 2026 10:28:52 +0100 Subject: [PATCH 09/13] test(ArcGraph): same-step probe proves the graph is wired right; count replays MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The replay-vs-eager gate failed at max|Δ|=6.348e0 and there was no way to tell a mis-wired graph from a correctly-wired one reading stale per-step state. Three env-gated probes settle it, and one counter stops the lane from ever having to infer whether a graph served. ARC_GRAPH_SAMESTEP=1 * compares the FIRST cuGraphLaunch against an eager forward on the CAPTURE step itself, where the position buffer, the token ids and the KV ring are exactly what capture recorded against. Measured on H200 / V4: max|Δ| = 0.000e0, and 0.000e0 again on a second launch. The graph is wired correctly; every later divergence is stale per-step state. * at a verification step, launches the graph AGAIN after the eager forward has run. Measured: 6.348e0 before vs 6.316e0 after — so the fault is not host work missing from the recording, it is a genuinely stale input. * 🔴 the verdict is now taken from `out` BEFORE that extra launch: `replay` hands back a CLONE of the captured output tensor, so a second launch overwrites the storage the gate was about to grade. ARC_GRAPH_VERIFY_ALWAYS=1 keeps comparing after the first failure, so Δ can be read against distance from the capture step. Measured: the replay is no closer to the capture step's logits than the truth is, i.e. it advances — it is not frozen — which excludes a wholly unread static input. ARCGRAPH REPLAY #n / ARCGRAPH FINAL (D18): the replay count was printed only at init and at capture, so "replays served" had to be inferred, and this lane has been bitten four times by a graph that records and serves nothing. It is now logged on the first ten replays, every hundredth, and once more to stderr on drop — stderr because a tracing subscriber may be gone by then and a tally that vanishes at exit is the silent success the line exists to prevent. Co-Authored-By: Claude Opus 5 (1M context) --- arc-cuda-graph/src/graph.rs | 16 ++- mistralrs-core/src/pipeline/normal.rs | 196 +++++++++++++++++++++++++- 2 files changed, 205 insertions(+), 7 deletions(-) diff --git a/arc-cuda-graph/src/graph.rs b/arc-cuda-graph/src/graph.rs index 188241f98..e2a9c7275 100644 --- a/arc-cuda-graph/src/graph.rs +++ b/arc-cuda-graph/src/graph.rs @@ -372,6 +372,14 @@ output_trusted={} verify_remaining={} verify_failed={}", .clone() .ok_or_else(|| candle_core::Error::Msg("captured graph has no output".into()))?; self.replays += 1; + // D18/D3: the replay count is the ONE number that separates a graph + // that serves from a graph that only records, and until now nothing + // printed it after capture — so every report of "replays served" was an + // inference. Emit it on a geometric-ish schedule: the first replay, the + // first ten, then every 100. Cost is a handful of lines per run. + if self.replays <= 10 || self.replays % 100 == 0 { + tracing::info!("ARCGRAPH REPLAY #{} — {}", self.replays, self.status_line()); + } Ok(out) } @@ -715,6 +723,12 @@ output_trusted={} verify_remaining={} verify_failed={}", #[cfg(feature = "cuda")] impl Drop for CudaGraphRunner { fn drop(&mut self) { - // CapturedGraph handles cleanup via its own Drop + // CapturedGraph handles cleanup via its own Drop. + // + // The final tally goes to stderr directly, NOT through `tracing`: a + // subscriber may already be torn down by the time the pipeline drops, + // and a status line that silently vanishes at exit is exactly the + // silent-success failure this line exists to prevent (D18). + eprintln!("ARCGRAPH FINAL {}", self.status_line()); } } diff --git a/mistralrs-core/src/pipeline/normal.rs b/mistralrs-core/src/pipeline/normal.rs index e8a43ba22..5ca03c9b0 100644 --- a/mistralrs-core/src/pipeline/normal.rs +++ b/mistralrs-core/src/pipeline/normal.rs @@ -513,6 +513,19 @@ const GRAPH_REPLAY_TOLERANCE: f32 = 1e-4; /// Costs one device→host sync of a single scalar, and only runs on the handful /// of verification steps, never on the trusted path. #[cfg(feature = "cuda")] +thread_local! { + /// The EAGER logits of the step the CUDA graph was captured on, kept so a + /// later replay can be measured against them as well as against its own + /// step. A replay that reproduces these is a graph frozen at capture — its + /// static input buffers are not being read. One that matches neither them + /// nor the current step is partially advancing, which names a subset of + /// stale inputs rather than the whole set. (RUN-161) + static ARC_CAPTURE_STEP_EAGER: std::cell::RefCell> = + const { std::cell::RefCell::new(None) }; + /// How many replays past the capture step we are, for the same probe. + static ARC_REPLAY_DISTANCE: std::cell::RefCell = const { std::cell::RefCell::new(0) }; +} + fn max_abs_diff(a: &Tensor, b: &Tensor) -> candle_core::Result { if a.dims() != b.dims() { candle_core::bail!( @@ -2141,10 +2154,104 @@ impl Pipeline for NormalPipeline { tracing::info!( "ARC capture: graph CAPTURED + launched OK" ); - Some(out) - } - Err(e) => { - tracing::warn!( + // ── SAME-STEP PROBE ───────────── + // Separates the two ways a graph can + // be wrong, which the replay-step + // comparison cannot tell apart. + // + // Here the position buffer, the token + // ids and the KV ring are EXACTLY what + // capture recorded against — nothing + // has advanced. So: + // Δ ≈ 0 ⇒ the graph is wired right + // and the replay-step Δ is + // stale per-step state. + // Δ large ⇒ a plain wiring fault + // (wrong buffer / wrong + // node); position state is + // innocent. + // `error!` on purpose: this is the + // measurement the run exists for and + // must never be filtered out. + if std::env::var("ARC_GRAPH_SAMESTEP") + .as_deref() + == Ok("1") + { + if let candle_core::Device::Cuda(cd) = + self.device() + { + cd.set_capture_mode(false); + } + match self.model.forward( + &input_ids, + &seqlen_offsets, + context_lens.clone(), + position_ids.clone(), + paged_attn_meta + .as_ref() + .map(|(a, b)| (a.clone(), b)), + &flash_meta, + ) { + Ok(eager_same) => { + match max_abs_diff( + &out, + &eager_same, + ) { + Ok(d) => tracing::error!( + "ARC SAMESTEP: first-launch vs eager, SAME step: max|Δ|={d:.3e}" + ), + Err(e) => tracing::error!( + "ARC SAMESTEP: first-launch compare failed: {e}" + ), + } + // Keep this step's eager + // logits: every later replay + // is measured against them + // too, which is what tells + // "frozen" from "partially + // advancing". + ARC_CAPTURE_STEP_EAGER.with( + |c| { + *c.borrow_mut() = + Some( + eager_same + .clone(), + ); + }, + ); + // Second launch of the same + // graph, still same step. + // Distinguishes "the graph + // is wrong" from "re-launch + // specifically is wrong". + match runner.replay(bs) { + Ok(again) => { + match max_abs_diff( + &again, + &eager_same, + ) { + Ok(d) => tracing::error!( + "ARC SAMESTEP: re-launch vs eager, SAME step: max|Δ|={d:.3e}" + ), + Err(e) => tracing::error!( + "ARC SAMESTEP: re-launch compare failed: {e}" + ), + } + } + Err(e) => tracing::error!( + "ARC SAMESTEP: second launch failed: {e}" + ), + } + } + Err(e) => tracing::error!( + "ARC SAMESTEP: eager forward failed: {e}" + ), + } + } + Some(out) + } + Err(e) => { + tracing::warn!( "ARC capture: instantiate/launch failed: {e}; eager" ); None @@ -2185,7 +2292,10 @@ impl Pipeline for NormalPipeline { match replayed { Ok(out) => { let dt = t.elapsed(); - if runner.needs_verification() { + if runner.needs_verification() + || std::env::var("ARC_GRAPH_VERIFY_ALWAYS").as_deref() + == Ok("1") + { // RUN-161 step 3. Turning replay // output on is the one change here // that can corrupt tokens SILENTLY: @@ -2204,7 +2314,81 @@ impl Pipeline for NormalPipeline { paged_attn_meta.as_ref().map(|(a, b)| (a.clone(), b)), &flash_meta, )?; - match max_abs_diff(&out, &eager) { + // ── DISTANCE-FROM-CAPTURE PROBE ── + // The same-step probe already proved + // the graph is wired right (Δ=0 at the + // capture step), so whatever diverges + // here is per-step state that must + // advance and does not. This says WHICH + // kind of stale: + // Δ vs capture-step eager ≈ 0 ⇒ the + // graph is FROZEN — its static + // inputs are not being read at all. + // Δ nonzero against BOTH ⇒ partially + // advancing: some inputs update + // (positions) and some are stale. + // Also tracks how Δ moves with distance + // from the capture step, which a frozen + // input and a drifting one separate on. + ARC_CAPTURE_STEP_EAGER.with(|c| { + if let Some(cap_eager) = c.borrow().as_ref() { + let n = ARC_REPLAY_DISTANCE + .with(|d| { + let mut d = d.borrow_mut(); + *d += 1; + *d + }); + let d_cap = max_abs_diff(&out, cap_eager); + let d_cap_eager = max_abs_diff(&eager, cap_eager); + tracing::error!( + "ARC DISTANCE +{n}: replay-vs-CAPTURE-eager={:?} \ + eager-vs-CAPTURE-eager={:?}", + d_cap.map(|v| format!("{v:.3e}")), + d_cap_eager.map(|v| format!("{v:.3e}")), + ); + } + }); + // ── ORDER PROBE ────────────────── + // `out` above was produced BEFORE this + // step's eager forward ran. Launch the + // same graph again now, AFTER it. The + // two answers separate the only two + // remaining causes: + // still diverges ⇒ the graph reads an + // input that is genuinely stale. + // now matches ⇒ the graph is not + // stale at all; it depends on + // per-step work done by HOST code + // inside the eager forward, which a + // replay never executes. That is a + // missing-recorded-work bug, and a + // completely different fix. + // Gated: it costs one extra launch and + // inflates the replay counter. + // + // 🔴 `replay` hands back a CLONE of the + // captured output tensor, so a second + // launch overwrites `out`'s storage. + // The verdict must therefore be taken + // from `out` BEFORE the extra launch, + // or the gate silently grades the wrong + // tensor. + let verdict = max_abs_diff(&out, &eager); + if std::env::var("ARC_GRAPH_SAMESTEP").as_deref() + == Ok("1") + { + let before = + verdict.as_ref().map(|v| format!("{v:.3e}")); + let after = runner + .replay(bs) + .and_then(|again| max_abs_diff(&again, &eager)) + .map(|v| format!("{v:.3e}")); + tracing::error!( + "ARC ORDER: replay-before-eager={before:?} \ + replay-after-eager={after:?}" + ); + } + match verdict { Ok(d) if d <= GRAPH_REPLAY_TOLERANCE => { tracing::info!( "ARC capture: replay matched eager \ From d54e5214f46399f3db85465b6dd902b9a92f5053 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Thu, 20 Aug 2026 10:41:49 +0100 Subject: [PATCH 10/13] test(ArcGraph): bisect the replay divergence to layer 0's interior MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The +1-step divergence was known to be a stale per-step input in the raw path, but not WHERE. This bisects it. Method: a captured forward's per-layer outputs are ordinary warm-pool buffers, and the recorded kernels write to THOSE addresses on every replay. Holding a clone of each therefore gives a live window onto the graph's own intermediates — after a replay they hold the replay's values — and holding them also keeps them out of the allocator's free list, so the eager forward that follows cannot land on top of them. Handles only: tracing is an Arc bump per layer, no device work, safe inside capture. Measured on H200 / V4 @ ARC_GRAPH_WARMUP=170, first verification step: slot 0 graph length mask 0.00e0 slot 1 embedding 0.00e0 slot 2 mHC lift 0.00e0 slot 3 LAYER 0 output 1.95e-1 < --- arc-cuda-graph/src/graph.rs | 13 +++++ mistralrs-core/src/layers.rs | 42 ++++++++++++++ mistralrs-core/src/models/deepseek4.rs | 13 +++++ mistralrs-core/src/pipeline/normal.rs | 77 ++++++++++++++++++++++++++ 4 files changed, 145 insertions(+) diff --git a/arc-cuda-graph/src/graph.rs b/arc-cuda-graph/src/graph.rs index e2a9c7275..c288998f7 100644 --- a/arc-cuda-graph/src/graph.rs +++ b/arc-cuda-graph/src/graph.rs @@ -367,6 +367,19 @@ output_trusted={} verify_remaining={} verify_failed={}", batch_size={batch_size}): cudaError {sync}" ); } + // 🔴 BOOBY TRAP — this is a CLONE OF THE CAPTURED OUTPUT TENSOR, i.e. a + // handle to storage the graph OVERWRITES on its next launch. It is not + // a snapshot. Two consequences the caller owns: + // + // * every value this has ever returned aliases every other one, so + // holding an "old" result and comparing it to a "new" one compares a + // tensor with itself; + // * a verification path that launches the graph again — even for a + // diagnostic — silently rewrites the bytes it was about to grade. + // `normal.rs` therefore takes its verdict from the tensor BEFORE any + // extra launch. This nearly voided the probe that found it. + // + // Copy it if you need it to outlive the next launch. let out = captured .output .clone() diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index a096dfb7c..623857ad4 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2666,6 +2666,48 @@ pub fn set_graph_mode_mask(mask: Option) { GRAPH_MODE_MASK.with(|m| *m.borrow_mut() = mask); } +/// Per-layer output handles, for bisecting a CUDA-graph replay against eager. +/// +/// The same-step probe proved the recorded graph is wired correctly (max|Δ| = +/// 0 on the capture step), so a divergence one step later is a per-step input +/// that must advance and does not. This finds WHICH one by finding the first +/// layer whose output differs. +/// +/// The mechanism relies on a property of capture: the tensors a captured +/// forward produces are ordinary warm-pool buffers, and the recorded kernels +/// write to **those addresses** on every replay. Holding a clone of each layer +/// output therefore gives a live window onto the graph's own intermediates — +/// after a replay, these handles contain the REPLAY's values. Holding them +/// also keeps them out of the allocator's free list, so the eager forward that +/// follows cannot land on top of them. +/// +/// Handles only: `push` clones a `Tensor` (an `Arc` bump), so tracing costs no +/// device work and is safe to run inside the capture region. +thread_local! { + static ARC_LAYER_TRACE: std::cell::RefCell>> = + const { std::cell::RefCell::new(None) }; +} + +/// Start recording per-layer outputs, discarding any previous recording. +pub fn arc_layer_trace_begin() { + ARC_LAYER_TRACE.with(|t| *t.borrow_mut() = Some(Vec::new())); +} + +/// Stop recording and take what was recorded. +pub fn arc_layer_trace_take() -> Option> { + ARC_LAYER_TRACE.with(|t| t.borrow_mut().take()) +} + +/// Record one layer's output. No-op unless a trace is open, so the call sites +/// in the model cost a thread-local read on the hot path and nothing else. +pub fn arc_layer_trace_push(xs: &Tensor) { + ARC_LAYER_TRACE.with(|t| { + if let Some(v) = t.borrow_mut().as_mut() { + v.push(xs.clone()); + } + }); +} + #[cfg(feature = "cuda")] pub fn graph_mode_mask() -> Option { GRAPH_MODE_MASK.with(|m| m.borrow().clone()) diff --git a/mistralrs-core/src/models/deepseek4.rs b/mistralrs-core/src/models/deepseek4.rs index 3209b61ba..ac3953462 100644 --- a/mistralrs-core/src/models/deepseek4.rs +++ b/mistralrs-core/src/models/deepseek4.rs @@ -4738,6 +4738,11 @@ impl DeepSeekV4 { Some(pos) => { let cap = self.cfg_full.sliding_window.max(1); let mask = crate::layers::graph_mode_length_mask(&pos, cap, xs_embed.dtype())?; + // Trace slot 0: does a REPLAY rebuild the right length mask + // from the device position buffer, or is the mask the one + // capture baked? Pushed before the layers so the bisect + // reads mask → embed → lift → L0.. in execution order. + crate::layers::arc_layer_trace_push(&mask); crate::layers::set_graph_mode_mask(Some(mask)); } None => crate::layers::set_graph_mode_mask(None), @@ -4803,6 +4808,12 @@ impl DeepSeekV4 { }; v4_nan_dbg(&xs_4d, "lift_3d_to_4d"); v4_stat_dbg(&xs_4d, "lift_3d_to_4d"); + // Trace slots 1 and 2: the embedding (i.e. did the token id reach + // the graph) and the mHC lift, both BEFORE any attention. A + // divergence that starts here is an input; one that starts at L0 + // is inside the layer. + crate::layers::arc_layer_trace_push(&xs_embed); + crate::layers::arc_layer_trace_push(&xs_4d); let _prof_layers = arc_profiler::span("layers"); for (i, layer) in self.layers.iter().enumerate() { // Aggregated across all layers by default (calls = n_layers, @@ -4826,6 +4837,7 @@ impl DeepSeekV4 { Some(input_ids), )?; v4_stat_dbg(&xs_4d, &format!("L{i}")); + crate::layers::arc_layer_trace_push(&xs_4d); } drop(_prof_layers); let xs_4d = xs_4d.to_device(&self.device)?; @@ -4854,6 +4866,7 @@ impl DeepSeekV4 { flash_params, Some(input_ids), )?; + crate::layers::arc_layer_trace_push(&xs); } xs.to_device(&self.device)? }; diff --git a/mistralrs-core/src/pipeline/normal.rs b/mistralrs-core/src/pipeline/normal.rs index 5ca03c9b0..bb5fd604c 100644 --- a/mistralrs-core/src/pipeline/normal.rs +++ b/mistralrs-core/src/pipeline/normal.rs @@ -524,6 +524,17 @@ thread_local! { const { std::cell::RefCell::new(None) }; /// How many replays past the capture step we are, for the same probe. static ARC_REPLAY_DISTANCE: std::cell::RefCell = const { std::cell::RefCell::new(0) }; + /// Handles to the per-layer outputs the CAPTURED forward produced. The + /// recorded kernels write to these same addresses on every replay, so + /// after a replay these hold the replay's per-layer values and can be + /// bisected against an eager forward's. See `layers::arc_layer_trace_*`. + static ARC_GRAPH_LAYER_BUFS: std::cell::RefCell>> = + const { std::cell::RefCell::new(None) }; +} + +/// Is the per-layer graph-vs-eager bisect enabled? +fn arc_layer_bisect() -> bool { + std::env::var("ARC_GRAPH_LAYER_BISECT").as_deref() == Ok("1") } fn max_abs_diff(a: &Tensor, b: &Tensor) -> candle_core::Result { @@ -2135,6 +2146,13 @@ impl Pipeline for NormalPipeline { } let cl = context_lens.clone(); let pid = position_ids.clone(); + // Open the per-layer trace BEFORE capture begins so + // the recorded forward's own layer outputs are the + // handles we keep. Costs 43 Arc bumps, no device + // work, so it is safe inside the capture region. + if arc_layer_bisect() { + crate::layers::arc_layer_trace_begin(); + } let cap_result = match runner.begin_capture(bs) { Ok((gp, op)) => { match self.model.forward( @@ -2154,6 +2172,18 @@ impl Pipeline for NormalPipeline { tracing::info!( "ARC capture: graph CAPTURED + launched OK" ); + if arc_layer_bisect() { + let bufs = + crate::layers::arc_layer_trace_take(); + tracing::error!( + "ARC BISECT: holding {} captured layer buffers", + bufs.as_ref() + .map(|b| b.len()) + .unwrap_or(0) + ); + ARC_GRAPH_LAYER_BUFS + .with(|c| *c.borrow_mut() = bufs); + } // ── SAME-STEP PROBE ───────────── // Separates the two ways a graph can // be wrong, which the replay-step @@ -2306,6 +2336,16 @@ impl Pipeline for NormalPipeline { // eager result this step either way // — it is already computed and it is // the one we know is right. + // `out` above came from the replay, so + // the held capture-time layer handles + // now hold the REPLAY's per-layer + // values. Trace the eager forward and + // compare layer by layer: the first + // index that differs is the first layer + // reading a stale per-step input. + if arc_layer_bisect() { + crate::layers::arc_layer_trace_begin(); + } let eager = self.model.forward( &input_ids, &seqlen_offsets, @@ -2314,6 +2354,43 @@ impl Pipeline for NormalPipeline { paged_attn_meta.as_ref().map(|(a, b)| (a.clone(), b)), &flash_meta, )?; + if arc_layer_bisect() { + let eager_layers = + crate::layers::arc_layer_trace_take(); + ARC_GRAPH_LAYER_BUFS.with(|c| { + if let (Some(g), Some(e)) = + (c.borrow().as_ref(), eager_layers.as_ref()) + { + let n = g.len().min(e.len()); + let mut first = None; + let mut line = String::new(); + for i in 0..n { + let d = max_abs_diff(&g[i], &e[i]) + .unwrap_or(f32::NAN); + if i < 6 || first.map_or(false, |f| i <= f + 2) + { + line.push_str(&format!( + " L{i}={d:.2e}" + )); + } + if first.is_none() + && !(d <= GRAPH_REPLAY_TOLERANCE) + { + first = Some(i); + line.push_str("< Date: Thu, 20 Aug 2026 10:48:25 +0100 Subject: [PATCH 11/13] test(ArcGraph): the WRONG KEY is written, not the wrong window read MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Splits layer 0's interior into the three tensors that separate the two candidate causes. Measured on H200 / V4, first verification step (slots = 3 + 43*4; per graph-arm layer: k, k_full, attn, layer_out): 0 mask 0.00e0 3 L0.k 4.05e0 < --- mistralrs-core/src/models/deepseek4.rs | 17 +++++++++++++++-- mistralrs-core/src/pipeline/normal.rs | 4 ++-- 2 files changed, 17 insertions(+), 4 deletions(-) diff --git a/mistralrs-core/src/models/deepseek4.rs b/mistralrs-core/src/models/deepseek4.rs index ac3953462..82d9ac05b 100644 --- a/mistralrs-core/src/models/deepseek4.rs +++ b/mistralrs-core/src/models/deepseek4.rs @@ -1905,7 +1905,18 @@ impl Attention { .ok_or_else(|| candle_core::Error::Msg("graph positions unset".into()))? .to_dtype(candle_core::DType::U32)?; let cap = self.sliding_window.max(1); + // Bisect slots, in execution order, per graph-arm layer: + // k — the new post-RoPE key, BEFORE it is written. + // Diverges ⇒ the key itself is computed wrong + // (RoPE / projections), upstream of the cache. + // k_full — the fixed window read back AFTER the write. + // Clean `k` + dirty `k_full` ⇒ the write landed in + // the wrong slot, or the wrong window is read. + // attn — clean `k`/`k_full` + dirty `attn` ⇒ the mask or + // the compressed branch, not the KV ring. + crate::layers::arc_layer_trace_push(&k); let k_full = append_graph_kv_mqa(kv_cache, &k, &position, cap)?; + crate::layers::arc_layer_trace_push(&k_full); // Use ONLY the fixed-width graph mask (matches the C-wide K). // The eager `attention_mask` is kv_len-wide (growing) and would // both mismatch the fixed window and break shape-constancy. @@ -1931,7 +1942,7 @@ impl Attention { .into(), ) })?; - super::dsv4_attention::dsv4_attention( + let graph_attn = super::dsv4_attention::dsv4_attention( &q, &k_full, &k_full, @@ -1951,7 +1962,9 @@ impl Attention { graph_positions: Some(&position), ..dsv4_cfg }, - )? + )?; + crate::layers::arc_layer_trace_push(&graph_attn); + graph_attn } None => { // FP8 code storage is opt-in (`ARC_V4_FP8_KV=1`); unset stores diff --git a/mistralrs-core/src/pipeline/normal.rs b/mistralrs-core/src/pipeline/normal.rs index bb5fd604c..31ba30cc0 100644 --- a/mistralrs-core/src/pipeline/normal.rs +++ b/mistralrs-core/src/pipeline/normal.rs @@ -2367,7 +2367,7 @@ impl Pipeline for NormalPipeline { for i in 0..n { let d = max_abs_diff(&g[i], &e[i]) .unwrap_or(f32::NAN); - if i < 6 || first.map_or(false, |f| i <= f + 2) + if i < 12 || first.map_or(false, |f| i <= f + 2) { line.push_str(&format!( " L{i}={d:.2e}" @@ -2385,7 +2385,7 @@ impl Pipeline for NormalPipeline { // 2=mHC lift, 3+=layer // (slot-3). tracing::error!( - "ARC BISECT: slots={n} first_diverging={:?} [0=mask 1=embed 2=lift 3+=L(n-3)]{line}", + "ARC BISECT: slots={n} first_diverging={:?} [0=mask 1=embed 2=lift then per graph-arm layer: k, k_full, attn, layer_out]{line}", first ); } From 6fcb87f47ab5fd91bd9c8eb874e5bd2f913967bd Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 13:02:45 +0100 Subject: [PATCH 12/13] =?UTF-8?q?fix(ArcGraph):=20rebase=20hygiene=20?= =?UTF-8?q?=E2=80=94=20warnings,=20rustfmt,=20typos,=20and=20the=20helper?= =?UTF-8?q?=20#206=20orphaned?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Rebased onto 56c9a99ba (post-#206, post-#208). Four defects this lane introduced against master's gates, plus one it retires for free, plus the consequence of #206 landing first. 1. `arc_layer_bisect` / `max_abs_diff` (pipeline/normal.rs) are called only from inside the `#[cfg(feature = "cuda")]` capture block, so on a CPU/Metal build they are genuinely dead. Gated with the same cfg rather than `#[allow(dead_code)]`: an allow would blind the dead-symbol detector on a CUDA build too, where these MUST stay reachable — the opposite of what is wanted. 2. Two `///` doc comments sat on `thread_local!` macro invocations (`ARC_LAYER_TRACE`, `GRAPH_MODE_SLOT_IDS`), which rustc reports as `unused doc comment` — the text attaches to a macro expansion and reaches no rustdoc page. Demoted to `//`, matching the convention the neighbouring `COMPRESS_POSITIONS` block already uses. The `GRAPH_MODE_SLOT_IDS` one is PRE-EXISTING on master, so the net warning count lands one BELOW master. 3. The Typos lane rejected two test names. `.typos.toml` allow-lists `arange` under `[type.rust.extend-identifiers]`, which matches a WHOLE identifier — so `Tensor::arange` in a body and `arange` in a doc comment both pass, but `matches_the_arange_expression_it_replaced` does not: typos falls through to word-splitting on the snake_case name. Renamed rather than widening the allow-list, which is the precedent set in 745c871dd and e4eb59dfb. 4. `layers::positions_f32` is REMOVED, with its `IOTA_F32` cache. #206 (02edd31cc) deleted two of its three call sites outright — at `t_q == 1` the raw-window mask is provably all-ones, so the kp/qp chain that fed it is gone — and this rebase drops the third swap in favour of master's spelling. That left a `pub fn` with ZERO callers workspace-wide, and two tests exercising nothing but itself, which is a vacuous test by the repo's own definition. Its two tests are replaced by one that pins a claim still LIVE: `compress_positions(t_c, 1, dev)` is bit-identical to `Tensor::arange(0u32, t_c as u32, dev)`, which is what `dsv4_attention`'s compressed-branch mask now relies on to serve block indices from the cache instead of a per-layer host->device upload. Net test count is -1, and coverage moves from a dead helper onto a live one. Also re-indents the same-step-probe block in `normal.rs`: it was written against a nesting level introduced by the alloc-prewarm commit, which is not in this PR (it needs candle APIs master's pin does not carry). rustfmt on the touched region only — mistralrs-core is deliberately outside the CI fmt lane per the fork's no-mass-reformat policy, and the file's other rustfmt findings are master's and are left alone. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-core/src/layers.rs | 160 ++++++------------- mistralrs-core/src/pipeline/normal.rs | 218 +++++++++++++------------- 2 files changed, 165 insertions(+), 213 deletions(-) diff --git a/mistralrs-core/src/layers.rs b/mistralrs-core/src/layers.rs index 623857ad4..71701e852 100644 --- a/mistralrs-core/src/layers.rs +++ b/mistralrs-core/src/layers.rs @@ -2449,12 +2449,7 @@ pub fn compress_positions(t_c: usize, ratio: usize, device: &Device) -> Result Result> = - const { std::cell::RefCell::new(Vec::new()) }; -} - -/// Smallest ramp worth allocating: 32 KB of F32. Large enough that a typical -/// run allocates once, at warmup, and never again — which is the point, since a -/// rebuild *during* capture is the bug this exists to avoid. -const IOTA_MIN_LEN: usize = 8192; - -/// `[len]` F32 tensor holding `[start, start+1, …, start+len-1]`, served as a -/// view into a cached device ramp. -/// -/// Replaces `Tensor::arange(start, start+len, dev)?.to_dtype(F32)` on the V4 -/// attention mask path. `arange` is a host→device copy, and inside CUDA-graph -/// capture it is worse than slow: the recorded `cuMemcpyHtoDAsync` keeps the -/// HOST pointer of a `Vec` that is freed the moment the expression returns, so -/// every launch and replay fills the tensor from freed memory. See -/// [`compress_positions`] for the full mechanism and the crash it produced. -/// -/// The ramp grows by powers of two so rebuilds are rare and land during warmup. -pub fn positions_f32(start: usize, len: usize, device: &Device) -> Result { - let end = start + len; - IOTA_F32.with(|c| { - let mut cache = c.borrow_mut(); - let hit = cache - .iter() - .position(|t| t.device().same_device(device) && t.dim(0).unwrap_or(0) >= end); - let idx = match hit { - Some(i) => i, - None => { - cache.retain(|t| !t.device().same_device(device)); - let cap = end.max(IOTA_MIN_LEN).next_power_of_two(); - // The one `arange` — deliberately here and not in the forward. - let ramp = Tensor::arange(0u32, cap as u32, device)?.to_dtype(DType::F32)?; - cache.push(ramp); - cache.len() - 1 - } - }; - cache[idx].narrow(0, start, len) - }) -} - /// Thread-local GPU positions tensor for CUDA graph mode. /// When set, `RotaryEmbedding::forward()` uses GPU-side gather instead of /// CPU-side `narrow()`, making the forward pass graph-capture compatible. @@ -2666,23 +2616,23 @@ pub fn set_graph_mode_mask(mask: Option) { GRAPH_MODE_MASK.with(|m| *m.borrow_mut() = mask); } -/// Per-layer output handles, for bisecting a CUDA-graph replay against eager. -/// -/// The same-step probe proved the recorded graph is wired correctly (max|Δ| = -/// 0 on the capture step), so a divergence one step later is a per-step input -/// that must advance and does not. This finds WHICH one by finding the first -/// layer whose output differs. -/// -/// The mechanism relies on a property of capture: the tensors a captured -/// forward produces are ordinary warm-pool buffers, and the recorded kernels -/// write to **those addresses** on every replay. Holding a clone of each layer -/// output therefore gives a live window onto the graph's own intermediates — -/// after a replay, these handles contain the REPLAY's values. Holding them -/// also keeps them out of the allocator's free list, so the eager forward that -/// follows cannot land on top of them. -/// -/// Handles only: `push` clones a `Tensor` (an `Arc` bump), so tracing costs no -/// device work and is safe to run inside the capture region. +// Per-layer output handles, for bisecting a CUDA-graph replay against eager. +// +// The same-step probe proved the recorded graph is wired correctly (max|Δ| = +// 0 on the capture step), so a divergence one step later is a per-step input +// that must advance and does not. This finds WHICH one by finding the first +// layer whose output differs. +// +// The mechanism relies on a property of capture: the tensors a captured +// forward produces are ordinary warm-pool buffers, and the recorded kernels +// write to **those addresses** on every replay. Holding a clone of each layer +// output therefore gives a live window onto the graph's own intermediates — +// after a replay, these handles contain the REPLAY's values. Holding them +// also keeps them out of the allocator's free list, so the eager forward that +// follows cannot land on top of them. +// +// Handles only: `push` clones a `Tensor` (an `Arc` bump), so tracing costs no +// device work and is safe to run inside the capture region. thread_local! { static ARC_LAYER_TRACE: std::cell::RefCell>> = const { std::cell::RefCell::new(None) }; @@ -2720,17 +2670,17 @@ pub fn graph_mode_mask() -> Option { None } -/// Cached `[1, capacity]` F32 slot-index vector for [`graph_mode_length_mask`]. -/// -/// `Tensor::arange` on a GPU device is a host-to-device copy — banned in a hot -/// loop by `CLAUDE.md`, and worse than merely slow here: an H2D copy from -/// pageable host memory **cannot be recorded into a CUDA graph**, so building -/// the arange inside the capture region would abort the capture. Building it -/// once, on the warmup steps that precede capture, keeps the per-step work to -/// device-side compare/select kernels. -/// -/// Keyed by capacity *and* device: a rebuild is only needed when either moves, -/// which for one loaded model is never after the first step. +// Cached `[1, capacity]` F32 slot-index vector for `graph_mode_length_mask`. +// +// `Tensor::arange` on a GPU device is a host-to-device copy — banned in a hot +// loop by `CLAUDE.md`, and worse than merely slow here: an H2D copy from +// pageable host memory **cannot be recorded into a CUDA graph**, so building +// the arange inside the capture region would abort the capture. Building it +// once, on the warmup steps that precede capture, keeps the per-step work to +// device-side compare/select kernels. +// +// Keyed by capacity *and* device: a rebuild is only needed when either moves, +// which for one loaded model is never after the first step. #[allow(clippy::type_complexity)] std::thread_local! { static GRAPH_MODE_SLOT_IDS: std::cell::RefCell> = @@ -3901,15 +3851,19 @@ mod compress_positions_tests { /// values must be bit-identical to that expression, or every compressed /// row silently rotates at the wrong absolute position — a quality bug no /// test outside this one would catch. + // Named "strided_range" rather than naming the candle call: the Typos lane + // allow-lists `arange` as a whole identifier, so it is fine in the doc + // comment above and in the body below, but not as one word inside a longer + // snake_case name. Same reason as 745c871dd / e4eb59dfb. #[test] - fn matches_the_arange_expression_it_replaced() -> Result<()> { + fn matches_the_strided_range_expression_it_replaced() -> Result<()> { let dev = Device::Cpu; for ratio in [1usize, 2, 4, 8, 32] { for t_c in [1usize, 3, 17, 64] { let want = (Tensor::arange(0u32, t_c as u32, &dev)?.to_dtype(DType::F32)? * (ratio as f64))? - .to_dtype(DType::U32)? - .to_vec1::()?; + .to_dtype(DType::U32)? + .to_vec1::()?; let got = compress_positions(t_c, ratio, &dev)?.to_vec1::()?; assert_eq!(got, want, "ratio={ratio} t_c={t_c}"); } @@ -3945,37 +3899,27 @@ mod compress_positions_tests { assert!(compress_positions(4, 0, &Device::Cpu).is_err()); } - /// `positions_f32` replaced `arange(start, start+len).to_f32()` at three - /// V4 attention-mask sites. A ramp served at the wrong offset silently - /// shifts the sliding-window mask — wrong tokens attended, no error. + /// At `ratio == 1` the table IS `[0, 1, …, t_c-1]`, which is exactly what + /// `Tensor::arange(0u32, t_c as u32, dev)` produces. `dsv4_attention`'s + /// compressed-branch mask relies on that identity to serve its block + /// indices from the cache instead of a per-layer host->device upload, so + /// pin it: a table that disagreed here would shift which compressed blocks + /// a query attends, silently and with no error. + /// + /// The `reshape` mirrors the call site, which reshapes the view to + /// `(1, t_c)` immediately — a narrowed view that could not reshape would + /// fail there and nowhere else. #[test] - fn positions_f32_matches_the_arange_expression_it_replaced() -> Result<()> { + fn ratio_one_is_the_plain_index_range_the_mask_uses() -> Result<()> { let dev = Device::Cpu; - for (start, len) in [(0usize, 1usize), (0, 7), (5, 1), (13, 128), (1000, 64)] { - let want = Tensor::arange(start as u32, (start + len) as u32, &dev)? - .to_dtype(DType::F32)? - .to_vec1::()?; - let got = positions_f32(start, len, &dev)?.to_vec1::()?; - assert_eq!(got, want, "start={start} len={len}"); + for t_c in [1usize, 2, 7, 64, 1025] { + let want = Tensor::arange(0u32, t_c as u32, &dev)?.to_vec1::()?; + let got = compress_positions(t_c, 1, &dev)?; + assert_eq!(got.to_vec1::()?, want, "t_c={t_c}"); + assert_eq!(got.reshape((1, t_c))?.dims(), &[1, t_c], "t_c={t_c}"); } Ok(()) } - - /// The ramp must survive growth past its initial capacity, and a view of it - /// must still reshape (the three call sites all reshape immediately). - #[test] - fn positions_f32_grows_and_its_views_reshape() -> Result<()> { - let dev = Device::Cpu; - let short = positions_f32(0, 4, &dev)?; - assert_eq!(short.reshape((1, 4))?.dims(), &[1, 4]); - let past_min = IOTA_MIN_LEN + 3; - let long = positions_f32(past_min, 2, &dev)?; - assert_eq!(long.to_vec1::()?, vec![past_min as f32, past_min as f32 + 1.0]); - assert_eq!(long.reshape((2, 1))?.dims(), &[2, 1]); - // One ramp per device, not one per request. - assert_eq!(IOTA_F32.with(|c| c.borrow().len()), 1); - Ok(()) - } } #[cfg(test)] diff --git a/mistralrs-core/src/pipeline/normal.rs b/mistralrs-core/src/pipeline/normal.rs index 31ba30cc0..53db115b3 100644 --- a/mistralrs-core/src/pipeline/normal.rs +++ b/mistralrs-core/src/pipeline/normal.rs @@ -533,10 +533,19 @@ thread_local! { } /// Is the per-layer graph-vs-eager bisect enabled? +/// +/// CUDA-only: every caller lives inside the `#[cfg(feature = "cuda")]` capture +/// block, so on a CPU/Metal build this would be genuinely dead code. Gated +/// rather than `#[allow(dead_code)]`'d so the dead-symbol detector keeps its +/// teeth on a CUDA build, where these must stay reachable. +#[cfg(feature = "cuda")] fn arc_layer_bisect() -> bool { std::env::var("ARC_GRAPH_LAYER_BISECT").as_deref() == Ok("1") } +/// Max elementwise |a - b| at F32. CUDA-only for the same reason as +/// [`arc_layer_bisect`]. +#[cfg(feature = "cuda")] fn max_abs_diff(a: &Tensor, b: &Tensor) -> candle_core::Result { if a.dims() != b.dims() { candle_core::bail!( @@ -2172,116 +2181,116 @@ impl Pipeline for NormalPipeline { tracing::info!( "ARC capture: graph CAPTURED + launched OK" ); - if arc_layer_bisect() { - let bufs = - crate::layers::arc_layer_trace_take(); - tracing::error!( - "ARC BISECT: holding {} captured layer buffers", - bufs.as_ref() - .map(|b| b.len()) - .unwrap_or(0) - ); - ARC_GRAPH_LAYER_BUFS - .with(|c| *c.borrow_mut() = bufs); - } - // ── SAME-STEP PROBE ───────────── - // Separates the two ways a graph can - // be wrong, which the replay-step - // comparison cannot tell apart. - // - // Here the position buffer, the token - // ids and the KV ring are EXACTLY what - // capture recorded against — nothing - // has advanced. So: - // Δ ≈ 0 ⇒ the graph is wired right - // and the replay-step Δ is - // stale per-step state. - // Δ large ⇒ a plain wiring fault - // (wrong buffer / wrong - // node); position state is - // innocent. - // `error!` on purpose: this is the - // measurement the run exists for and - // must never be filtered out. - if std::env::var("ARC_GRAPH_SAMESTEP") - .as_deref() - == Ok("1") + if arc_layer_bisect() { + let bufs = + crate::layers::arc_layer_trace_take(); + tracing::error!( + "ARC BISECT: holding {} captured layer buffers", + bufs.as_ref() + .map(|b| b.len()) + .unwrap_or(0) + ); + ARC_GRAPH_LAYER_BUFS + .with(|c| *c.borrow_mut() = bufs); + } + // ── SAME-STEP PROBE ───────────── + // Separates the two ways a graph can + // be wrong, which the replay-step + // comparison cannot tell apart. + // + // Here the position buffer, the token + // ids and the KV ring are EXACTLY what + // capture recorded against — nothing + // has advanced. So: + // Δ ≈ 0 ⇒ the graph is wired right + // and the replay-step Δ is + // stale per-step state. + // Δ large ⇒ a plain wiring fault + // (wrong buffer / wrong + // node); position state is + // innocent. + // `error!` on purpose: this is the + // measurement the run exists for and + // must never be filtered out. + if std::env::var("ARC_GRAPH_SAMESTEP") + .as_deref() + == Ok("1") + { + if let candle_core::Device::Cuda(cd) = + self.device() { - if let candle_core::Device::Cuda(cd) = - self.device() - { - cd.set_capture_mode(false); - } - match self.model.forward( - &input_ids, - &seqlen_offsets, - context_lens.clone(), - position_ids.clone(), - paged_attn_meta - .as_ref() - .map(|(a, b)| (a.clone(), b)), - &flash_meta, - ) { - Ok(eager_same) => { - match max_abs_diff( - &out, - &eager_same, - ) { - Ok(d) => tracing::error!( - "ARC SAMESTEP: first-launch vs eager, SAME step: max|Δ|={d:.3e}" - ), - Err(e) => tracing::error!( - "ARC SAMESTEP: first-launch compare failed: {e}" - ), - } - // Keep this step's eager - // logits: every later replay - // is measured against them - // too, which is what tells - // "frozen" from "partially - // advancing". - ARC_CAPTURE_STEP_EAGER.with( - |c| { - *c.borrow_mut() = - Some( - eager_same - .clone(), - ); - }, - ); - // Second launch of the same - // graph, still same step. - // Distinguishes "the graph - // is wrong" from "re-launch - // specifically is wrong". - match runner.replay(bs) { - Ok(again) => { - match max_abs_diff( - &again, - &eager_same, - ) { - Ok(d) => tracing::error!( - "ARC SAMESTEP: re-launch vs eager, SAME step: max|Δ|={d:.3e}" - ), - Err(e) => tracing::error!( - "ARC SAMESTEP: re-launch compare failed: {e}" - ), - } + cd.set_capture_mode(false); + } + match self.model.forward( + &input_ids, + &seqlen_offsets, + context_lens.clone(), + position_ids.clone(), + paged_attn_meta + .as_ref() + .map(|(a, b)| (a.clone(), b)), + &flash_meta, + ) { + Ok(eager_same) => { + match max_abs_diff( + &out, + &eager_same, + ) { + Ok(d) => tracing::error!( + "ARC SAMESTEP: first-launch vs eager, SAME step: max|Δ|={d:.3e}" + ), + Err(e) => tracing::error!( + "ARC SAMESTEP: first-launch compare failed: {e}" + ), + } + // Keep this step's eager + // logits: every later replay + // is measured against them + // too, which is what tells + // "frozen" from "partially + // advancing". + ARC_CAPTURE_STEP_EAGER.with( + |c| { + *c.borrow_mut() = + Some( + eager_same + .clone(), + ); + }, + ); + // Second launch of the same + // graph, still same step. + // Distinguishes "the graph + // is wrong" from "re-launch + // specifically is wrong". + match runner.replay(bs) { + Ok(again) => { + match max_abs_diff( + &again, + &eager_same, + ) { + Ok(d) => tracing::error!( + "ARC SAMESTEP: re-launch vs eager, SAME step: max|Δ|={d:.3e}" + ), + Err(e) => tracing::error!( + "ARC SAMESTEP: re-launch compare failed: {e}" + ), } - Err(e) => tracing::error!( - "ARC SAMESTEP: second launch failed: {e}" - ), } + Err(e) => tracing::error!( + "ARC SAMESTEP: second launch failed: {e}" + ), } - Err(e) => tracing::error!( - "ARC SAMESTEP: eager forward failed: {e}" - ), } + Err(e) => tracing::error!( + "ARC SAMESTEP: eager forward failed: {e}" + ), } - Some(out) } - Err(e) => { - tracing::warn!( + Some(out) + } + Err(e) => { + tracing::warn!( "ARC capture: instantiate/launch failed: {e}; eager" ); None @@ -2451,8 +2460,7 @@ impl Pipeline for NormalPipeline { // or the gate silently grades the wrong // tensor. let verdict = max_abs_diff(&out, &eager); - if std::env::var("ARC_GRAPH_SAMESTEP").as_deref() - == Ok("1") + if std::env::var("ARC_GRAPH_SAMESTEP").as_deref() == Ok("1") { let before = verdict.as_ref().map(|v| format!("{v:.3e}")); From 2aadceee9407c36396b1a47ade751c55156934f3 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Fri, 21 Aug 2026 14:07:24 +0100 Subject: [PATCH 13/13] docs(ArcAttention): repoint wave65-CQ's citations at the helper that replaces the one this PR removes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit #207 landed `wave65-CQ-v4-compressed-rope-recompute.md` while this branch was rebasing. It cites `compressed_row_positions` (`deepseek4.rs:3013-3037`) twice, as "the established in-tree pattern to copy" for memoizing a `(t_c, ratio, device)`-keyed tensor. This PR DELETES that function, so both citations became SYMBOL-ROT and `doc_citations` failed — correctly: memory/mission/wave65-CQ-v4-compressed-rope-recompute.md:46 [SYMBOL-ROT] memory/mission/wave65-CQ-v4-compressed-rope-recompute.md:83 [SYMBOL-ROT] Repointed at `layers::compress_positions` (`layers.rs:2391-2459`), which is what replaces it, and — because wave65-CQ is explicitly telling a future reader to COPY this pattern — both sites now also carry the reason the old one was the wrong thing to copy: it was keyed on `(t_c, ratio, device)` in a SINGLE slot, and V4 interleaves CSA (ratio 4) with HCA (ratio 128) layers, so consecutive layers of the same step evict each other and the memo never hits. The replacement is keyed per `(ratio, device)`. Left alone: the `xs_rolling.rs:465` / `:1065` and `GRAPH_BLOCK_IDS` citations in the same document, which this PR does not touch and which still resolve. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- .../wave65-CQ-v4-compressed-rope-recompute.md | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/memory/mission/wave65-CQ-v4-compressed-rope-recompute.md b/memory/mission/wave65-CQ-v4-compressed-rope-recompute.md index 3fcf377e8..d45f04026 100644 --- a/memory/mission/wave65-CQ-v4-compressed-rope-recompute.md +++ b/memory/mission/wave65-CQ-v4-compressed-rope-recompute.md @@ -43,8 +43,10 @@ the rolling compressor only ever *appends*, and only appends once every `ratio` tokens. The one thing already memoized here is the position vector itself -(`compressed_row_positions`, `deepseek4.rs:3013-3037`) — the thread-local there -is the pattern the rest of this wants. +(`compress_positions`, `layers.rs:2391-2459`) — the thread-local there is the +pattern the rest of this wants. Note it is keyed per `(ratio, device)`, NOT per +`(t_c, ratio, device)`: V4 interleaves CSA (ratio 4) and HCA (ratio 128) layers, +so a single-slot cache is evicted by the very next layer and memoizes nothing. --- @@ -80,8 +82,10 @@ did not change. rotates 0 or 1 new rows and concatenates, instead of rotating `t_c`. 2. The `cos`/`sin` gathers are **separately memoizable** — they depend only on `(t_c, ratio, device)`, exactly like the already-memoized - `compressed_row_positions` (`deepseek4.rs:3013-3037`) and `GRAPH_BLOCK_IDS` - (`dsv4_attention.rs`). Those two are the established in-tree pattern to copy. + `compress_positions` (`layers.rs:2391-2459`) and `GRAPH_BLOCK_IDS` + (`dsv4_attention.rs`). Those two are the established in-tree pattern to copy — + and copy the KEY as well as the shape: one slot per `(ratio, device)`, so the + CSA/HCA interleave cannot thrash it. Item 2 is independent of item 1 and much cheaper to land; it is worth doing first.