perf(qtip): arch-specialise the trellis grouped GEMM (D16) + close the silent-failure holes (D18)

arc-tools/cuda_compile_check.sh

@@ -75,8 +75,8 @@ fi
 step "2/5 compile Arc CUDA kernels (mistralrs-quant + arc-cuda-graph) [features: cuda]"
 # `--lib`: build.rs runs nvcc on every .cu, but no binary is linked. Keeps this
 # step GPU-less-safe (arc-cuda-graph's gemv_bench bin would otherwise link -lcuda).
-cargo build -p mistralrs-quant --lib --features cuda 2>&1 | tail -5 || fail "mistralrs-quant cuda kernels did not compile"
-cargo build -p arc-cuda-graph --lib --features cuda 2>&1 | tail -5 || fail "arc-cuda-graph cuda kernels did not compile"
+cargo build -p mistralrs-quant --lib --features cuda 2>&1 | tail -40 || fail "mistralrs-quant --features cuda FAILED. This step does TWO things: nvcc compiles the .cu files, then rustc type-checks the cuda-gated Rust. A 'Compiling N of N kernels' line above means the KERNELS were fine and the RUST did not compile -- read the rustc errors, not the nvcc output."
+cargo build -p arc-cuda-graph --lib --features cuda 2>&1 | tail -40 || fail "arc-cuda-graph --features cuda FAILED (nvcc kernels AND/OR cuda-gated Rust -- see above)"
 ok "Arc CUDA kernels compiled for sm_${CUDA_COMPUTE_CAP}"
 
 step "3/5 compile-check CUDA-gated tests (--no-run)"

arc-tools/wave64_keystone_gate.sh

@@ -0,0 +1,168 @@
+#!/usr/bin/env bash
+#
+# Arc — wave64 keystone gate: trellis grouped-GEMM arch specialisation.
+#
+# What this answers, in order, so a cheap box can answer the first half and an
+# H200 the second:
+#
+#   STAGE 1  COMPILE (no GPU needed)  — the new SM90+ kernel and the
+#            status-returning launchers compile for sm_80, sm_90a and sm_100a.
+#            nvcc cross-compiles, so ANY box with CUDA >= 12.8 can run this;
+#            an A30/A6000 at ~$0.40/hr is enough. This is the gate that has
+#            never been run for this change.
+#   STAGE 2  PARITY (needs sm_90+)    — the grouped-GEMM CUDA parity tests.
+#            On an sm_90 box these exercise the NEW kernel; on sm_80 they
+#            exercise the old one. Both must pass.
+#   STAGE 3  A/B TIMING (needs sm_90+) — qtip_grouped_curve on this branch and
+#            on the baseline ref, same box, same fixture.
+#
+# D18: three distinct exit codes. 0 = pass. 1 = a genuine failure, the only
+# result that should change a decision. 2 = the environment could not answer
+# (no CUDA, no GPU, wrong arch) — NEVER conflate this with 1. Every terminal
+# outcome, including a timeout, writes a line to the status file, so silence
+# can never be read as success.
+#
+# Usage (main runs this; agents do not call runcrate — D15):
+#   GATE=/tmp/arc-gate-wave64-keystone
+#   mkdir -p $GATE && setsid nohup bash arc-tools/wave64_keystone_gate.sh \
+#       < /dev/null > $GATE/run.log 2>&1 &
+#
+# Every artefact lives under GATE_DIR, which is namespaced per D19. Shared,
+# un-namespaced log paths have already caused one agent's output to be read as
+# another's, and two gates running concurrently on the same box would otherwise
+# overwrite each other's status file and curve logs.
+#
+# Env:
+#   BASELINE_REF   git ref for the A/B baseline (default: origin/master)
+#   GATE_DIR       artefact dir (default: /tmp/arc-gate-wave64-keystone)
+#   STATUS         status file (default: $GATE_DIR/status.txt)
+#   SKIP_AB        set to 1 to run stages 1-2 only
+#   CURVE_ARGS     extra args for qtip_grouped_curve
+
+set -uo pipefail
+
+GATE_DIR=${GATE_DIR:-/tmp/arc-gate-wave64-keystone}
+mkdir -p "$GATE_DIR"
+STATUS=${STATUS:-$GATE_DIR/status.txt}
+BASELINE_REF=${BASELINE_REF:-origin/master}
+REPO=$(git rev-parse --show-toplevel 2>/dev/null) || {
+  echo "STAGE0 env  FAIL not a git repo" | tee -a "$STATUS"; exit 2; }
+cd "$REPO" || exit 2
+
+: > "$STATUS"
+say() { echo "$*" | tee -a "$STATUS"; }
+# Any unexpected exit still writes a terminal line.
+trap 'rc=$?; [ "$rc" -ne 0 ] && say "TERMINAL unexpected-exit rc=$rc"; exit $rc' EXIT
+
+say "wave64 keystone gate  branch=$(git rev-parse --abbrev-ref HEAD)  head=$(git rev-parse --short HEAD)"
+
+# ---------------------------------------------------------------------------
+# Environment probe. Distinguishes "cannot answer" from "failed".
+# ---------------------------------------------------------------------------
+if ! command -v nvcc >/dev/null 2>&1; then
+  export PATH=${CUDA_HOME:-/usr/local/cuda}/bin:$PATH
+fi
+if ! command -v nvcc >/dev/null 2>&1; then
+  say "STAGE1 compile  ENV-CANNOT-ANSWER no nvcc on PATH"
+  say "TERMINAL exit=2"
+  exit 2
+fi
+NVCC_VER=$(nvcc --version | sed -n 's/.*release \([0-9.]*\).*/\1/p' | head -1)
+GPU_CC=$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader 2>/dev/null \
+         | head -1 | tr -d '.' || true)
+say "env nvcc=$NVCC_VER gpu_cc=${GPU_CC:-none}"
+
+# ---------------------------------------------------------------------------
+# STAGE 1 — compile for every arch this change claims to support.
+# sm_100a needs CUDA >= 12.8; when the toolkit is older that is an
+# environment answer, not a failure of the code.
+# ---------------------------------------------------------------------------
+STAGE1_FAIL=0
+for CC in 80 90 100; do
+  if [ "$CC" -ge 100 ]; then
+    case "$NVCC_VER" in
+      12.[0-7]*|12.[0-7]|11.*) say "STAGE1 sm_${CC}  ENV-CANNOT-ANSWER nvcc $NVCC_VER predates sm_100a"; continue ;;
+    esac
+  fi
+  if CUDA_COMPUTE_CAP=$CC FEATURES="cuda" RUN_GPU_TESTS=no \
+     timeout 3600 bash arc-tools/cuda_compile_check.sh > "$GATE_DIR/compile_${CC}.log" 2>&1; then
+    say "STAGE1 sm_${CC}  PASS"
+  else
+    rc=$?
+    [ "$rc" -eq 124 ] && say "STAGE1 sm_${CC}  FAIL timeout" || say "STAGE1 sm_${CC}  FAIL rc=$rc (see $GATE_DIR/compile_${CC}.log)"
+    STAGE1_FAIL=1
+  fi
+done
+[ "$STAGE1_FAIL" -eq 1 ] && { say "TERMINAL exit=1 compile gate failed"; exit 1; }
+
+# ---------------------------------------------------------------------------
+# STAGE 2 — grouped-GEMM parity on real silicon.
+# ---------------------------------------------------------------------------
+if [ -z "${GPU_CC:-}" ]; then
+  say "STAGE2 parity  ENV-CANNOT-ANSWER no GPU visible"
+  say "STAGE3 ab      ENV-CANNOT-ANSWER no GPU visible"
+  say "TERMINAL exit=0 compile-only (stages 2-3 unanswered)"
+  exit 0
+fi
+
+export CUDA_COMPUTE_CAP=$GPU_CC
+PARITY_LOG=$GATE_DIR/parity.log
+if timeout 5400 cargo test --release -p mistralrs-quant --features cuda \
+      cuda_grouped_gemm -- --nocapture --test-threads=1 > "$PARITY_LOG" 2>&1; then
+  # D18 rule 1: a green result must PROVE work happened. `cargo test` reports
+  # success for zero tests run, which is exactly the shape being guarded here.
+  NRUN=$(grep -Eo '[0-9]+ passed' "$PARITY_LOG" | head -1 | grep -Eo '[0-9]+' || echo 0)
+  if [ "${NRUN:-0}" -lt 2 ]; then
+    say "STAGE2 parity  FAIL only ${NRUN} grouped tests ran (expected >= 2); a filter that matches nothing reports success"
+    say "TERMINAL exit=1"; exit 1
+  fi
+  say "STAGE2 parity  PASS sm_${GPU_CC} tests_passed=${NRUN}"
+else
+  rc=$?
+  [ "$rc" -eq 124 ] && say "STAGE2 parity  FAIL timeout" || say "STAGE2 parity  FAIL rc=$rc (see $PARITY_LOG)"
+  say "TERMINAL exit=1"; exit 1
+fi
+
+if [ "${SKIP_AB:-0}" = "1" ]; then
+  say "STAGE3 ab      SKIPPED by request"
+  say "TERMINAL exit=0"; exit 0
+fi
+if [ "$GPU_CC" -lt 90 ] 2>/dev/null; then
+  say "STAGE3 ab      ENV-CANNOT-ANSWER sm_${GPU_CC} takes the Ampere path; the change under test is the SM90+ kernel"
+  say "TERMINAL exit=0 (parity green, delta unmeasured)"; exit 0
+fi
+
+# ---------------------------------------------------------------------------
+# STAGE 3 — A/B against the baseline on the SAME box and fixture.
+# ---------------------------------------------------------------------------
+run_curve() {  # $1 = label, $2 = dir
+  local out="$GATE_DIR/curve_$1.log"
+  if ! ( cd "$2" && timeout 5400 cargo run --release -p mistralrs-quant \
+           --features cuda --example qtip_grouped_curve -- ${CURVE_ARGS:-} ) > "$out" 2>&1; then
+    local rc=$?
+    [ "$rc" -eq 124 ] && say "STAGE3 $1  FAIL timeout" || say "STAGE3 $1  FAIL rc=$rc (see $out)"
+    return 1
+  fi
+  # The harness has previously printed a headline while every row was empty.
+  if ! grep -qE '^[[:space:]]*[0-9]+[[:space:]]*\|' "$out"; then
+    say "STAGE3 $1  FAIL ran clean but produced no data rows (see $out)"
+    return 1
+  fi
+  say "STAGE3 $1  PASS rows=$(grep -cE '^[[:space:]]*[0-9]+[[:space:]]*\|' "$out")  (see $out)"
+  return 0
+}
+
+run_curve head "$REPO" || { say "TERMINAL exit=1"; exit 1; }
+
+BASE_WT=$GATE_DIR/baseline-worktree
+rm -rf "$BASE_WT"
+if ! git worktree add --detach "$BASE_WT" "$BASELINE_REF" >/dev/null 2>&1; then
+  say "STAGE3 baseline  ENV-CANNOT-ANSWER could not create a worktree at $BASELINE_REF"
+  say "TERMINAL exit=0 (head measured, baseline unmeasured)"; exit 0
+fi
+run_curve baseline "$BASE_WT" || { say "TERMINAL exit=1"; exit 1; }
+git worktree remove --force "$BASE_WT" >/dev/null 2>&1 || true
+
+say "STAGE3 ab      PASS — compare $GATE_DIR/curve_head.log vs $GATE_DIR/curve_baseline.log"
+say "TERMINAL exit=0"
+exit 0

mistralrs-quant/examples/qtip_grouped_curve.rs

@@ -76,7 +76,7 @@
 use candle_core::{DType, Device, Result, Tensor};
 use mistralrs_quant::{
     qtip_expected_distinct_experts, Qtip2bLayer, QtipMode, QuantMethod, QTIP_GATHER_GEMV_MAX_PAIRS,
-    QTIP_GROUPED_TILE_K, QTIP_GROUPED_TILE_M, QTIP_ONDEVICE_MOE_MAX_TOKENS_ENV,
+    QTIP_GROUPED_TILE_K, QTIP_ONDEVICE_MOE_MAX_TOKENS_ENV,
 };
 use std::sync::Arc;
 use std::time::Instant;
@@ -313,11 +313,15 @@ fn assert_config_is_measurable(e: usize, batches: &[usize]) -> Result<()> {
     // expert read once per step" — the model behind the floor column — only
     // holds while the average pairs per woken expert stays within one tile.
     let per_expert = max_pairs as f64 / qtip_expected_distinct_experts(max_b, TOPK, e);
-    if per_expert > QTIP_GROUPED_TILE_M as f64 {
+    // D16: the m-tile is arch-dependent (16 on Ampere, 64 on SM90+). Ask the
+    // device -- assuming the Ampere tile on an H200 would bail four times too
+    // early and mis-state the traffic model below.
+    let tile_m = tile_m_for_device()?;
+    if per_expert > tile_m as f64 {
         candle_core::bail!(
-            "at B={max_b} each woken expert takes {per_expert:.1} pairs, past GROUPED_TILE_M=\
-             {QTIP_GROUPED_TILE_M}: the grouped kernel re-stages weights per m-tile beyond that, \
-             so the unique-expert floor column would understate its real traffic."
+            "at B={max_b} each woken expert takes {per_expert:.1} pairs, past this device's \
+             grouped m-tile {tile_m}: the grouped kernel re-stages weights per m-tile beyond \
+             that, so the unique-expert floor column would understate its real traffic."
         );
     }
     Ok(())
@@ -437,14 +441,34 @@ struct Shape {
     calls_per_layer: f64,
 }
 
+/// The grouped kernel's m-tile ON THIS DEVICE (D16: 16 on Ampere, 64 on
+/// SM90+). Every traffic model in this harness is bounded by it, so it is
+/// queried, never assumed -- a wrong tile silently mis-states the grouped
+/// path's advantage.
+fn tile_m_for_device() -> Result<usize> {
+    #[cfg(feature = "cuda")]
+    {
+        mistralrs_quant::qtip_grouped_tile_m()
+    }
+    #[cfg(not(feature = "cuda"))]
+    {
+        Ok(mistralrs_quant::QTIP_GROUPED_TILE_M)
+    }
+}
+
 /// Packed bytes per expert: 2-bit codes (n*k/4) + F32 row scales.
 fn bytes_per_expert(s: &Shape) -> f64 {
     s.n as f64 * s.k as f64 * 0.25 + s.n as f64 * 4.0
 }
 
 /// Build the routing draws once per batch size and share them across paths, so
 /// grouped and GEMV are compared on byte-identical traffic.
-fn build_routing(dev: &Device, b: usize, e: usize) -> Result<(Vec<Tensor>, f64, f64, f64)> {
+fn build_routing(
+    dev: &Device,
+    b: usize,
+    e: usize,
+    tile_m: usize,
+) -> Result<(Vec<Tensor>, f64, f64, f64)> {
     let mut rng = Rng(ROUTING_SEED ^ (b as u64).wrapping_mul(0x9E37_79B9));
     let mut sets = Vec::with_capacity(ROUTING_WINDOWS);
     let (mut sum_distinct, mut sum_tiles) = (0.0f64, 0.0f64);
@@ -454,7 +478,7 @@ fn build_routing(dev: &Device, b: usize, e: usize) -> Result<(Vec<Tensor>, f64,
         sum_distinct += counts.iter().filter(|c| **c > 0).count() as f64;
         sum_tiles += counts
             .iter()
-            .map(|c| c.div_ceil(QTIP_GROUPED_TILE_M) as f64)
+            .map(|c| c.div_ceil(tile_m) as f64)
             .sum::<f64>();
         sets.push(Tensor::from_vec(idx, (b, TOPK), dev)?);
     }
@@ -479,7 +503,7 @@ fn time_curve(
     path.select();
     for &b in batches {
         assert_cap_unchanged(cap_at_startup)?;
-        let (idx_sets, mean_distinct, mean_tiles, pairs) = build_routing(dev, b, e)?;
+        let (idx_sets, mean_distinct, mean_tiles, pairs) = build_routing(dev, b, e, tile_m_for_device()?)?;
         assert_routing_matches_model(b, e, mean_distinct)?;
         let issued_units = if path == Path::Gemv {
             pairs

mistralrs-quant/kernels/qtip/qtip2b_common.cuh

@@ -63,6 +63,30 @@ __device__ __forceinline__ uint32_t q2b_state_from_window(uint32_t win16) {
     return ((r & 0x5555u) << 1) | ((r >> 1) & 0x5555u);
 }
 
+// ---------------------------------------------------------------------------
+// SEQUENTIAL RUN ADVANCE — the cheap half of the same identity.
+//
+// `q2b_state_from_window` costs ~11 ALU ops (two smem words, two shifts, an
+// OR, __brev, and the adjacent-bit swap). That is the price of decoding an
+// ISOLATED position. Along a CONTIGUOUS run of symbols the original trellis
+// recurrence is far cheaper — one SHF plus one LOP3:
+//
+//     state(t+1) = ((state(t) << 2) | sym[t+1]) & 0xFFFF
+//
+// so a run of R symbols costs `window + (R-1)*advance` instead of
+// `R * window`. At R = 32 that is ~11.3 ops/weight -> ~2 ops/weight of state
+// reconstruction. This matters because the trellis GEMM is ALU-issue-bound on
+// decode, not bandwidth-bound (98/98 gen-2 GEMV variants measured: latency
+// hiding does not pay, per-symbol decode is the ceiling).
+//
+// The window identity is still what SEEDS each run, so runs stay independent:
+// no cross-thread state chain, no warm-up replay. Mirrored and property-tested
+// on the Rust side by `grouped::run_states_2b`.
+// ---------------------------------------------------------------------------
+__device__ __forceinline__ uint32_t q2b_advance_state(uint32_t state, uint32_t sym) {
+    return ((state << Q2B_K) | (sym & (Q2B_ALPHABET - 1u))) & Q2B_STATE_MASK;
+}
+
 // ---------------------------------------------------------------------------
 // cp.async helpers (sm_80+). `valid == false` uses the src-size-0 form, which
 // zero-fills the destination without touching the source address — that is