ArcGraph: land the CUDA-graph capture lane — 3,961 -> 3.8 launch APIs/token, 10,114 -> 131 allocs/step

.claude/run_capture.sh

@@ -0,0 +1,74 @@
+#!/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 <legname> <gen_len> <env-string...>
+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"
+# 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, 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 ==="
+  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"

arc-cuda-graph/src/graph.rs

@@ -367,11 +367,32 @@ 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()
             .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)
     }
 
@@ -542,22 +563,68 @@ 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");
+
+        // 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
-        // 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 +638,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 +654,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
@@ -667,6 +736,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());
     }
 }

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.
 

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)
     }

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)?