From 0595ba07829463334bf838d2e99e71576646456a Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:15:49 +0100 Subject: [PATCH 1/8] =?UTF-8?q?fix(ArcGate):=20the=20citation=20gate=20had?= =?UTF-8?q?=20never=20opened=20a=20.rs=20file=20=E2=80=94=20310=20citation?= =?UTF-8?q?s=20ungated?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `doc_citations.rs` says, in its own header, that it verifies the file-and-line evidence a reader is expected to follow. It scanned `SCAN_ROOTS = ["memory", "arc-tools"]` and filtered to `.md`. So for its whole life it checked the mission record and never once opened a source file — while ~310 citations sat in Rust doc comments, pointing into the fastest-moving code in the repo, completely ungated. Its own doc named the fix: "Adding a root here is a one-line change plus whatever waivers it brings." This is that change, plus the two things it turned out to need. * `SCAN_ROOTS` gains the first-party crates; `SCANNED_DOC_EXTS` gains `.rs`. Census 972 -> 1,244 citations. `CITATION_FLOOR` 700 -> 1,100, above the `.md`-only total on purpose: a scope guard has to notice scope *shrinking*, which is exactly how the Rust corpus stayed invisible. * **A `.rs` line cannot be scanned whole.** Citation-shaped text lives in string literals — this file's own BASELINE is a table of them, and its fixtures build citations that are deliberately wrong. Scanning those manufactures `SymbolRot`, which has no waiver path, so the gate would have been unlandable against its own source. `split_rust` lexes comments out: nested `/* */`, `\`-continued and multi-line raw strings, `'x'` char literals kept distinct from `'a` lifetimes. A per-line scanner is not enough and this is not theoretical — the first fixture written here was a multi-line raw string, and a per-line scanner reported every deliberately broken citation in it as real. * **Widening silently disarmed the one unwaivable verdict.** Once `.rs` files were documents, `symbol_exists_anywhere` started finding a cited symbol in the very comment that cited it. Every `SymbolRot` degraded to the waivable `SymbolMoved`, for the entire new corpus, and nothing would have said so. The symbol search now reads comment-free code — the same rule `.md` already gets from `CODE_EXTS`: a symbol surviving only in prose is the rot. Two new tests, because a capability nobody proves is a capability nobody has: `rust_comments_are_scanned_and_code_is_not` pins the exact citations extracted from a fixture holding an escaped quote, a raw string containing `//`, and a lifetime; `rust_comment_citations_are_gated` proves a dead symbol cited from a Rust comment still hard-fails. Result: 0 live findings. The 130 that surfaced are all UNRESOLVED-PATH into upstream trees Arc reads but does not vendor — sglang, vLLM, TensorRT-LLM, DeepSeek's `inference/`, the QTIP reference, candle, cudarc — pinned as 79 BASELINE rows (one per doc/path, generated from the run, not hand-guessed). **No SymbolRot accompanied the widening**: all ~310 Rust-comment citations name symbols that still exist. Two self-inflicted findings were fixed rather than waived, since this file's prose about citations is now itself scanned. cargo test -p mistralrs-core --test doc_citations: 12 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-core/tests/doc_citations.rs | 915 +++++++++++++++++++++++++- 1 file changed, 902 insertions(+), 13 deletions(-) diff --git a/mistralrs-core/tests/doc_citations.rs b/mistralrs-core/tests/doc_citations.rs index 3a0272b4e..8b585be82 100644 --- a/mistralrs-core/tests/doc_citations.rs +++ b/mistralrs-core/tests/doc_citations.rs @@ -8,7 +8,7 @@ //! nothing goes red — the compiler has no opinion about prose. //! //! Prose cannot be checked mechanically. But the load-bearing part of Arc's -//! docs is not prose, it is **evidence**: `path/to/file.rs:123` citations that +//! docs is not prose, it is **evidence**: file-and-line citations that //! a reader is expected to follow. Those *are* mechanically checkable, and when //! they rot the surrounding claim is unfalsifiable — the reader lands on //! unrelated code and cannot tell whether the doc was ever right. @@ -29,6 +29,29 @@ //! (`deepseek4.rs:2211`) ``, by far the dominant shape in these docs — the //! identifier must appear near the cited line. //! +//! ## Where it looks: `.md` **and Rust comments** +//! +//! For its first life this gate read `.md` files only. That was a hole big +//! enough to drive the whole point through: Arc's densest citation corpus is +//! not the mission record, it is the **module and item docs in the crates +//! themselves** — several hundred `file.rs:N` pointers written by the same +//! agents, into the same moving code, and *none* of them were gated. A doc +//! comment that says "the mask is dropped at `sinks.rs:214`" rots exactly like +//! a `.md` one, and rots faster, because the file it cites is the file it +//! lives next to. +//! +//! So [`SCAN_ROOTS`] now names the first-party crate directories too, and +//! [`analyze`] reads `.rs` alongside `.md`. +//! +//! **In a `.rs` file only comment text is scanned** — see [`rust_comment`]. +//! Code is not prose: a `.rs` line may hold a path-and-number inside a string +//! literal (this very file's [`BASELINE`] is a table of them, and its fixtures +//! build citations that are *deliberately* wrong). Scanning those would +//! manufacture [`Kind::SymbolRot`] findings, which have no waiver path, and +//! the gate would be unlandable against its own source. The extractor is +//! deliberately conservative: when it cannot tell whether a `//` is a comment +//! or string content, it treats the line as code and skips it. +//! //! ## Drift versus rot //! //! These are not the same defect and are not treated the same way. @@ -86,7 +109,8 @@ use std::path::{Path, PathBuf}; // Configuration // --------------------------------------------------------------------------- -/// Directories scanned for `.md` files carrying citations. +/// Directories scanned for files carrying citations — `.md` anywhere, and +/// `.rs` **comments** in the first-party crates. /// /// `docs/` and `research/` are deliberately **excluded for now**: 212 of their /// 336 citations point into external reference checkouts (the v4 reference @@ -94,7 +118,32 @@ use std::path::{Path, PathBuf}; /// waiver rows and almost no verification. `memory/` is the mission record and /// `arc-tools/` the runbooks — that is where a false citation does damage. /// Adding a root here is a one-line change plus whatever waivers it brings. -const SCAN_ROOTS: &[&str] = &["memory", "arc-tools"]; +/// +/// The crate roots below carry the ~310 citations that live in Rust comments. +/// They are listed individually rather than as a bare `"."` for two reasons: +/// `.` would sweep in `docs/` and `research/` by the back door, and it would +/// index the vendored/generated trees that [`SKIP_DIRS`] only partly covers. +/// Crates with no citations *today* are still named, so the first one written +/// tomorrow is gated on arrival rather than on someone remembering this list. +const SCAN_ROOTS: &[&str] = &[ + "memory", + "arc-tools", + "mistralrs-core", + "mistralrs-quant", + "arc-engine", + "arc-cli", + "arc-cuda-graph", + "arc-bench", + "arc-profiler", + "arc-turbo", +]; + +/// Extensions read as citation-bearing documents by [`analyze`]. +/// +/// Distinct from [`CITED_EXTS`] (what a citation may *point at*) and from +/// [`CODE_EXTS`] (where a symbol may *live*). A `.rs` entry here means "read +/// this file's comments", never its code — see [`rust_comment`]. +const SCANNED_DOC_EXTS: &[&str] = &["md", "rs"]; /// Extensions a citation may point at. Anything else is not a citation. const CITED_EXTS: &[&str] = &[ @@ -135,8 +184,15 @@ const SKIP_DIRS: &[&str] = &[".git", "target", "node_modules", ".venv", "__pycac const DRIFT_WINDOW: i64 = 15; /// Floor on citations found, so an empty or mis-rooted scan fails loudly -/// instead of reading as a pass. Measured at 798 when this landed. -const CITATION_FLOOR: usize = 700; +/// instead of reading as a pass. +/// +/// Measured at 798 when this landed, 1,244 once `.rs` comments were added — +/// the crates carry roughly a quarter of Arc's citation corpus. The floor is +/// set above the `.md`-only total on purpose: reverting [`SCANNED_DOC_EXTS`] +/// to `["md"]`, or dropping the crate [`SCAN_ROOTS`], drops the census back to +/// ~970 and this fails. A scope guard has to notice scope *shrinking*, which +/// is the exact way the Rust corpus went ungated for as long as it did. +const CITATION_FLOOR: usize = 1_100; // --------------------------------------------------------------------------- // Findings @@ -214,9 +270,11 @@ impl Finding { /// `cite` matching is deliberately two-mode, so the waiver is no broader than /// the debt it records: /// -/// * contains `:` — matches that **exact citation** (`device.rs:754-908`). +/// * contains `:` — matches that **exact citation**, line numbers and all. /// Used where the line number is the thing that is wrong, so a second bad -/// citation into the same file is not silently covered. +/// citation into the same file is not silently covered. (No example is +/// spelled out here: since this gate reads `.rs` comments, a citation +/// written in prose *about* the matcher would be scanned as a real one.) /// * no `:` — matches **any citation of that path** in that doc. Used for /// external references, where every line of `server_args.py` is equally /// unresolvable and pinning each would be noise. @@ -808,6 +866,503 @@ static BASELINE: &[Waiver] = &[ kind: Kind::Unresolved, why: "candle upstream, a registry dependency, not vendored here", }, + // -- Rust-comment citations, added when SCAN_ROOTS grew to cover the + // crates (see the module docs). Every row below is a pointer into an + // upstream reference checkout — sglang, vLLM, TensorRT-LLM, DeepSeek's + // own `inference/`, the QTIP reference, candle, cudarc, toktrie — that + // Arc reads beside the code but does not vendor, so none of them can + // ever resolve here. **No `SymbolRot` accompanied this widening**: the + // ~310 citations in Rust comments all name symbols that still exist. + Waiver { + doc: "arc-engine/src/weight_schema.rs", + cite: "deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/attention/backends/flash.rs", + cite: "candle-flash-attn-v3/src/lib.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/attention/backends/flash.rs", + cite: "hkernel/flash_fwd_launch_template.h", + kind: Kind::Unresolved, + why: "flash-attention is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-core/src/kv_cache/single_cache.rs", + cite: "candle-core/src/tensor_cat.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/kv_sharing/evict.rs", + cite: "evict_policy.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/kv_sharing/evict.rs", + cite: "python/sglang/srt/mem_cache/evict_policy.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/kv_sharing/radix.rs", + cite: "radix_cache.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/layers.rs", + cite: "candle-core/src/tensor_cat.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/layers.rs", + cite: "deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/layers.rs", + cite: "inference/model.py", + kind: Kind::Unresolved, + why: "DeepSeek's own inference/ reference is an upstream reference \ + checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/layers.rs", + cite: "srt/models/deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "candle-core/src/tensor_cat.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "deepseek_v4_memory_pool.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "deepseek_v4_nextn.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "deepseek_v4_rope.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "eagle_utils.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "eagle_worker.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "inference/model.py", + kind: Kind::Unresolved, + why: "DeepSeek's own inference/ reference is an upstream reference \ + checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "mem_cache/deepseek_v4_memory_pool.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "model.py", + kind: Kind::Unresolved, + why: "DeepSeek's own inference/ reference is an upstream reference \ + checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "router/gate_linear.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "srt/layers/deepseek_v4_rope.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "srt/models/deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "srt/models/deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/deepseek4.rs", + cite: "topk.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_attention.rs", + cite: "model_executor/pool_configurator.py", + kind: Kind::Unresolved, + why: "TensorRT-LLM is an upstream reference checkout, not vendored \ + here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_indexer.rs", + cite: ".../sglang/srt/models/deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_indexer.rs", + cite: "attention/dsv4/compressor.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_indexer.rs", + cite: "compressor.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_indexer.rs", + cite: "deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_indexer.rs", + cite: "indexer.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_kv_fp8.rs", + cite: "inference/model.py", + kind: Kind::Unresolved, + why: "DeepSeek's own inference/ reference is an upstream reference \ + checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_mhc.rs", + cite: "deepseek_v4_nextn.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_mhc.rs", + cite: "sglang/srt/layers/mhc.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/models/dsv4_mhc.rs", + cite: "sglang/srt/models/deepseek_v4_nextn.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/expert_parallel.rs", + cite: "sglang/python/sglang/srt/eplb/eplb_algorithms/deepseek.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/expert_parallel.rs", + cite: "tensorrt_llm/examples/wide_ep/ep_load_balancer/README.md", + kind: Kind::Unresolved, + why: "TensorRT-LLM is an upstream reference checkout, not vendored \ + here", + }, + Waiver { + doc: "mistralrs-core/src/moe/experts.rs", + cite: "inference/model.py", + kind: Kind::Unresolved, + why: "DeepSeek's own inference/ reference is an upstream reference \ + checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/experts.rs", + cite: "sglang/jit_kernel/csrc/deepseek_v4/silu_and_mul_masked_post_quant.cuh", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/experts.rs", + cite: "silu_and_mul_masked_post_quant.cuh", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/experts.rs", + cite: "sm100_fp8_fp4_mega_moe.cuh", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/moe/experts.rs", + cite: "srt/models/deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/llg.rs", + cite: "toktrie-1.7.0/src/toktree.rs", + kind: Kind::Unresolved, + why: "toktrie is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/llg.rs", + cite: "toktrie_hf_tokenizers-1.7.0/src/lib.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/loaders/normal_loaders.rs", + cite: "deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mod.rs", + cite: "cuda_backend/mod.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mod.rs", + cite: "tensor.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "config/speculative.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "deepseek_v4_nextn.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "eagle_info.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "eagle_utils.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "eagle_worker.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "eagle_worker_common.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "logits_processor.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "pyexecutor/resource_manager.py", + kind: Kind::Unresolved, + why: "TensorRT-LLM is an upstream reference checkout, not vendored \ + here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "scheduler_components/batch_result_processor.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "server_args.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "spec_utils.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "speculative/eagle_worker_common.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "v1/core/sched/scheduler.py", + kind: Kind::Unresolved, + why: "vLLM is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/mtp_pipeline.rs", + cite: "v1/spec_decode/metrics.py", + kind: Kind::Unresolved, + why: "vLLM is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/pipeline/normal.rs", + cite: "candle-core/src/cuda_backend/device.rs", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/prefix_cacher.rs", + cite: "block_pool.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/prefix_cacher.rs", + cite: "radix_cache.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/src/sampler.rs", + cite: "candle-kernels/src/reduce.cu", + kind: Kind::Unresolved, + why: "candle is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/tests/synthetic_load_smoke.rs", + cite: "deepseek_v4_nextn.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-core/tests/synthetic_load_smoke.rs", + cite: "eagle_worker.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/cuda_peer.rs", + cite: "src/driver/result.rs", + kind: Kind::Unresolved, + why: "cudarc is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/cuda_peer.rs", + cite: "src/driver/safe/core.rs", + kind: Kind::Unresolved, + why: "cudarc is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/cuda_peer.rs", + cite: "src/driver/sys/mod.rs", + kind: Kind::Unresolved, + why: "cudarc is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "lib/algo/ldlq.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "lib/codebook/bitshift.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "lib/utils/math_utils.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "math_utils.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "quantize_llama/input_hessian_llama.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "research/code/01_weight_compression/qtip/lib/utils/data_utils.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, + Waiver { + doc: "mistralrs-quant/src/qtip/viterbi.rs", + cite: "research/code/01_weight_compression/qtip/lib/utils/math_utils.py", + kind: Kind::Unresolved, + why: "the QTIP reference is an upstream reference checkout, not \ + vendored here", + }, ]; // --------------------------------------------------------------------------- @@ -822,6 +1377,7 @@ struct Repo { files: Vec, by_suffix: BTreeMap>, bodies: HashMap>>, + code_bodies: HashMap>>, } impl Repo { @@ -844,6 +1400,7 @@ impl Repo { files, by_suffix, bodies: HashMap::new(), + code_bodies: HashMap::new(), } } @@ -890,10 +1447,37 @@ impl Repo { } } - /// Does `sym` appear as a whole word in any code file? + /// The file's lines with Rust comments blanked out, so a symbol that + /// survives only in prose does not read as a symbol that still exists. + /// + /// Non-`.rs` files are returned unchanged: `.md` never reaches here (it is + /// absent from [`CODE_EXTS`]) and the rest — `.py`, `.cu`, `.toml` — have + /// their own comment syntaxes that this gate has no reason to model. + fn code_lines(&mut self, idx: usize) -> Option<&Vec> { + if !self.code_bodies.contains_key(&idx) { + let path = self.files[idx].clone(); + let stripped = self.lines(idx).map(|body| { + if path.ends_with(".rs") { + split_rust(&body.join("\n")).code + } else { + body.clone() + } + }); + self.code_bodies.insert(idx, stripped); + } + self.code_bodies.get(&idx).and_then(|o| o.as_ref()) + } + + /// Does `sym` appear as a whole word in the **code** of any code file? /// /// Only reached when a cited symbol is missing from its own cited file, so /// the whole-repo read stays rare. + /// + /// Comments are excluded — see [`RustSplit`]. Without that, a citation + /// written in a Rust doc comment is its own witness: the symbol it names + /// is present in the comment naming it, every [`Kind::SymbolRot`] softens + /// to a waivable [`Kind::SymbolMoved`], and the one verdict with no waiver + /// path stops being reachable for the entire crate corpus. fn symbol_exists_anywhere(&mut self, sym: &str) -> bool { let code: Vec = (0..self.files.len()) .filter(|&i| has_ext(&self.files[i], CODE_EXTS)) @@ -904,7 +1488,11 @@ impl Repo { }) .collect(); for i in code { - if !self.symbol_lines(i, sym).is_empty() { + let hit = self + .code_lines(i) + .map(|body| body.iter().any(|l| contains_word(l, sym))) + .unwrap_or(false); + if hit { return true; } } @@ -1049,6 +1637,223 @@ fn identifier_before(chars: &[char], cite_start: usize) -> Option { Some(last) } +/// A Rust file cut into the two halves this gate treats differently. +/// +/// Both halves matter and they matter in *opposite* directions, which is why +/// one pass produces both: +/// +/// * [`RustSplit::comments`] is where citations are allowed to live. Scanning +/// code instead would manufacture [`Kind::SymbolRot`] out of string +/// literals — this file's own [`BASELINE`] is a table of citation-shaped +/// strings, and rot has no waiver path, so the gate would be unlandable +/// against its own source. +/// * [`RustSplit::code`] is where a symbol is allowed to *exist*. Once `.rs` +/// files became documents, `symbol_exists_anywhere` started finding symbols +/// in the very comment that cited them — every rot downgraded to the +/// waivable [`Kind::SymbolMoved`] and the one unwaivable verdict quietly +/// stopped being reachable for the whole new corpus. Blanking comments +/// restores it, and it is the same rule `.md` already gets from +/// [`CODE_EXTS`]: a symbol surviving only in prose is exactly the rot this +/// is looking for. +struct RustSplit { + /// `(1-based line, comment text)`, one entry per comment span. + comments: Vec<(usize, String)>, + /// The body line for line, with every comment span blanked to spaces. + code: Vec, +} + +/// Split Rust source into comment spans and comment-free code. +/// +/// A line-at-a-time scanner cannot do this: a multi-line raw string +/// (`r#"…"#`), a `\`-continued string literal and a `/* */` block all carry +/// state across newlines, and a `//!` sitting inside any of them is not a +/// comment. Getting that wrong is not cosmetic — the first fixture written for +/// this feature was a multi-line raw string full of deliberately-broken +/// citations, and a per-line scanner reported every one of them as real. +/// +/// Handled: `//` line comments (so `///` and `//!` too), nested `/* */` +/// blocks, `"…"` with `\` escapes, `r"…"` / `r#…#"…"#…` raw strings including +/// the `b`-prefixed forms, and `'x'` char literals kept distinct from `'a` +/// lifetimes. +fn split_rust(body: &str) -> RustSplit { + #[derive(Clone, Copy)] + enum St { + Code, + Line, + Block(usize), + Str, + Raw(usize), + } + + let chars: Vec = body.chars().collect(); + let mut st = St::Code; + let mut comments: Vec<(usize, String)> = Vec::new(); + let mut code: Vec = Vec::new(); + let mut line_no = 1usize; + let mut cur_code = String::new(); + let mut cur_comment = String::new(); + let mut i = 0usize; + + let flush_comment = |buf: &mut String, line: usize, out: &mut Vec<(usize, String)>| { + if !buf.trim().is_empty() { + out.push((line, std::mem::take(buf))); + } else { + buf.clear(); + } + }; + + while i < chars.len() { + let c = chars[i]; + if c == '\n' { + if matches!(st, St::Line) { + st = St::Code; + } + flush_comment(&mut cur_comment, line_no, &mut comments); + code.push(std::mem::take(&mut cur_code)); + line_no += 1; + i += 1; + continue; + } + match st { + St::Line => { + cur_comment.push(c); + cur_code.push(' '); + i += 1; + } + St::Block(depth) => { + if c == '*' && chars.get(i + 1) == Some(&'/') { + st = if depth == 1 { + St::Code + } else { + St::Block(depth - 1) + }; + cur_comment.push_str("*/"); + cur_code.push_str(" "); + i += 2; + } else if c == '/' && chars.get(i + 1) == Some(&'*') { + st = St::Block(depth + 1); + cur_comment.push_str("/*"); + cur_code.push_str(" "); + i += 2; + } else { + cur_comment.push(c); + cur_code.push(' '); + i += 1; + } + } + St::Str => { + cur_code.push(c); + if c == '\\' { + if let Some(&n) = chars.get(i + 1) { + // A `\` at end of line continues the literal; the + // newline arm above must still see the `\n`. + if n != '\n' { + cur_code.push(n); + i += 2; + continue; + } + } + i += 1; + continue; + } + if c == '"' { + st = St::Code; + } + i += 1; + } + St::Raw(hashes) => { + cur_code.push(c); + if c == '"' { + let closed = (1..=hashes).all(|k| chars.get(i + k) == Some(&'#')); + if closed { + for _ in 0..hashes { + cur_code.push('#'); + } + st = St::Code; + i += 1 + hashes; + continue; + } + } + i += 1; + } + St::Code => { + if c == '/' && chars.get(i + 1) == Some(&'/') { + st = St::Line; + cur_comment.push_str("//"); + cur_code.push_str(" "); + i += 2; + } else if c == '/' && chars.get(i + 1) == Some(&'*') { + st = St::Block(1); + cur_comment.push_str("/*"); + cur_code.push_str(" "); + i += 2; + } else if c == '"' { + st = St::Str; + cur_code.push(c); + i += 1; + } else if c == 'r' && raw_open(&chars, i).is_some() { + let (hashes, span) = raw_open(&chars, i).expect("checked"); + st = St::Raw(hashes); + cur_code.extend(chars[i..i + span].iter()); + i += span; + } else if c == '\'' { + // `'x'` / `'\n'` are literals; `'a` in `&'a str` is not. + let escaped = + chars.get(i + 1) == Some(&'\\') && chars.get(i + 3) == Some(&'\''); + let plain = chars.get(i + 1) != Some(&'\\') && chars.get(i + 2) == Some(&'\''); + let span = if escaped { + 4 + } else if plain { + 3 + } else { + 1 + }; + cur_code.extend(chars[i..(i + span).min(chars.len())].iter()); + i += span; + } else { + cur_code.push(c); + i += 1; + } + } + } + } + flush_comment(&mut cur_comment, line_no, &mut comments); + if !cur_code.is_empty() { + code.push(cur_code); + } + RustSplit { comments, code } +} + +/// If a raw-string literal opens at `i` (the `r`), its hash count and the +/// width of the opening token. `br"…"` reaches here with `i` on the `r`. +fn raw_open(chars: &[char], i: usize) -> Option<(usize, usize)> { + let mut k = i + 1; + let mut hashes = 0usize; + while chars.get(k) == Some(&'#') { + hashes += 1; + k += 1; + } + if chars.get(k) == Some(&'"') { + Some((hashes, k + 1 - i)) + } else { + None + } +} + +/// Every `(line, text)` in a document that a citation may live in. +/// +/// `.md` is prose end to end. `.rs` is prose only inside comments. +fn prose_spans(doc: &str, body: &str) -> Vec<(usize, String)> { + if doc.ends_with(".rs") { + split_rust(body).comments + } else { + body.lines() + .enumerate() + .map(|(n, l)| (n + 1, l.to_string())) + .collect() + } +} + /// Extract every citation on one line of a doc. fn citations_in_line(line: &str) -> Vec { let chars: Vec = line.chars().collect(); @@ -1123,7 +1928,7 @@ fn analyze(root: &Path, scan_roots: &[&str]) -> Census { let base = root.join(scan); let mut found = Vec::new(); walk(root, &base, &mut found); - docs.extend(found.into_iter().filter(|p| p.ends_with(".md"))); + docs.extend(found.into_iter().filter(|p| has_ext(p, SCANNED_DOC_EXTS))); } docs.sort(); // So a nested scan root ("memory" plus "memory/mission") counts each doc @@ -1142,9 +1947,8 @@ fn analyze(root: &Path, scan_roots: &[&str]) -> Census { Ok(b) => b, Err(_) => continue, }; - for (n, line) in body.lines().enumerate() { - let doc_line = n + 1; - for c in citations_in_line(line) { + for (doc_line, prose) in prose_spans(doc, &body) { + for c in citations_in_line(&prose) { census.total += 1; let mk = |kind: Kind, detail: String| Finding { doc: doc.clone(), @@ -1439,6 +2243,9 @@ impl Scratch { fn run(&self) -> Census { analyze(&self.0, &["memory"]) } + fn run_roots(&self, roots: &[&str]) -> Census { + analyze(&self.0, roots) + } } impl Drop for Scratch { @@ -1537,6 +2344,88 @@ fn fixtures_prove_each_violation_class_fails() { assert!(!Kind::SymbolRot.is_waivable()); } +/// **Proof that reading `.rs` is not decorative, and that it reads only +/// comments.** +/// +/// Two failures are possible when a citation gate grows a second file type, +/// and both are silent: it can scan `.rs` and find nothing (the corpus stays +/// ungated while the census looks bigger), or it can scan `.rs` code and +/// manufacture unwaivable [`Kind::SymbolRot`] out of string literals. The +/// fixture below contains one of each shape, plus the three lines that +/// [`rust_comment`] has to get right — an escaped quote, a raw string holding +/// `//`, and a lifetime — and asserts exactly which citations come back. +#[test] +fn rust_comments_are_scanned_and_code_is_not() { + let s = Scratch::new("rustcomments"); + s.write("src/foo.rs", &fixture_source()); + s.write( + "src/claims.rs", + r####" +//! Module doc: `alpha_beta` (`src/foo.rs:10`) is the entry point. +/// Item doc citing a range: `alpha_beta` (`src/foo.rs:5-15`). +const CITE: &str = "`ghost_symbol` (src/foo.rs:10)"; // a literal, not a claim +let msg = "he said \"see src/foo.rs:10\""; // escaped quotes stay in the string +let url = r#"https://x/y//src/foo.rs:10"#; // a raw string holding a comment marker +fn f<'a>(x: &'a str) {} // lifetime, then a real one: `alpha_beta` (`src/foo.rs:10`) +"####, + ); + let c = s.run_roots(&["src"]); + + // Every citation the extractor found, in order. + let body = fs::read_to_string(s.0.join("src/claims.rs")).expect("fixture readable"); + let found: Vec = prose_spans("src/claims.rs", &body) + .into_iter() + .flat_map(|(_, p)| citations_in_line(&p).into_iter().map(|c| c.text)) + .collect(); + assert_eq!( + found, + vec![ + "src/foo.rs:10", // module doc + "src/foo.rs:5-15", // item doc + "src/foo.rs:10", // trailing comment on the lifetime line + ], + "comment text must be scanned and code text must not" + ); + + // And end to end: nothing in that file is a violation. If string literals + // leaked in, `ghost_symbol` would be an unwaivable SymbolRot. + assert!( + c.findings.is_empty(), + "a .rs fixture whose comments all check out must be clean:\n{}", + c.findings + .iter() + .map(|f| f.render()) + .collect::>() + .join("\n") + ); + assert_eq!(c.total, 3, "3 citations live in comments in that fixture"); +} + +/// A `.rs` file whose **comment** carries a rotten citation is caught, so the +/// widening is a gate and not just a bigger census. +#[test] +fn rust_comment_citations_are_gated() { + let s = Scratch::new("rustgate"); + s.write("src/foo.rs", &fixture_source()); + s.write( + "src/claims.rs", + "//! `gamma_delta_gone` (`src/foo.rs:10`) does the work.\n\ + //! Also `alpha_beta` (`src/foo.rs:999`).\n", + ); + let c = s.run_roots(&["src"]); + let kinds: BTreeSet = c.findings.iter().map(|f| f.kind).collect(); + assert!( + kinds.contains(&Kind::SymbolRot), + "a dead symbol cited from a Rust comment must be rot: {:?}", + c.findings.iter().map(|f| f.render()).collect::>() + ); + assert!( + kinds.contains(&Kind::OutOfRange), + "a past-the-end line cited from a Rust comment must be out of range: {:?}", + c.findings.iter().map(|f| f.render()).collect::>() + ); +} + /// The other half of "prove it red": the same checker must stay **green** on /// citations that are merely churned. A gate that fires on everything is as /// useless as one that fires on nothing. From a1870391d7dc8b5e2b19675b4202fa9170b25395 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:19:33 +0100 Subject: [PATCH 2/8] =?UTF-8?q?perf(ArcKernels):=20lm=5Fhead=20re-read=20i?= =?UTF-8?q?ts=201.06=20GB=20weight=20256=20times=20=E2=80=94=20flatten=20r?= =?UTF-8?q?ank-3=20in=20UnquantLinear?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit V4's `lm_head` is 67 ms/step at B=256, 8.41% of a 794 ms decode step, in a single launch. Its bound is `max(271 GFLOP / 989 TFLOP/s, 1.06 GB / 4.8 TB/s) = 0.27 ms`. That is **244x over, and 244 ~= B = 256**, which is the entire explanation: back-solve the traffic and it is 256 x 1.06 GB = 271 GB per step at 4.05 TB/s — the kernel is running at 84% of peak HBM bandwidth while doing 256x the necessary reads. It is not a slow kernel. It is a correct kernel being asked for the wrong work. Every activation in this engine is rank-3 `[B, T, hidden]`. `UnquantLinear:: forward` matched on that and took `w.broadcast_left(B)` — a stride-0 view — after which both dispatch arms treat `B` as a GEMM batch dimension with `stride_b = 0`. At decode `T = 1`, so it is B independent `m = 1` GEMVs over one weight: a batched GEMV wearing a GEMM's name. The 1.06 GB weight does not fit in L2, so there is no reuse across rows to save it. Flattening `[B, T, hidden]` to `[B*T, hidden]` makes it one `[256, 4096] x [4096, 129280]` GEMM: weight streamed once, `m = 256` instead of `m = 1`. It also drops the trailing `.t()` the cuBLASLt arm returned, so the result is contiguous for whatever reads it next. Why here and not at the call site: 445b92063 fixed exactly this by flattening inside `BlockwiseFP8Linear` before delegating. That covers activations routed through blockwise FP8, and `lm_head` is not one — DeepSeek ships `lm_head.weight` in BF16 with no `weight_scale_inv`, so `blockwise_fp8::linear_no_bias` short-circuits straight to `UnquantLinear`. The defect is in `UnquantLinear`; fixing it there is the general case and subsumes the caller-side workaround. Every other `bias: None` rank-3 linear that reaches this method is fixed by the same three lines. One latent arm had to be closed on the way. The bias-on-CUDA branch called `cublaslt.batch_matmul` unconditionally, and `batch_matmul` calls `dims3()` on both operands — so a 2-D activation errored instead of falling back. That was unreachable-by-accident only because every caller passed rank 3. It now carries the same `rank() >= 3` guard the no-bias arm already had. The GEMV pre-check deliberately stays ahead of the flatten. `should_use_gemv` reads `product(dims[..rank-1])`, which flattening does not change, so the ordering is belt-and-braces rather than load-bearing — but it makes that a fact instead of a coincidence. **UNVERIFIED ON HARDWARE.** The 67 ms is measured (CUDA-event `device_span`); the bound, the 244x and the 271 GB are derived from shapes and H200 peaks. No speedup is claimed here — only that the weight is read once instead of 256 times. Three CPU tests pin what the change owes: a rank-3 forward must agree with its own flattened 2-D result element-for-element across four `[B, T]` splittings, with and without bias; rank 4 must keep its leading dims; and a non-contiguous rank-3 input must still agree, since `reshape` is the only thing between a transposed view and `matmul`. cargo test -p mistralrs-quant --lib: 331 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-quant/src/unquantized/mod.rs | 185 +++++++++++++++++++++++-- 1 file changed, 172 insertions(+), 13 deletions(-) diff --git a/mistralrs-quant/src/unquantized/mod.rs b/mistralrs-quant/src/unquantized/mod.rs index cd46e6847..ed0b1fb71 100644 --- a/mistralrs-quant/src/unquantized/mod.rs +++ b/mistralrs-quant/src/unquantized/mod.rs @@ -65,12 +65,62 @@ impl QuantMethod for UnquantLinear { // Batch matrix multiplication maybe_init_cublas_lt_wrapper(a.device().clone()); - // Try custom GEMV for single-token decode (batch_size=1) + // Try custom GEMV for single-token decode (batch_size=1). + // + // Deliberately tested on the *original* shape, before the flatten + // below: `should_use_gemv` reads `product(dims[..rank-1])` against + // `MAX_GEMV_BATCH_SIZE`, so flattening first would not change which + // side of that threshold a tensor lands on, but keeping the order + // makes that a fact rather than a coincidence. #[cfg(feature = "cuda")] if crate::gemv::should_use_gemv(a, &self.w) { return crate::gemv::gemv(a, &self.w, self.b.as_ref()); } + // 🔴 A rank-3 activation must be flattened to 2-D, or the weight is + // re-read from HBM once per batch row. + // + // Every activation in this engine is rank-3 `[B, T, hidden]`. Left + // that way, the `match` below took `broadcast_left(B)` — a stride-0 + // view of the weight — and both dispatch arms then treated `B` as a + // GEMM batch dimension with `stride_b = 0`. At decode `T = 1`, so that + // is **B independent `m = 1` GEMVs over the same weight**: a batched + // GEMV wearing a GEMM's name. + // + // The cost is not subtle. V4's `lm_head` is `[256, 1, 4096] x + // [4096, 129280]` BF16, whose bound is + // `max(271 GFLOP / 989 TFLOP/s, 1.06 GB / 4.8 TB/s) = 0.27 ms`. + // Measured on an H200 it is **67 ms/step, 8.41% of a 794 ms B=256 + // decode step** — 244x over, and `244 ~= B = 256` is the whole + // explanation. Back-solving, 256 x 1.06 GB = 271 GB of weight traffic + // per step at 4.05 TB/s: the kernel is running at 84% of peak + // bandwidth while doing 256x the necessary reads. The 1.06 GB weight + // does not fit in L2, so nothing is reused across rows. + // + // Flattened, the same work is one `[256, 4096] x [4096, 129280]` GEMM: + // the weight is streamed once and the tensor cores see `m = 256` + // instead of `m = 1`. + // + // The identical defect was fixed one level up in `BlockwiseFP8Linear` + // (445b92063) by flattening before delegating here. That fix only + // covers activations routed through blockwise FP8 — `lm_head` is not, + // because DeepSeek ships `lm_head.weight` in BF16 with no + // `weight_scale_inv` and `blockwise_fp8::linear_no_bias` short-circuits + // straight to `UnquantLinear`. Fixing it here is the general case and + // subsumes the caller-side workaround. + // + // UNVERIFIED ON HARDWARE: derived from shapes and H200 peaks, not + // measured. The arithmetic above is the claim; the speedup is not. + if a.rank() > 2 { + let batch_dims = a.dims()[..a.rank() - 1].to_vec(); + let rows: usize = batch_dims.iter().product(); + let features = a.dim(D::Minus1)?; + let out = self.forward(&a.reshape((rows, features))?)?; + let mut out_dims = batch_dims; + out_dims.push(out.dim(D::Minus1)?); + return out.reshape(out_dims); + } + let w = match *a.dims() { [b1, b2, _, _] => self.w.broadcast_left((b1, b2))?, [bsize, _, _] => self.w.broadcast_left(bsize)?, @@ -91,21 +141,34 @@ impl QuantMethod for UnquantLinear { match a.device().location() { DeviceLocation::Cuda { .. } => { - // Try to use cublaslt, otherwise fallback to gemm + // Try to use cublaslt, otherwise fallback to gemm. + // + // The `rank() >= 3` guard mirrors the no-bias arm below and + // is now load-bearing: `batch_matmul` calls `dims3()` on + // both operands (`cublaslt/api.rs`), so a 2-D activation + // errors out rather than falling back. Before the flatten + // above that was unreachable-by-accident — every caller + // passed rank 3 — which is exactly the kind of latent arm + // a shape change turns into a crash. if let (Device::Cuda(_), Some(cublaslt)) = (a.device(), CUBLASLT_CONTROLLER.get_for_device(a.device())) { - cublaslt - .batch_matmul( - a, - &w, - Some(&b.t()?.contiguous()?), - None, - Some(1.0), - None, - None, - )? - .t() + if a.rank() >= 3 && w.rank() >= 3 { + cublaslt + .batch_matmul( + a, + &w, + Some(&b.t()?.contiguous()?), + None, + Some(1.0), + None, + None, + )? + .t() + } else { + let matmul_result = MatMul.matmul(a, &w.t()?)?; + matmul_result.broadcast_add(&b) + } } else { let matmul_result = a.matmul(&w.t()?)?; matmul_result.broadcast_add(&b) @@ -664,3 +727,99 @@ impl QuantizedSerde for UnquantLinear { )) } } + +#[cfg(test)] +mod tests { + use super::*; + + fn lin(out: usize, inp: usize, bias: bool) -> Result { + let dev = Device::Cpu; + // Deterministic, non-symmetric weights: a constant or symmetric matrix + // would make a transposed or broadcast result compare equal by luck. + let w: Vec = (0..out * inp) + .map(|i| ((i % 17) as f32 - 8.0) / 16.0) + .collect(); + let w = Tensor::from_vec(w, (out, inp), &dev)?; + let b = if bias { + let v: Vec = (0..out).map(|i| (i % 5) as f32 - 2.0).collect(); + Some(Tensor::from_vec(v, out, &dev)?) + } else { + None + }; + UnquantLinear::new(QuantMethodConfig::Unquantized(Linear::new(w, b))) + } + + fn act(rows: usize, inp: usize) -> Result { + let v: Vec = (0..rows * inp) + .map(|i| ((i % 23) as f32 - 11.0) / 32.0) + .collect(); + Tensor::from_vec(v, (rows, inp), &Device::Cpu) + } + + /// **The flatten must not change the answer.** + /// + /// `forward` now reshapes any rank-3 activation to `[B*T, hidden]` so the + /// weight is streamed once instead of once per batch row. That is a + /// dispatch change on the hottest path in the engine, so what it owes is + /// exact agreement with the 2-D result it is derived from — element for + /// element, not "close". + #[test] + fn rank3_forward_matches_the_flattened_2d_result() -> Result<()> { + for bias in [false, true] { + let l = lin(7, 5, bias)?; + let flat = act(6, 5)?; + let want = l.forward(&flat)?; + + for shape in [(6usize, 1usize), (3, 2), (2, 3), (1, 6)] { + let a = flat.reshape((shape.0, shape.1, 5))?; + let got = l.forward(&a)?; + assert_eq!( + got.dims(), + &[shape.0, shape.1, 7], + "batch dims must survive the round trip (bias={bias})" + ); + let diff = (got.reshape((6, 7))? - &want)?.abs()?.max_all()?; + assert_eq!( + diff.to_scalar::()?, + 0.0, + "rank-3 {shape:?} disagreed with the 2-D result (bias={bias})" + ); + } + } + Ok(()) + } + + /// Rank 4 is the `[b1, b2, T, hidden]` shape the old `match` had its own + /// arm for; it flattens to the same three leading dims. + #[test] + fn rank4_forward_keeps_its_leading_dims() -> Result<()> { + let l = lin(3, 4, false)?; + let a = act(8, 4)?.reshape((2, 2, 2, 4))?; + let got = l.forward(&a)?; + assert_eq!(got.dims(), &[2, 2, 2, 3]); + let want = l.forward(&act(8, 4)?)?; + let diff = (got.reshape((8, 3))? - &want)?.abs()?.max_all()?; + assert_eq!(diff.to_scalar::()?, 0.0); + Ok(()) + } + + /// A non-contiguous activation reaches `reshape`, which candle services + /// with a strided copy rather than an error. Pinning it here because the + /// flatten is the only thing standing between a transposed view and + /// `matmul`. + #[test] + fn non_contiguous_rank3_input_still_agrees() -> Result<()> { + let l = lin(3, 4, false)?; + // [2, 4, 3] -> transpose to [2, 3, 4]: last dim is no longer the + // fastest-varying axis. + let base = act(8, 3)?.reshape((2, 4, 3))?; + let a = base.transpose(1, 2)?; + assert!(!a.is_contiguous()); + let got = l.forward(&a)?; + assert_eq!(got.dims(), &[2, 3, 3]); + let want = l.forward(&a.contiguous()?.reshape((6, 4))?)?; + let diff = (got.reshape((6, 3))? - &want)?.abs()?.max_all()?; + assert_eq!(diff.to_scalar::()?, 0.0); + Ok(()) + } +} From 10448bbe41dbb7f47daa6abb2392ab8b5effc94f Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:24:40 +0100 Subject: [PATCH 3/8] fix(ArcGate): the rental gate certified "V4 fits in 60 GB" from a run that loaded on CPU MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `arc validate --target-hbm` measures residency as `used_bytes = free_before.saturating_sub(free_after)`. That expression is `0` whenever the model did not become resident — loaded on the CPU with no device map, snapshots straddling a CUDA driver state reset, or free memory simply not going down. The verdict was `pass = total_gb <= target`, so zero gave `0.0 <= 60.0` — **true**. The tool printed `result: PASS`, wrote `"pass": true` into the JSON the rental step reads, and exited `0`. The zero case was already detected. It pushed a free-form `note`. A note changes no exit code and nothing in CI reads one. A gate that cannot tell "it fits" from "we did not look" is worse than no gate: it manufactures the exact number a buying decision rests on. Three changes: * `hbm_verdict` returns `(pass, unmeasured)`, and `pass` is `false` whenever the delta is zero. * `HbmReport` carries `unmeasured` (`#[serde(default)]`, so reports written before this field still parse). The printed verdict gains a third line. * `exit_code` is three-valued: **0 fits, 1 measured-and-does-not-fit, 2 we do not know**. `2` is reserved for the environment and never `1`, matching what `main.rs` already exits with for a bad `--compression-stack` or a load error — so callers need no new rule: non-zero-and-not-1 means look at the box, not the model. Collapsing "unmeasured" into `1` would read as "V4 does not fit", the opposite error to the one this code used to make and just as wrong. Both new functions live **outside** `#[cfg(feature = "cuda")]`, and that is the structural half of the fix. The defect sat in CUDA-gated code that no CPU CI job compiles, let alone runs, so a test next to the bug would have been exactly as unreachable as the bug. Moving the decision into always-compiled code is what lets `zero_hbm_delta_is_unmeasured_not_a_pass` go red on a laptop. Four tests: the zero case is not a pass; a real 40 GiB still passes and a real 80 GiB still fails, so the check did not swallow the measurement it guards; the exit mapping stays three-valued with `unmeasured` dominating `pass`; and a legacy JSON report without the new field still parses. cargo test -p arc-cli --bins validate: 15 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- arc-cli/src/validate.rs | 196 ++++++++++++++++++++++++++++++++++++++-- 1 file changed, 190 insertions(+), 6 deletions(-) diff --git a/arc-cli/src/validate.rs b/arc-cli/src/validate.rs index ecb09b895..ee3a78db4 100644 --- a/arc-cli/src/validate.rs +++ b/arc-cli/src/validate.rs @@ -18,6 +18,24 @@ //! //! The output JSON is the contract that the rental verification step (and //! downstream CI) reads. Schema is intentionally narrow and stable. +//! +//! # Exit codes +//! +//! | code | meaning | who acts | +//! |---|---|---| +//! | `0` | measured, and it fits | nobody | +//! | `1` | measured, and it does not fit | the model owner | +//! | `2` | **not measured** — the run proves nothing | the operator | +//! +//! The third row is the whole point of [`hbm_verdict`]. A zero HBM delta — +//! model loaded on the CPU, snapshots straddling a driver reset, free memory +//! not going down — used to satisfy `0.0 <= target` and exit `0`, so this gate +//! would certify "V4 fits in 60 GB on this H100" off a run where V4 never +//! reached the H100. **A gate that cannot distinguish "it fits" from "we did +//! not look" is worse than no gate**, because it manufactures the exact +//! number a rental decision is made on. `2` is reserved for the environment, +//! never `1`, so that a red lane always says whether to look at the model or +//! at the box. use anyhow::{anyhow, Context, Result}; use serde::{Deserialize, Serialize}; @@ -93,8 +111,22 @@ pub struct HbmReport { pub target_gb: f64, /// Measured (or mocked) memory breakdown. pub measured: MeasuredFootprint, - /// True when measured.total_gb <= target_gb. + /// True when the run produced a usable measurement **and** + /// `measured.total_gb <= target_gb`. + /// + /// Never true when [`HbmReport::unmeasured`] is set: a run that measured + /// nothing has not shown that the model fits, and `0.0 <= target` is not + /// evidence of anything. pub pass: bool, + /// The run did not produce a usable measurement — the model never became + /// resident on the device. + /// + /// This is the difference between "V4 does not fit in 60 GB" and "we do + /// not know whether V4 fits in 60 GB", and it is the difference between + /// exit 1 and exit 2. `#[serde(default)]` so reports written before this + /// field existed still parse, as `false`. + #[serde(default)] + pub unmeasured: bool, /// "mock" or "gpu". pub mode: String, /// GPU descriptor (None for mock path). @@ -230,6 +262,46 @@ fn estimate_kv_bytes() -> u64 { const GB: f64 = 1024.0 * 1024.0 * 1024.0; /// Convert bytes to gigabytes (binary GB). +/// 🔴 Decide `(pass, unmeasured)` from one post-load HBM delta. +/// +/// **A zero delta is an environment failure, not a pass.** +/// +/// The real-GPU path computes `used_bytes = free_before.saturating_sub( +/// free_after)`. If the model loaded on the CPU, or the two snapshots straddled +/// a CUDA driver state reset, or free memory simply did not go down, that +/// expression is `0` — and the verdict used to be `total_gb <= target`, i.e. +/// `0.0 <= 60.0`, i.e. **true**. The gate printed `PASS` and exited `0`. That +/// is this tool certifying "V4 fits in 60 GB on this H100" from a run in which +/// V4 never touched the H100, and it is the worst failure a rental gate can +/// have: the number it invents is the number a buying decision rests on. The +/// zero case *was* detected — and pushed a free-form `note`. A note changes no +/// exit code and nothing in CI reads one. +/// +/// This lives outside `#[cfg(feature = "cuda")]` **on purpose**. The defect sat +/// in CUDA-gated code that no CPU CI job ever compiled, let alone ran, so a +/// test next to it would have been as unreachable as the bug. Keeping the +/// decision in always-compiled code is what makes +/// [`tests::zero_hbm_delta_is_unmeasured_not_a_pass`] able to fail on a laptop. +pub fn hbm_verdict(used_bytes: u64, total_gb: f64, target_hbm_gb: f64) -> (bool, bool) { + let unmeasured = used_bytes == 0; + (!unmeasured && total_gb <= target_hbm_gb, unmeasured) +} + +/// Process exit code for a finished report: `0` fits, `1` measured and does not +/// fit, `2` we do not know. +/// +/// Kept separate from [`run`] so the mapping is testable without a GPU, a +/// model download, or stdout capture. +pub fn exit_code(report: &HbmReport) -> i32 { + if report.unmeasured { + 2 + } else if report.pass { + 0 + } else { + 1 + } +} + fn to_gb(bytes: u64) -> f64 { bytes as f64 / GB } @@ -280,6 +352,9 @@ pub fn run_mock(opts: &HbmValidateOptions) -> Result { kv_estimate_gb: to_gb(kv_bytes), }, pass, + // The mock path computes rather than measures, so it always produces a + // number. It can be wrong; it cannot be absent. + unmeasured: false, mode: "mock".to_string(), gpu: None, notes, @@ -441,7 +516,7 @@ pub fn run_real_gpu(opts: &HbmValidateOptions) -> Result { let weight_bytes = used_bytes.saturating_sub(workspace_bytes + kv_bytes); let total_gb = to_gb(used_bytes); - let pass = total_gb <= opts.target_hbm_gb; + let (pass, unmeasured) = hbm_verdict(used_bytes, total_gb, opts.target_hbm_gb); let mut notes = vec![ format!( @@ -457,10 +532,11 @@ pub fn run_real_gpu(opts: &HbmValidateOptions) -> Result { .unwrap_or("none") ), ]; - if used_bytes == 0 { + if unmeasured { notes.push( - "WARN: post-load HBM delta is 0 — either the model loaded entirely on CPU \ + "UNMEASURED: post-load HBM delta is 0 — either the model loaded entirely on CPU \ (no device map?) or the snapshots straddle a CUDA driver state reset. \ + This run proves nothing about residency; exit code is 2, not 0 or 1. \ Inspect logs." .to_string(), ); @@ -485,6 +561,7 @@ pub fn run_real_gpu(opts: &HbmValidateOptions) -> Result { kv_estimate_gb: to_gb(kv_bytes), }, pass, + unmeasured, mode: "gpu".to_string(), gpu: Some(GpuInfo { name: detect_gpu_name(), @@ -574,7 +651,9 @@ pub fn run(opts: HbmValidateOptions) -> Result { println!(" total: {:>7.2} GB", report.measured.total_gb); println!(" target: {:>7.2} GB", report.target_gb); println!(); - if report.pass { + if report.unmeasured { + println!(" result: UNMEASURED (no HBM was consumed — this run proves nothing)"); + } else if report.pass { println!(" result: PASS (measured <= target)"); } else { println!( @@ -592,7 +671,16 @@ pub fn run(opts: HbmValidateOptions) -> Result { println!(); println!("Report written to: {}", opts.output_path.display()); - Ok(if report.pass { 0 } else { 1 }) + // 0 = fits. 1 = measured, and does not fit. 2 = we do not know. + // + // The three must stay distinct. `1` is a verdict about the model and CI is + // entitled to act on it; `2` is a verdict about the run and means the + // operator has to look at the box. Collapsing "unmeasured" into `1` would + // read as "V4 does not fit" — the opposite error to the one this code used + // to make, and just as wrong. `2` also matches what `main.rs` already exits + // with for a bad `--compression-stack` or a load error, so callers need no + // new rule: **non-zero-and-not-1 means the environment, not the model.** + Ok(exit_code(&report)) } #[cfg(test)] @@ -776,6 +864,7 @@ mod tests { kv_estimate_gb: 5.0, }, pass: true, + unmeasured: false, mode: "mock".into(), gpu: None, notes: vec![], @@ -816,6 +905,101 @@ mod tests { assert_eq!(parsed.mode, "mock"); } + /// 🔴 **The regression this file exists to never repeat.** + /// + /// `used_bytes = free_before.saturating_sub(free_after)` is `0` whenever + /// the model did not become resident. The old verdict was + /// `pass = total_gb <= target` — `0.0 <= 60.0` — so the gate printed PASS + /// and exited 0 off a run in which nothing loaded onto the GPU. + /// + /// This test is in always-compiled code deliberately: the defect lived + /// inside `#[cfg(feature = "cuda")]`, which no CPU CI job compiles, so a + /// test next to the bug would have been as unreachable as the bug. + #[test] + fn zero_hbm_delta_is_unmeasured_not_a_pass() { + let (pass, unmeasured) = hbm_verdict(0, 0.0, 60.0); + assert!( + !pass, + "a zero HBM delta must never pass — 0.0 <= 60.0 is not evidence" + ); + assert!( + unmeasured, + "a zero HBM delta must be reported as unmeasured" + ); + } + + /// The zero check must not swallow real results in either direction. + #[test] + fn nonzero_hbm_delta_is_judged_against_the_target() { + // 40 GiB measured against a 60 GB target: fits. + let used = 40 * 1024 * 1024 * 1024; + let (pass, unmeasured) = hbm_verdict(used, to_gb(used), 60.0); + assert!(pass && !unmeasured, "40 GiB must pass a 60 GB target"); + + // 80 GiB measured against the same target: a real failure, and it must + // stay distinguishable from "we did not look". + let used = 80 * 1024 * 1024 * 1024; + let (pass, unmeasured) = hbm_verdict(used, to_gb(used), 60.0); + assert!(!pass, "80 GiB must fail a 60 GB target"); + assert!( + !unmeasured, + "80 GiB was measured — it is a model verdict, not an env one" + ); + + // Exactly on the line still passes; the contract is `<=`. + let (pass, _) = hbm_verdict(1, 60.0, 60.0); + assert!(pass, "the target is inclusive"); + } + + /// Exit codes must stay three-valued: 0 fits, 1 does not fit, **2 unknown**. + /// + /// Collapsing unknown into 1 would read as "V4 does not fit" — the opposite + /// error to the one this code used to make, and just as wrong. + #[test] + fn exit_code_separates_the_model_verdict_from_the_environment() { + let mk = |pass: bool, unmeasured: bool| HbmReport { + model: "m".into(), + compression_stack: "bf16".into(), + target_gb: 60.0, + measured: MeasuredFootprint { + total_gb: 0.0, + weight_gb: 0.0, + workspace_gb: 0.0, + kv_estimate_gb: 0.0, + }, + pass, + unmeasured, + mode: "gpu".into(), + gpu: None, + notes: vec![], + }; + assert_eq!(exit_code(&mk(true, false)), 0, "fits"); + assert_eq!(exit_code(&mk(false, false)), 1, "measured, does not fit"); + assert_eq!(exit_code(&mk(false, true)), 2, "unmeasured"); + // Belt and braces: even a report that somehow claims both must not + // report success, because `unmeasured` is the stronger statement. + assert_eq!(exit_code(&mk(true, true)), 2, "unmeasured wins over pass"); + } + + /// An older report on disk, written before `unmeasured` existed, must still + /// parse — the JSON file is a cross-version contract with the rental step. + #[test] + fn reports_without_the_unmeasured_field_still_parse() { + let legacy = r#"{ + "model": "deepseek-ai/DeepSeek-V4-Flash", + "compression_stack": "qtip2+td-moe", + "target_gb": 60.0, + "measured": {"total_gb": 55.0, "weight_gb": 50.0, "workspace_gb": 4.0, "kv_estimate_gb": 1.0}, + "pass": true, + "mode": "gpu", + "gpu": null, + "notes": [] + }"#; + let parsed: HbmReport = serde_json::from_str(legacy).expect("legacy report parses"); + assert!(parsed.pass); + assert!(!parsed.unmeasured, "absent means false, not unknown"); + } + #[test] fn run_end_to_end_returns_one_when_target_too_tight() { let tmp = TmpDir::new(); From 2fde72a891d88fc1e428754bf488110be0cdf3c6 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:26:55 +0100 Subject: [PATCH 4/8] docs(ArcGate): commit the seven mission docs the code cites but nobody else could read MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Rust comments across `mistralrs-core`, `mistralrs-quant` and `mistralrs-core/src/paged_attention` cite seven mission documents as the evidence behind load-bearing decisions. None of them was in the repository. They were sitting untracked in one laptop's `~/.claude/projects/...` directory — not committed, not gitignored, just absent. `memory/mission/` already tracks 47 files, so nothing was stopping them; they were never added. `CEILINGS.json` is the file this project's own rules order an agent to read before quoting any speed number or saying something is not achievable. Anyone cloning this repo has been unable to read it. Committed at the cited paths, so no citation had to be rewritten to reach them: * `CEILINGS.json` — physics vs implementation bounds, cited from `mtp_pipeline.rs` and `qtip_grouped_curve.rs` * `BUDGET_V4_PREFILL.md` — cited from `inputs_processor.rs`, `gather_policy.rs` * `BUDGET_V4_B1.md` — cited from `00_RESUME_HERE.md`, `COMPETITIVE_TEARDOWN.md` * `EXTERNAL_FINDINGS.md` — cited twice from `mtp_pipeline.rs` * `wave13-AH-config.md` — cited from `deepseek4.rs:7517` * `wave16-AF-beam-perf.md` — cited from `qtip/search_bench.rs:743` * `wave45-BW-mtp-batched.md` — see below The seventh citation was rot, not absence. `segmented_allocator.rs:15` cited `memory/mission/wave59-CJ-mtp-batch.md` §3, and no such file exists anywhere — not in the repo, not on the laptop, not in any worktree. The claim it carries is real: it is in `wave45-BW-mtp-batched.md` **§2**, which states the shared cache rolls back to `C + min_i(u_i + accepted_i)`, "the one length every sequence can prove". So the pointer was wrong twice over — wrong file, wrong section — and pointed at a claim that does exist. Both fixed, and the quoted formula now matches the document verbatim. Worth naming, because this session's other commit widened the citation gate and it did **not** catch this one: the gate matches `path.ext:N`, and all seven of these are cited by bare filename or with a `§` section. A path citation with no line number is invisible to it. That is a real hole in the widening and a follow-up, not something to fix by loosening a rule that currently has zero false positives. The seven docs bring 17 citations of their own into scope, all into sglang, CUDA headers and DeepSeek configs that Arc reads and does not vendor; pinned in `BASELINE`. Census 1,244 -> 1,323, still 0 live findings. cargo test -p mistralrs-core --test doc_citations: 12 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- memory/mission/BUDGET_V4_B1.md | 664 ++++++++++++++++++ memory/mission/BUDGET_V4_PREFILL.md | 329 +++++++++ memory/mission/CEILINGS.json | 296 ++++++++ memory/mission/EXTERNAL_FINDINGS.md | 99 +++ memory/mission/wave13-AH-config.md | 455 ++++++++++++ memory/mission/wave16-AF-beam-perf.md | 427 +++++++++++ memory/mission/wave45-BW-mtp-batched.md | 223 ++++++ .../paged_attention/segmented_allocator.rs | 9 +- mistralrs-core/tests/doc_citations.rs | 106 +++ 9 files changed, 2604 insertions(+), 4 deletions(-) create mode 100644 memory/mission/BUDGET_V4_B1.md create mode 100644 memory/mission/BUDGET_V4_PREFILL.md create mode 100644 memory/mission/CEILINGS.json create mode 100644 memory/mission/EXTERNAL_FINDINGS.md create mode 100644 memory/mission/wave13-AH-config.md create mode 100644 memory/mission/wave16-AF-beam-perf.md create mode 100644 memory/mission/wave45-BW-mtp-batched.md diff --git a/memory/mission/BUDGET_V4_B1.md b/memory/mission/BUDGET_V4_B1.md new file mode 100644 index 000000000..7dd134e8c --- /dev/null +++ b/memory/mission/BUDGET_V4_B1.md @@ -0,0 +1,664 @@ +# THE V4 B=1 DECODE BUDGET — where the 98× goes + +**This file answers "why is V4 at ~1% of the roofline". It is MEASURED, not estimated. +Do not re-derive any number here. Do not reason about V4 decode speed without reading it.** + +Provenance: box `arc-v4-stack` (H200), SHA `05af600e7` (= `origin/master`, PR #120 confirmed +ancestor), `--features "cuda flash-attn"`, model `aeonmind/DeepSeek-V4-Flash-UQFF-qtip2` over +`deepseek-ai/DeepSeek-V4-Flash`. Device-timer selftest PASS, ratio 26.8×. Session 8, 2026-08-18. + +--- + +## The headline + +| | ms/step | % | +|---|---|---| +| **Measured step** (uninstrumented `mistralrs bench`, B=1, **15.0 tok/s**) | **66.68** | 100% | +| **GPU executing kernels** | **34.04** | **51%** | +| **GPU idle, waiting on the host** | **32.64** | **49%** | +| memory-bandwidth floor (13B active × 2 bit ÷ 4.8 TB/s) | 0.677 | — | + +**66.68 / 0.677 = 98×.** It factors as **~2× (GPU idle) × ~50× (the kernels themselves)**. + +Profiler-accounted step wall 70.1 ms agrees with the clean 66.68 within 1.5% — that agreement is +what makes the span budget below trustworthy. `ARC_PROFILE` overhead measured at **5–8%**, not 4×. + +Span budget (device stream-elapsed, 320 recorded steps, µs/step): `mla_attn` 27,689 (40.7%) · +`moe` 23,571 (34.7%) · `mhc_ffn_pre` 7,085 (10.4%) · `mhc_attn_pre` 5,993 (8.8%) · +`mix_post_*` 2,521 · `lm_head` 281. The `layer` node reconciles exactly: `device_self_ns = 0`, +zero violations, zero misnested, zero unresolved. + +--- + +## THE DIAGNOSIS: op count, not slow kernels + +- **9,131 kernel launches per token** (212 per layer). Step count N=189 pinned from three agreeing + anchors (517,120 B logits D2H × 189; `topk_kernel` 40.1/step = 40 MoE layers; `sinkhorn` + 86.2/step = 2 × 43). +- **11,436 `cuMemAllocAsync` + 11,436 `cuMemFreeAsync` per token — 1.25 allocations per kernel.** + > ✅ **SOLVED (mostly), 2026-08-18** — branch `run-161-arena-counters` @ `dac70e6`, binary `0f1a6125`: + > | per token | baseline | arena | + > |---|---|---| + > | `cuMemFreeAsync` | 11,406.2 | **0.0 (−100%)** | + > | `cuMemAllocAsync` | 11,406.2 | **~288/step (interval) — −97.5%** | + > | layout H2D | 2,361.6 | **6.0 (−99.7%)** | + > | **total host calls** | **25,622** | **~737 (−97%)** | + > + > ### 🔑 A CUMULATIVE AVERAGE CANNOT SHOW A STEADY STATE — use the INTERVAL rate + > The first-quoted `181.4 → 234.5` allocs/step were **run averages**, and *"a run average over a + > warming cache is the same trap as a distinct-size count"* — **warm-up allocations stay in it + > forever, decaying like 1/steps toward a floor they never reach.** The honest steady-state figure + > is the rate **between consecutive reports**: **287.6/step over steps 20→40, and rising.** + > Interval extractor **falsified both ways**: reports **0/step** where an interval is genuinely + > empty (the average would have claimed 2.50/step) and **300/step** for a real ongoing cost. + > ⇒ **Same family as "distinct count ≠ occurrence count" and "absent ≠ stale": a statistic that + > mixes history with the present cannot answer a question about the present.** + > + > **Bit-identical over all 52 steps** (FNV over raw IEEE logit bits), with `A1 vs F DIFFER` proving + > the comparison returns the negative and `A1 vs A2 MATCH` proving the harness is deterministic. + > 8,554 cache hits/step, `accounting OK`, `puts_refused 0`, arena high-water **956 MB**. + > **`driver_frees` is EXACTLY 0** — with the arena on, every `CudaStorage::drop` becomes a cache + > put, so `cuMemFreeAsync` is never called. **Half the 1,535 `MEM_ALLOC`/`MEM_FREE` capture gate + > has no remaining source.** + > + > ⚠️ **The allocation side is RISING, not decaying — the chain's own prediction refuted, and it + > said so.** Cumulative 181.4/step at 20 steps → 234.5/step at 40 ⇒ the **interval** rate over + > 20→40 was **287.6/step**. Bucket-decay model treated as wrong pending the `--gen-len 3000` run. + > **Two levers:** coarser bucketing for large buffers (1/2-octave: ~45 → ~12 distinct sizes per + > growing tensor, costs only memory against ~65 GB headroom); and **graph mode's fixed-capacity KV + > slot makes shapes step-invariant, removing the growing-size family at source** — the standing + > prediction for how `driver_allocs == 0` is reached *in the capture window specifically*, and the + > only part still unmeasured. + > + > ### 🔴 GATE FAILS IN EAGER MODE — THE ARENA LEAKS, mechanism named (500-token run) + > | step | arena MB | buffers | `puts_refused` | `driver_frees`/tok | + > |---|---|---|---|---| + > | 100 | 2,098 | 44,150 | 0 | **0.0** | + > | 200 | 4,527 | 97,716 | 0 | **0.0** | + > | 300 | 6,878 | 156,102 | 0 | **0.0** | + > | 400 | **8,192 (CAP)** | 212,305 | 15,822 | 39.6 | + > | 500 | **8,192 (CAP)** | 261,153 | **119,828** | **239.7** | + > + > **Interval `driver_allocs` over steps 400→500 = 1,401/step. `GATE driver_allocs == 0: FAIL`.** + > + > **Mechanism:** the arena grows **~24 MB and ~520 buffers per step and never returns any.** + > 1/8-octave bucketing does **not** collapse the growing-`kv_len` family — it hands each step a + > fresh set of sizes that are **cached forever and never requested again.** + > **🔑 THE CAP POLICY IS EXACTLY BACKWARDS: under pressure it refuses NEW puts while retaining the + > OLD buffers — and the old ones, at sizes that will never be requested again, ARE the dead + > weight.** Past the cap the arena degrades **below its own purpose**: `cuMemFreeAsync` returns + > from 0.0 to **239.7/token**. + > **Fix (neither speculative): EVICTION, not refusal** — when the cap binds, drop stale entries and + > keep fresh ones; plus **power-of-two bucketing for large buffers** (~45 → ~7 distinct sizes per + > growing tensor). + > + > **What still stands:** `cuMemFreeAsync` **0.0/token** and layout H2D **6.0/token** held **flat + > across all 500 steps until the cap bound**, `accounting OK` throughout, output bit-identical. + > ⚠️ **And the earlier 287.6/step reading was itself an early-warmup artefact of a RISING curve** — + > the reason run-averages were abandoned applies to short interval samples too. + > + > **`puts_refused` is the only reason this was visible.** Without it the arena reads healthy — + > `accounting OK`, hits still ~8,200/step — **while silently reverting to driver frees.** Sixth + > instance of a metric that must be structurally impossible being the one that carries the signal. + > + > **Two defects found in this work by instruments, not inspection:** symmetric bucketing filed + > buffers under keys **larger** than physical size ⇒ pre-arena buffers handed out short ⇒ prefill + > `CUDA_ERROR_INVALID_VALUE` (**a short buffer is a wrong-numbers bug by nature; it surfaced as a + > hard error only by luck**); then that fix returned `0` below 128 B ⇒ every small buffer allocated + > *and* freed through the driver **while the arena reported itself healthy**. Tell: `driver_frees` + > non-zero at all — a quantity that is supposed to be structurally impossible. +- **2,818 H2D copies per token averaging 72 bytes** — candle's per-op dims/strides upload + (`clone_htod`). Not model data. +- Kernel duration **mean 3.73 µs, median 1.18 µs**, p90 4.58 µs. +- **~5,300 of the 9,131 launches (58%) are dtype casts, copies and elementwise ops** at 1–3 µs each + ⇒ **≥8.0 ms/step, ≥24% of all GPU time**. `cast_bf16_f32` + `cast_f32_bf16` alone are + **1,571 launches/token**, largely MHC's F32 round trip in `hc_pre`/`hc_post`. +- Corroboration on the clean run: **memory-controller utilization 4%**, board power **204 W of + 700 W**. The GPU is barely touching HBM. + +--- + +## THE SYNCS ARE INVISIBLE TO THE OBVIOUS GREP + +**`ALL *Synchronize* = 0.0 calls/step.** Anyone counting `cudaStreamSynchronize` concludes +"no syncs" and is wrong.** + +The real synchronisation is **44 `cuMemcpyDtoHAsync_v2` per step at ~109 µs each = 4.81 ms**. +An "async" call taking 109 µs of host time is blocking — pageable-memory staged copy. +**43 per layer + 1 logits.** + +Source: `dsv4_kv_fp8.rs:260` `e4m3_codes_cpu` — a deliberate device→host→device round trip +**because candle has no CUDA F8E4M3 cast**. At B=1 it moves exactly 448 F32 = **1792 bytes**. +Trace shows exactly two DtoH sizes: 1792 B × 8144 and 517,120 B × 189. `KvQuantMode` defaults to +`CpuExact`, whose own doc says *"at the price of one device sync per layer."* + +### 🔑 Removing the 43 syncs makes it SLOWER — and that is the proof of the diagnosis + +Interleaved A/B (`ARC_GPU_ACT_QUANT=1`, i.e. `GpuApprox`): +`A1 66.99 · B1 67.71 · A2 68.05 · B2 70.30 ms/T` — **monotonic 3.3% drift exceeding the arm +difference**, so always interleave A-B-A-B on this box and state the drift. +Engagement proven per D32: `kv_fp8_quant` **73.49 → 134.18 µs/call (+83%)**, `kv_fp8_dequant` +50.09 → 37.54. The sync-free path swaps one blocking copy for **~19 extra elementwise kernels +per layer**. + +> **In a machine bottlenecked on op count, a sync is cheaper than the 800 launches it takes to +> avoid one.** ⇒ **`GpuApprox` is NOT the fix and must not ship as one.** The fix is a single +> **fused quantize+dequantize device kernel**: the round trip costs 5.31 ms/step (8.0%) under the +> default, 11.1% under GpuApprox, and should cost **~2 µs**. + +## ✅ SOLVED — the fused FP8 kernel shipped, and the ~2 µs prediction landed + +Branch `arckv-fp8-fused` (9 commits). Measured on `arc-v4-stack`, private target dir, binary md5 +`e15259dc…`, 118 symbol hits verified in the measured binary. + +| | before | after | +|---|---|---| +| **blocking D2H per step** | **43.89** | **1.00** | +| launches per step | 9,081.8 | **8,404.5** (−677.3) | +| the FP8 round trip | `cuMemcpyDtoHAsync_v2` **48.56 µs/call host = 2.131 ms/step BLOCKING** | quantize **1.98 µs/call** + dequantize **2.34 µs/call** @ 42.93 calls/step = **0.185 ms/step device** | + +The **1,792-byte transfer size is absent from the after-trace entirely.** The one surviving D2H is +the 517,120-byte logits readback — present in both arms, not this subsystem's. +`cuLaunchKernel` −773.5, `cuMemAllocAsync` −758.6. + +### 🏆 The parity proof — exhaustive over the ENTIRE input domain +**All 2^32 f32 values** swept against NVIDIA's software reference (what `float8` ports, hence +candle's CPU cast) **and** the sm_90 hardware path: **0 mismatches**, with a `visited` counter +confirming all **4,294,967,296** were actually touched. Negative control catches **123,731,850 +(2.88%)**. Full-tensor GPU test passes on codes + side bytes, crosswise reconstruction, and the +verbatim pre-packing implementation. +**D33 discharged on the shipped path itself:** pointing it at the truncating mutant produced +`assertion left == right failed: fused E4M3 codes differ from candle's CPU cast`; revert → passed, +worktree clean. The counter was also validated on a known answer first — it reproduced the recorded +44.09 / 9,131.5 / 11,436 / 0.00-Synchronize before being pointed at new data. + +### ⚠️ `--use_fast_math` SILENTLY BREAKS `__f*_rn` INTRINSICS +The chain assumed `__f*_rn` were IEEE. **The PTX showed nvcc 13.1 rewriting them to `.ftz`** under +`--use_fast_math` — flush-to-zero denormals inside a *numeric format conversion*. Replaced with +inline PTX; the audit is now **`grep -c '\.ftz\.f32'` = 0**. Third `--use_fast_math` surprise of the +session (it had also already collapsed the trellis transcendentals). **Check the PTX, not the +intrinsic's name.** + +### No end-to-end tok/s, deliberately +The chain **deleted its A/B driver rather than ship a number a biased estimator can't support** +(D36). Per-call kernel timings + launch counts only — which is why this result survives a night +that destroyed three end-to-end measurements. + +**ArcGraph unblocked from this side:** zero blocking host work remains in the FP8 path; kernels +launch on `dev.cuda_stream().cu_stream()` so they record into a graph captured on candle's stream; +a warm decode step issues no `cuMemAllocAsync` from this op. **The remaining capture blocker is the +NULL-stream binding + the PagedAttention gate.** `GpuApprox` untouched, still reachable, documented +in-code as not the fix. + +--- + +## WHY CUDA GRAPH REPLAY HAS NEVER LAUNCHED — root cause, finally named + +> `arc-cuda-graph/src/graph.rs:71` — `cuda_dev.cuda_stream().cu_stream()` is **NULL**, i.e. CUDA's +> **legacy default stream, on which capture is impossible.** Logged live: +> `CUDA graph: NULL stream, capture disabled`. + +Three independent confirmations it is off: +1. `Decode path extraction failed: tensor_device_ptr requires CUDA tensor` +2. `paged_attention` **UNREACHABLE** (`pipeline/mod.rs:1088`, DefaultInstructions) — this takes + `graph_wrapped_forward`, the dedicated decode path, and graph replay with it +3. `cuda_graph.capture_probe` **UNREACHABLE** + +**Prerequisite chain (ordering matters — graphs cannot capture a blocking D2H):** +**(a)** move the engine off the legacy default stream → **(b)** remove the per-layer D2H *as a +fused device kernel, NOT as `GpuApprox`* → **(c)** capture. + +--- + +--- + +## ✅ THE CPU **CAN** BE TAKEN OUT OF THE LOOP — MEASURED ON THIS H200, 2026-08-18 + +Jish's instruction was *"kill the CPU out of the loop."* **It is achievable on our own hardware +and the number is in.** Measured at `/root/arc-wt/probe`, interleaved A-C-A-C ×4, bench lock held: + +| route | **host calls per decode step** | +|---|---| +| today (graph + sync per step) | **2.0** | +| host-side K-replay | 1.0078 | +| **device-side WHILE conditional node** | **0.0156** at N=128 | + +The 0.0156 is **2 calls per *generation*** (1 launch + 1 sync) — **not per step** — so it goes to +zero as N grows. Device-loop overhead is **constant in body size**: +2.90 µs (M=1), +2.39 (M=64), +**+1.98 µs (M=256)** — against a 66.68 ms step that is **0.004%**. Drift 0.1–1.3%. +**Falsifiable control that makes it real:** one instantiated graph run with device-side `target=1` +then `target=N` returned counter 1 then N — proving the **device**, not the host, chose the trip +count. Three redundant device counters (counter/checksum/filler) agreed in every arm. + +**Residual host work per step once this lands: ZERO issued calls.** Token streaming is already a +host *poll* of pinned memory (`decode_loop.cu:61-66`, `__threadfence_system`), not work issued to +the GPU. + +### What the installed toolkit actually supports (quoted, not remembered) +`/usr/local/cuda-13.1/include/` · H200 **cc 9.0** · driver 580.173.02 · `nvcc` V13.1.115 · +`cuda.h:264 CUDA_VERSION 13010`. +- **WHILE conditional nodes**: `cuda.h:1951 CU_GRAPH_COND_TYPE_WHILE = 1` + device-side + `cuda_device_runtime_api.h:480 cudaGraphSetConditional` — **measured working** +- body population via `cuda.h:15665 cuStreamBeginCaptureToGraph` (`cuda.h:1976` names it the + supported way) +- **device-side graph launch**: `cuda_device_runtime_api.h:290 cudaGraphLaunch` + + `cuda.h:5100 CUDA_GRAPH_INSTANTIATE_FLAG_DEVICE_LAUNCH = 4` — instantiate returned SUPPORTED, + **but the header bars combining it with `AUTO_FREE_ON_LAUNCH`, which `graph.rs:304` uses today** + +### 🔒 HARD DEPENDENCY, confirmed by header not guessed +`cuda.h:1971` permits only *"kernel nodes, empty nodes, child graphs, memsets, memcopies, and +conditionals"* inside a conditional body — **no alloc/free nodes.** ⇒ **the autonomous loop cannot +close until the 11,436 allocations per token are gone.** The workspace-arena chain is the gate. + +### Bugs fixed to get there (compile-verified, D33 negative control: injected type error → exit 101) +- **`cuGraphAddNode` declared with 5 args and the first two transposed** against `cuda.h:21829`'s + 6, whose first parameter is an out-pointer — **the driver wrote through a `CUgraph` it thought + was a `CUgraphNode*`.** ⇒ **this is the cause of the previously unexplained "host heap + corruption" capture failure already on our record.** +- conditional body was captured to a throwaway graph and **destroyed**, leaving the body empty +- missing `CU_GRAPH_COND_ASSIGN_DEFAULT`; missing `ctx` field +- `DecodeState` allocated **I64** while every kernel signature is `int32_t*` +- `reset()` reallocated buffers, **silently invalidating the captured graph's baked pointers** +- `tensor_device_ptr` had **no I32 arm** ⇒ see the sampling defect below + +### ✅✅ THE GATE IS OPEN — 2026-08-18, branch `wt/pagedgate` @ `b137f1e0`, 3 files + +| marker | control (`ARC_NO_AUTONOMOUS_DECODE=1`) | gate open | +|---|---|---| +| `cuda_graph.autonomous_decode` | **ABSENT** | **FIRED, calls=22** | +| `cuda_graph.capture_probe` | **UNREACHABLE, calls=0** | **FIRED, calls=22** | + +Three independent counters agree **live, not at teardown**: `engine_attempts=25 == pipeline_entered=25 +== capture_probe_entered=25`; profiler node calls=22 == decode `forward` calls=22. Binary md5-frozen +(`bff2a901…`), symbol-gated before each run, private `CARGO_TARGET_DIR`, builds under `host.lock -s`. +**Cost of the now-reachable attempt: 780 ns against a 66,680 µs step.** D33: poison injected into the +new code → `RC=101` naming it; restored → clean. Unit tests 5/5 incl. a paged-vs-contiguous negative +control. Xid delta 0. + +**VERDICT: INCIDENTAL, proven.** `AutonomousDecodeRunner::capture` is generic over any +`Fn() -> Result` and **mentions paging nowhere**. `block_size`/`max_blocks_per_seq` are +consumed only to size `block_tables` and feed `slot = block_tables[b][pos/bs]*bs + pos%bs`. +**A dense cache is the one-block case of that exact formula** — `block_size=max_seq_len`, +`max_blocks=1`, `block_tables[b]=b` collapses it to `slot = b*max_seq_len + pos`. +**Paging is a parameterisation, not a capability.** Route (b) separately disproven: V4-with-paging +measured **zero tokens (ILLEGAL_ADDRESS)**, and the runner isn't reached even then. + +**WHY V4 DEFAULTS TO `paged=false` — premise sound, test untouched and still passing.** +`flashinfer_mla_decode.cu` fixes `HEAD_DIM_CKV=512`/`KPE=64`; **V4's 448 trips +`static_assert(vec_size % 8 == 0)` and will not compile.** It computes dense causal attention with +**no sliding window, no sinks, no second key set — which no V4 layer computes.** Plus +`cache_write_and_gather` returns a varlen pack forcing `bs=1`, and V4's compressor `xs` history lives +in `NormalCache` slots the paged arm never clones. ⇒ **the test stayed and the gate moved.** + +**The fix names the real property:** new `DecodeKvGeometry` + `GeneralMetadata::decode_kv_geometry()`; +`normal.rs` gates on **geometry** instead of `cache_config`; `step()`'s DefaultInstructions arm now +routes through `graph_wrapped_forward`; the engine offers decode batches from **both** scheduler arms +via one shared `autonomous_decode_attempt`. Kill-switch/negative control `ARC_NO_AUTONOMOUS_DECODE=1` +and an `autonomous_gave_up` latch so a failed init cannot re-allocate every token. + +> ⚠️ **CORRECTION TO THIS FILE'S OWN RECORD — it was ONE cause with FOUR symptoms, and the outermost +> was never on file.** The engine called `autonomous_decode` **only** from +> `SchedulerOutput::PagedAttention`, so for V4 the method never ran and +> **`mark_unreachable("cuda_graph.autonomous_decode")` NEVER EXECUTED.** Falsifier: the marker is +> **ABSENT**, not "unreachable", from `/root/budget-chain/profiles/v4b1.json` and +> `v4b1_gpuactquant.json`. **The `normal.rs:1907` citation below was a SOURCE READ, not an +> observation** — an unreached `mark_unreachable` is invisible, and "absent" and "unreachable" are +> different states (D24, three-state guards). + +### 🔴🔴 THE KV-WRITE BUG — a **fake 4.9× speedup**, and it is upstream of TWO chains + +**Measured, same card / same session / same binary:** +``` +EAGER 600 tokens 17.4 T/s 57.55 ms/T no errors +GRAPHMODE 600 tokens 85.7 T/s 11.67 ms/T ERROR at prompt step — + kv-write: src [1,1,1,1] incompatible with cache [1,1,512,512] +``` +**The prompt step FAILS, decode proceeds from unwritten KV, and throughput RISES because the model +skipped its work.** ⇒ **Any A/B on this path measures how much work was skipped.** +🔑 **A replay-consistency check would have PASSED — both replay arms agree with each other while +both are wrong. Only replay-vs-EAGER detects it.** + +**Trigger is `ARC_CANDLE_ALLOC_CACHE=1`, NOT the capture probe.** Probe alone → zero kv-write errors +across 7 runs. That one env var does **two unrelated things** inside one `if` at `normal.rs:1605`: +`set_alloc_cache_enabled(true)` **and** `set_graph_mode_positions(...)`. **One flag, two behaviours — +that coupling is what hid this. File them as separate flags.** + +> ⚠️ **THE NODE COUNTS WERE MEASURED ON THE BROKEN FORWARD.** Corrected: +> | config | prompt step | nodes | memory nodes | +> |---|---|---|---| +> | probe + alloc-cache | **FAILS** | 14,274 | **1,535** ← *do not quote as the gate* | +> | probe only (clean) | OK | 32,621 | **20,930** | +> +> ✅ **The gate is SEPARABLE after all:** the arena hook at `normal.rs:1585` is **outside** that `if` +> (gated on `ARC_ARENA` + `seq_len == 1`). ⇒ **`ARC_V4_CAPTURE_PROBE=1 ARC_ARENA=1` with +> `ARC_CANDLE_ALLOC_CACHE` UNSET** = capture + allocator cache + **correct forward**. +> +> ### ✅✅ THE CENSUS ON A VALID FORWARD — 20,930 → **687**, and MEM_FREE is **0** +> | config | prompt step | nodes | MEM_ALLOC | MEM_FREE | +> |---|---|---|---|---| +> | probe + `ARC_CANDLE_ALLOC_CACHE` | **FAILS** | 14,274 | 1,535 combined | — | +> | probe only | OK | 32,621 | 20,930 combined | — | +> | **probe + `ARC_ARENA`** | **OK ✅** | **10,351** | **687** | **0** | +> +> **Prompt step asserted, not assumed: 0 kv-write errors, 0 `prompt step - Model failed` lines.** +> Arena engagement printed: `ARC arena: ON (cap 8192 MB, intern 16384 layouts)` · +> `warmup done after 3 steps — 3407 driver allocs, 97 interned layouts. Counters zeroed.` +> ⇒ **Valid-to-valid: 20,930 → 687 = 96.7% reduction** — *better* than the 93% claimed off the broken +> run, and the discredited 1,535 was never the right baseline. **`MEM_FREE = 0` because the arena +> defers frees entirely, so the ENTIRE residual is 687 allocation nodes. That is the distance to +> capture.** +> +> ⚠️ **RETRACTION by the graph chain, its own:** *"Xid unchanged 1590 → 1590"* held for the +> alloc-cache and probe-only paths, but **the arena run produced 1590 → 1591** — +> `Xid 31, name=mistralrs, MMU Fault … FAULT_PDE ACCESS_TYPE_VIRT_WRITE`, at its run's end. +> **Diagnosis corroborates the earlier one exactly: capture-time allocations are VIRTUAL and never +> materialise unless the graph launches; the eager fallback then writes into them** — plausibly worse +> under the arena because deferred frees leave more such buffers live. +> ⇒ **Refusing to launch is NECESSARY BUT NOT SUFFICIENT. The discarded body's buffers must be +> purged — or better, never created.** +> +> **Provenance (this build is NOT reproducible from git alone):** `arena/decode-zero-host` merged +> into `graphstream` (clean, zero file overlap), plus a `[patch]` pointing candle at +> `/root/arc-wt/candle` **which their `Cargo.toml` never committed**. That checkout was **dirty** +> during the build, so it is pinned by `candle_rev=698dec410`, `device.rs md5=4d5e5b2c`, +> binary `341ecfa2`, verified to contain both the census and `ARC_ARENA` before running. + +### Root cause: an OWNERLESS allocation decision (not a semantics disagreement) +⚠️ **RETRACTED first diagnosis — and it is the one everyone will guess:** *"eager writes real V, +graph writes a 1-wide marker, the phases disagree."* **Wrong. Both paths write the identical +1-wide marker** (`append_kv_mqa` and `append_graph_kv_mqa` both via `v4_v_marker(k)`). + +**Decode reads NOTHING from the V half** — syntactic, not behavioural: both call sites discard the V +readback (`let (k_cached, _marker_cached)`), the graph arm passes `&k_full` as **both** K and V to +`dsv4_attention`, and there is a test `append_kv_mqa_returns_k_for_both_sides`. V4's absorbed MQA +reconstructs V from the same latent as K. **The marker is NOT the bug.** + +**The actual fault:** +> `SingleCache::append` tolerates a mismatched buffer by **reallocating** (`preallocation_fits`). +> `SingleCache::append_graph` **never reallocates** and hard-fails on a width it didn't expect. +> **One design, two tolerances, no owner — and the tolerant one always runs first.** + +**Three allocator candidates, not two** — the third found in `append`'s own comment: on a prompt step +with no prefix hit the engine issues `CacheInstruction::Reset { load_preallocated_cache: true }` and +`NormalCacheManager::set_none_cache` installs the sequence's **preallocated** tensor as `all_data`, +**before any append**. Sized uniformly from `head_dim` with no knowledge of the marker convention ⇒ +**leading suspect.** + +**Structural principle (holds whichever candidate wins): the V half's width must be decided ONCE, +explicitly, at cache construction from model config — never implied by execution order. And +`append_graph` must fail loudly on an unexpected width rather than inherit it.** + +### ✅ FIXED AND PROVEN ON HARDWARE — `realloc_lines=0` +Binary md5 `ce7b8e15…` frozen and gated, `SYMBOL_OK ×3`, card `apps=[] free_mb=143156`, `bench rc=0`: +``` +prefill src=[1,1,64,512] buf=Some([1,1,512,512]) reuse +prefill src=[1,1,64, 1] buf=Some([1,1,512, 1]) reuse ← was REALLOCATE +decode src=[1,1, 1,512] buf=Some([1,1,512,512]) reuse +decode src=[1,1, 1, 1] buf=Some([1,1,512, 1]) reuse ← was REALLOCATE +realloc_lines=0 v_half_reuse_lines=3 CARD WAS EXCLUSIVE BEFORE AND AFTER fail=0 +``` +**Both V-half `REALLOCATE`s replaced by `reuse` against a 1-wide preallocated buffer, in both +phases. The run COMPLETED — this is a trace from a working forward, not from a corpse.** + +**Discrimination, not just counting:** the two remaining `ALLOCATE` lines are **rank-3** +(`[1,16,512]`, `[1,1,512]`) — the xs-history/compressor slots that `set_none_cache` deliberately +resets to `None` because they have no preallocated KV-shaped buffer. **Not KV halves, correctly not +flagged.** A guard that counted `ALLOCATE`s would have failed here; one that classifies them passes +for the right reason. + +**Three changes shipped:** `ModelConfigLike::v_cache_head_dim()` + `KvCacheLayout::FusedMqaMarker`, +with `engine/add_request.rs` sizing the pair through one `preallocated_kv_shapes()` — and +**`kv_cache_elements_per_token` left byte-identical so the fix cannot smuggle a context-size change +in beside it**; `append_graph` now refuses a buffer it cannot write, naming both shapes; and +**`ARC_CANDLE_ALLOC_CACHE` (allocator) split from `ARC_GRAPH_MODE_POSITIONS` (device-slot KV write)**. +5 new tests pass, **2 go red when `v_cache_head_dim` reverts to mirroring `v_head_dim`**, and the +sentinel builds its config through the **production** `v4_model_config_metadata()` from V4's verbatim +`config.json` — **so restating the struct cannot keep it green.** 185 existing tests still pass. + +⏳ **Still running: the three-arm A/B (`eager` / `alloccc` / `graph`) that must show the 4.9× gone.** +**Pass condition is INVERTED and compiled into the script:** it fails if graph exceeds **1.2× eager**, +if the graph arm yields no decode row, or if the `kv-write: src … incompatible` refusal reappears. + +### 🔴 THE TWO REMAINING BLOCKERS — named at the branch that bails, and NEITHER IS PAGING +1. **`cuda_graph.capture_probe.capture`** — candle bound to the **legacy default NULL stream** + ⇒ **graph-stream chain** (`ARC_CAPTURE_STREAM`). +2. **`cuda_graph.autonomous_decode.capture`** — **no V4 forward closure**: + `Decode path extraction failed: tensor_device_ptr requires CUDA tensor`, and the dedicated path + models a **dense q/k/v/o + gate/up/down decoder that cannot express MLA + MoE + mHC** + ⇒ **autonomous-decode chain must pass the V4 candle forward to the generic `capture`.** + *(The runner is already generic — this is a plumbing job, not a redesign.)* +Then the **arena chain** (no alloc/free nodes in a conditional body). + +⚠️ **OPERATIONAL: V4 maps 140 GB of 143.7 GB.** Two control attempts OOM'd — **any neighbour holding +VRAM kills the model load.** V4 timing runs need the card essentially alone. + +--- + +### (historical) THE GATE THAT BLOCKED ALL OF THIS FOR V4 — found by two chains independently + +**V4 does not use PagedAttention.** It takes `DefaultInstructions`, so `cache_config` is `None`, +and **three separate fast paths are constructed only on the PagedAttention arm and therefore never +exist for V4**: `graph_wrapped_forward`, the dedicated decode path, and the autonomous runner +(`normal.rs:1907` → `mark_unreachable("cuda_graph.autonomous_decode")`). The current behaviour is +**pinned by a unit test**: `normal_loaders.rs:5687 v4_supports_paged_attention_defaults_to_false`. +Meanwhile `paged_attention` is itself recorded UNREACHABLE at `pipeline/mod.rs:1088` — the gate is +closed *and* the thing it gates on isn't running. + +⇒ **The measured 2.0 → 0.0156 host-calls-per-step win cannot be switched on for V4 until this is +resolved.** Open question for the owning chain: is the coupling **essential or incidental**? The +runners need stable device pointers and statically-shaped buffers; paged block tables are one way +to get that, not obviously the only way — and V4's KV path is ours (compressed KV, sliding window, +MQA `num_key_value_heads=1`). **Find out why V4 defaults to `paged=false` before touching it; do +not invert the test to make something pass.** + +### ⚠️ SAMPLING DEFECT — REAL, BUT **NOT LIVE IN SERVING** (proven by count, not by call graph) +- **`sampling.cu:150 fused_top_p_bf16` is NOT nucleus sampling.** It walks the vocab in + **token-id order** accumulating full-distribution mass to `top_p*u`, **biasing toward low token + ids**. Its own comment says *"Not quite right — proper top-p needs sorting"*. It documented + itself and shipped anyway. +- **`CudaSampler::sample()` returned `unsupported dtype I32` on EVERY call** ⇒ the correct + top-k/top-p sampler has **never once executed on a GPU**. Its tests are a CPU simulator that + **structurally cannot observe this** — D18 again: the test could not have caught it. +- ✅ **`rng_offset` baking — NOW MEASURED, prediction confirmed exactly.** Identical + capture-and-replay, 64 replays each: **OLD `fused_top_p_bf16` → 1 distinct token** (the by-value + `rng_offset` is baked at capture, so every step of an autonomous loop draws the same uniform). + **NEW `arc_launch_sampler_bf16` → 64 distinct tokens, 64/64 inside the probability mass**; + device-resident `rng_state` advances across replay. **This is the A/B where a broken instrument + could have shown "both fine" and did not.** + ↳ *Retracted by the chain itself:* an earlier claim that the old sampler drew **outside** the + probability mass — that branch never fired, and its single token (121637) was **inside** the + support. The vocab-id-order bias is real **by code reading**; this test did not isolate it. + +### 🛑 STOP-SHIP FOUND IN THE FIX ITSELF — the CORRECT sampler is 663 ms/token + +**`arc_launch_sampler_bf16` measured at 663,116 µs/step = 663 ms per token** — **~10× V4's entire +66.68 ms decode step** — at vocab 129,280 with a flat 12,928-token support and `top_k` disabled. + +Mechanism confirmed at **`sampling_kernel.cu:279`**: `cap = vocab` whenever `top_k <= 0`, and the +keep loop runs a **full O(vocab) argmax scan per kept token in a single block** ⇒ cost is +**O(vocab × kept)**. Two honest qualifiers from the chain: this is the **worst case** (real LLM +distributions are peaked, keeping `kept` in the tens), and the support-width cost curve +(1 → 12,928) that would locate the cliff is **compiled and queued, not returned**. +**But mistral.rs's default maps to `top_k = -1`, so nothing bounds `kept` today.** + +**⛔ Do NOT enable the autonomous sampler as a default until `kept` is bounded.** It is correct and +replay-safe; it is **not yet fast**. + +### 📉 THE COST CURVE — measured, and the fix is a REGRESSION in the common case + +⚠️ From the **unguarded** run: treat absolute µs as **provisional**; the *shape* is robust to any +uniform contamination factor. Guarded re-run pending. + +| support width | legacy µs | bisect µs | speedup | +|---|---|---|---| +| 1 | 134 | 1994 | **0.1×** | +| 8 | 503 | 1996 | **0.3×** | +| 64 | 3,132 | 1991 | 1.6× | +| 512 | 24,287 | 1997 | 12.2× | +| 4,096 | 205,446 | 1998 | 102.8× | +| 12,928 | 665,694 | 1998 | **333×** | + +**Bisect is FLAT at ~1995 µs across four orders of magnitude of support width** — the design +property, confirmed empirically, and the strongest internal-consistency signal in the table. +Legacy scales linearly, exactly as the O(vocab × kept) reading predicted. + +**But the crossover is at support ≈ 64, and real LLM distributions are peaked** ⇒ typical `kept` is +in the tens ⇒ **bisect is 3–15× SLOWER in the everyday regime.** Swapping unconditionally trades a +pathological cliff for a permanent regression. + +### ✅ THE ANSWER IS A HYBRID, NOT A WINNER (D21 — don't rank the novel system down) +Run **legacy's exact iterative argmax while `kept` stays small; fall back to bisection when it +exceeds a fixed iteration budget.** Peaked case stays at legacy speed; diffuse case is bounded at +**~2 ms instead of 663 ms**. And because the fallback is **exact rather than a truncation**, it +still introduces **no cap and no narrowing** — D4 intact. **Designed, not yet built.** +`CudaSampler::sample` stays on the legacy launcher meanwhile. + +### ✅ A2 tie boundary — measured, matches the prediction that was refused as a claim +Exactly-uniform support: legacy **15/16** reachable, bisect **16/16**, **TV = 0.0612** against the +analytic **0.0625**. It errs **wider** — the safe direction. The harness printed `MISMATCH` only +because that run used the old single-threshold bound: **the test artifact the chain had identified +and split out in advance precisely so it would not be mistaken for a sampler defect.** It isn't one. + +**SCOPE — proven with counts from the real V4 B=1 trace** (`/root/budget-chain/nsys/v4b1.sqlite`, +**4,913,454 kernel launches across 109 distinct kernels**): `fused_top_p*` **0**, `fused_argmax*` +**0**, `arc_sampler*` **0**, anything matching `sampl` **0**. Statically, zero references to +`CudaSampler`/`fused_top_p_bf16`/`launch_fused_argmax_bf16` outside `arc-cuda-graph`. +**Serving samples through `mistralrs-core/src/sampler.rs`, a separate implementation** — +consistent with the 517,120 B D2H per step (= 129,280 vocab × 4 B). The 21,796 `topk_kernel` +launches are the **MoE router** (40/step × 40 layers), not token selection. +⇒ **Confined to the gated-off autonomous path. NOT a stop-everything.** It *was* live for the +autonomous chain's own work, which is why it mattered. + +**FIXED:** `arc_launch_sampler_bf16` with device-resident `rng_state` now runs in the graph body; +`fused_top_p_bf16`/`fused_argmax_bf16` are out of it; the by-value `rng_offset` field is **deleted** +rather than left to bake. `top_k` added to `AutonomousDecodeConfig`; the single `normal.rs` +construction site maps mistral.rs's `<=0` to the kernel's `-1`. Penalties stay an in-place pre-pass +and are zeroed in config, because passing `freq_counts=None` with non-zero penalties **drops them +silently** (`sampling_kernel.cu:167`). `cargo check --features cuda` EXIT 0 for `arc-cuda-graph` +and `mistralrs-core`, with an injected-error control proving the check actually compiles them. + +**NEW code-read risk for whoever takes the sampler forward:** phase 4 does a **full O(vocab) argmax +scan per kept token** with `cap = vocab` when `top_k <= 0` (`sampling_kernel.cu:279`), and +mistral.rs's default maps to `top_k = -1`. Peaked distributions stay cheap; **a flat distribution +at 129K vocab is a cliff.** + +**AUTO_FREE / DEVICE_LAUNCH resolved explicitly:** the conditional path instantiates via flag-free +`cuGraphInstantiate_v2` with the reason in-comment (`autonomous.rs:518-522`); `graph.rs:276` keeps +`AUTO_FREE_ON_LAUNCH` for the *non-conditional* capture-once runner. **Those two instantiate calls +must never be merged** — the header forbids combining that flag with `DEVICE_LAUNCH`. + +**tok/s on V4 is NOT measured and must not be estimated.** Work is on `wt/autonomous-decode` in +`/root/arc-wt/autonomous` (6 files, +357/−95); probes at `/root/arc-wt/probe/`. + +**RETRACTED:** the five clashing extern declarations in `gemv_ffi.rs` are **NOT** the same class as +the `cuGraphAddNode` bug. `CUstream = *mut CUstream_st` (cudarc-0.19.4 `driver/sys:1759`), so all +five are *pointer-type* disagreements with **identical arity and identical register layout** — +cosmetic, worth a cleanup pass, **not a corruption risk**. `cuGraphAddNode` was different arity +*plus* a transposed out-pointer. Filing them together would have been a false alarm. + +--- + +## Runner-up by GPU time — the keystone GEMM (ArcKernels/Trellis, the moat) + +`qtip_gather_gemv_warp_kernel`: **9.85 ms/step, 129 calls, ~77.3 µs each — 29% of all kernel time.** + +> ⚠️ **CORRECTION: it is 29× above its own bandwidth bound, NOT 15×.** The 15× figure divided this +> kernel's ms/step by the *whole model's* 0.677 ms floor. **Its own floor is 129 × 2.62 µs = +> 0.338 ms.** Do not requote 15×. + +> ⚠️ **"near its instruction floor for this format" is now QUANTIFIED, not accepted.** The Gaussian +> codeword accounts for **13.0 of the 35.5 inst/weight**, and predicted time from inst/weight tracks +> measured time at **~70% issue efficiency in every arm** ⇒ **this kernel is ISSUE-BOUND and +> inst/weight is the control variable.** That is a transfer function, not a ceiling. + +### Measured on `arc-v4-stack`, SHA `05af600e7`, interleaved A-B-A-B ×4, **drift 0.16%** + +| | inst/weight | µs/call (gate/up) | ms/step | × its own bw bound | +|---|---|---|---|---| +| **ships today** | 35.52 | 77.31 | **9.85** | 29.2× | +| **+ bit-identical fixes** | 28.91 | 66.19 | **8.59** | 25.4× | +| + `sum2` codebook (**FORMAT CHANGE — see caveat**) | 15.85 | 33.47 | **~4.6** | ~13.6× | + +> ⚠️ **THE `sum2` ROW IS TIMING-VALID, NOT A QUALITY CLAIM. Quote it as "2.27× available, +> contingent on a re-bake and Jish's call" — NEVER as "2.27× measured end-to-end".** +> In that arm the harness decodes **Gaussian-baked bytes against the MCG codebook**, so the outputs +> are **numerically meaningless**. Control flow and instruction mix are identical, which is exactly +> why µs and inst/weight transfer — and exactly why accuracy does not. +> The only quality evidence that exists is an in-tree fixture +> (`computed_codebook_quality_is_neutral_as_shipped`: **+0.00017 cos, +0.37% weight NMSE**, against +> `split` at −0.00174 / +3.73%). **A serving-side quality number requires the re-bake +> (`ARC_QTIP_CODEBOOK=mcg`) and has NOT been run.** +> +> **Containment confirmed by construction, not by intent:** `git status --porcelain` shows exactly +> one changed file and it is a `.cu` — `mistralrs-quant/kernels/qtip/qtip_gather_gemv.cu`. No Rust +> file touched, so `QtipCodebook::DEFAULT = Gaussian` (`mod.rs:564`) is untouched and `cuda_mult()` +> still returns `CB_MULT_GAUSSIAN_LUT` = 0 for every shipped artifact. **Nothing in this work can +> flip the format**; the kernel merely keeps supporting both arms through the **pre-existing** +> `cb_mult` ABI. Every `sum2` number came from forcing `cb_mult` via a `GG_CB_MULT` env var that +> exists **only in `bench_gg.cu`** and never reached the model path. Both e2e A/B arms are Gaussian. + +**Both projections moved** — gate/up **77.31 → 66.19 µs/call (−14.4%)**, down **74.54 → 67.48 +(−9.5%)**, total **9.85 → 8.59 ms/step (−12.8%)**. Static SASS and ncu agree to **0.3%** on the +codeword cost. + +**Parity: 0 differing bytes** of 24,576 and 49,152, identical md5/FNV/sum — **and the check is +proven able to fail**: a one-token mutation produced 21,494 and 43,150 differing bytes (D33). +**`GROUP=8` is faster but changes summation order ⇒ deliberately EXCLUDED** — speed that costs +bit-identity is a different decision and doesn't ride along in a parity-preserving PR. + +**SHIPPED: PR #129**, one file (`mistralrs-quant/kernels/qtip/qtip_gather_gemv.cu`, +158/−100), +checkout restored to `master` with all seven other chains' modified and untracked work intact. + +### What the SASS showed that nobody had named +Loop fully unrolled to GROUP×ROWS = 8 symbol decodes; static count 561 = **35.06 inst/weight vs +ncu's measured 35.52**, so the static histogram is a faithful dynamic profile. +1. **16 × `LD.E.U8` — *generic* byte loads, not `LDS`.** `stage_packed` was a runtime `int`, so + `rp[r]` was a runtime-selected shared-or-global pointer and the compiler had to emit generic + loads plus a 64-bit address chain (`LEA.HI.X.SX32`×12, `IMAD.SHL.U32`×12, `IMAD.X`×6, + `IADD3`×23). **The shared-memory staging optimisation was being defeated at the ISA level** — + the code was written, and silently degraded by one runtime flag. +2. **Half of those loads are the warm-up replay** — 4 symbols replayed per 4 decoded, per row, + producing **no weights at all**. + +**The three bit-identical fixes:** `STAGED` as a *template* parameter (→ real `LDS`, 32-bit +addressing); the L=16 state **is** the last 4 nibbles of the stream, so the replay collapses to one +aligned `LDS.U16` window plus a nibble reversal; and the GROUP activations become one `LDG.E.128`. + +**🔑 The `sum2` codebook row is a FORMAT change — it halves inst/weight and needs Jish's call plus +a quality number and a re-bake.** D17: the byte format is the moat, which is exactly why changing +it is a decision, not an optimisation. + +### Trap avoided, worth keeping +The chain's first standalone harness ran **129.5 µs vs the model's 78.4** — it had not matched the +crate's nvcc flags. The crate builds with **`--use_fast_math`, which had already collapsed +`logf`/`sqrtf`/`sincosf` into single `MUFU` ops.** Matching the flags reproduced the in-model number +to **1.4%** (77.27 vs 78.35, regs 40 = trace). **Without that check it would have "optimised" +transcendentals that were not there** — the same shape as the earlier +12–20%-predicted/+3.3%- +delivered miss. **Reproduce the in-model number in your harness before optimising anything.** + +--- + +## SUSPECT SCORECARD (three hypotheses, all resolved) + +| # | Hypothesis | Verdict | +|---|---|---| +| 1 | mask `arange`/`zeros`/`full` per layer per step (CLAUDE.md pitfall #5) | **INNOCENT** — those are H2D and non-blocking. Whole `sdpa` span incl. all mask construction = 165.01 µs/call = 7.10 ms/step = 10.6%. The real mechanism is the FP8 CPU cast, and the naive fix for it is *negative*. | +| 2 | CUDA graph replay dead | **CONFIRMED**, root cause named: legacy default stream | +| 3 | unfused decode | **CONFIRMED but bounded at 10.6%** | + +## Corrections folded in + +- The earlier "**~215 µs unattributed** inside `mla_attn`" was **never unattributed** — the child + list was incomplete. True residual **21.90 µs/call = 3.4%**. The four missing spans: + `kv_fp8_quant` 73.49 + `kv_fp8_dequant` 50.09 + `compressed_kv_build` 49.97 + + `compressor_advance` 27.37 = **200.92**, + 21.90 = **222.8 µs**. Nothing is hiding. +- The earlier brief was internally inconsistent by 2.16× (623.32 µs/call × 43 = 26.8 ms, not the + stated 58.0 ms/step). Resolved in favour of the per-call figure: 643.94 × 43 = **27.69 ms**. +- **RETRACTED: "GPU is 2–7% busy"** — that sampled `nvidia-smi` during the *profiled* run. Clean + run reads **69–82%**. +- **RETRACTED: "ARC_PROFILE costs 4×"** — inferred from log timestamps; directly measured at 5–8%. + +## Surfaced, not shipped + +`v4_nan_dbg(&y, &format!("L{li}..."))` builds the `String` **unconditionally, before** the env gate +short-circuits — ~500 wasted allocations + env lookups per token. Assigned to the ArcGraph chain. diff --git a/memory/mission/BUDGET_V4_PREFILL.md b/memory/mission/BUDGET_V4_PREFILL.md new file mode 100644 index 000000000..b5ebbfaca --- /dev/null +++ b/memory/mission/BUDGET_V4_PREFILL.md @@ -0,0 +1,329 @@ +# THE V4 PREFILL BUDGET — a DIFFERENT disease from decode + +**Read this before proposing any prefill or TTFT work. MEASURED, not estimated.** +Companion to `BUDGET_V4_B1.md` (decode). **The two have opposite causes and opposite fixes.** + +Provenance: `arc-v4-stack` (H200), SHA `05af600e`, binary md5 `41d0bc20…`, exclusive card, both locks +held, exclusivity asserted per leg with a selftest proving the assertion can fail. +`arc_profiler` device selftest **PASS 43.0×**, overhead **0.034%**, reconciliation **0 violations / +0 misnested / 0 unresolved**, self-times sum to **100.0%**. Drift control: N=128 re-measured 14 slots +/ 17 min later at **1578 vs 1574 ms = 0.26%**. Session 8, 2026-08-18. + +--- + +## 🔴 THE HEADLINE: PREFILL IS 99.5% GPU-BUSY. DECODE IS 49% GPU-IDLE. + +| | decode (B=1) | **prefill** | +|---|---|---| +| GPU idle waiting on host | **49%** | **~1%** | +| kernel launches | 9,131/token | **15,744/step** (N=128) — only **1.78×** | +| but wall clock | 66.68 ms | **1,531.8 ms — 21×** | +| kernel mean duration | 3.81 µs | **92.9 µs — 24× longer** | +| `cuLaunchKernel` host cost | the bottleneck | **50.7 ms of 1,545 = 3.3%** | + +> **PREFILL IS NOT LAUNCH-BOUND. CUDA graphs, launch reduction and the arena buy prefill NOTHING.** +> Everything session 8 built for decode is irrelevant here. This needs a different fix. + +## Per-step attribution, N=128, B=1 (self-times sum to 100.0%) + +| component | calls/step | ms/step | % | +|---|---|---|---| +| **`experts.down/gate/up_proj` — QTIP MoE gather** | **129** | **1092.4** | **71.3%** | +| `q_proj` + `o_proj` (MLA) | 86 | 184.4 | 12.0% | +| `moe.shared_expert` | 43 | 67.4 | 4.4% | +| `sdpa` / `mhc_attn_pre` / `kv_fp8_quant` / `mhc_ffn_pre` | 43 ea | 100.8 | 6.6% | +| everything else | — | 86.8 | 5.7% | +| **step total** | | **1531.8** | **100%** | + +**Two independent instruments agree to 0.2%:** profiler says the three expert-gather spans total +**1092.4 ms**; nsys says `qtip_gather_gemv_warp_kernel` is **1090.5 ms across 129 launches**. +That is **8.45 ms per launch** reading **1.61 GB** of packed weights = **190 GB/s ≈ 4% of peak HBM**, +independently corroborated by `nvidia-smi` at **2% memory-controller, 98% SM, 450 W**. +**The SMs are pinned serially decoding the trellis, not moving weights.** + +## Counts per prefill step — with a decode control that reproduces the recorded budget first + +| | N=128 (fused GEMV) | N=1024 (dequant fallback) | decode control | +|---|---|---|---| +| kernel launches | 15,744 | **76,626** | 8,830 *(recorded: 9,131)* | +| kernel mean / median | 92.9 µs / 1.18 µs | — | 3.81 µs / 1.18 µs | +| `cuMemAllocAsync` | 14,008 | **137,976** | 10,992 *(recorded: 11,436)* | +| blocking `cuMemcpyDtoHAsync` | 44 | **176** | 43.8 *(recorded: 44)* | +| `cuMemcpyHtoDAsync` | 4,169 | **45,274** | 2,577 | +| **ALL `*Synchronize*`** | **0.00** | **0.00** | 0.00 | +| dominant kernel | `qtip_gather_gemv_warp` **74.6%** | `qtip_dequantize_v2_k4_l16` **68.1%**, ~20,307 launches | gather_gemv 29.7% | + +*(The control reproducing the recorded decode numbers before any prefill number was trusted is what +makes this table credible — the instrument was validated on a known answer.)* + +--- + +## 🎯 NAMED CAUSE: ArcQuant / QTIP **LUT rung** (`QtipLayer`) — what `qtip2-*.uqff` actually ships + +**`gather_forward` has exactly two prefill paths and NEITHER IS A BATCHED GEMM:** + +1. **< 683 tokens — fused GEMV**, one (token,slot) pair per `grid.y`, **no dedup**: + 768 pairs against ~244 distinct experts at N=128 ⇒ **3.15× redundant weight reads.** +2. **≥ 683 tokens — a host-synced, per-distinct-expert dequantize-materialize loop** + (`mistralrs-quant/src/qtip/mod.rs:2887`), with `Tensor::from_vec` in the hot loop and `out_flat` + reallocated per expert. **Measured launch signature matches the loop body exactly:** + `qtip_dequantize` **81,228** · `is_u32_bf16` (index_select) **81,641** · `ia_u32_bf16` + (index_add) **81,073** — all equal. + +**⛔ `ARC_QTIP_ONDEVICE_MOE_MAX_TOKENS` CANNOT FIX THIS.** Throughput is the same either side of the +crossover — **86.4 vs 84.7 tok/s** — so the policy is choosing between **two equally bad paths**. + +## ✅✅✅✅ CUMULATIVE TTFT — **5.93 → 1.80 s at N=512 (3.3×)**, **16.87 → 5.90 s at N=2048 (2.9×)** + +Two shipped PRs, both measured on an exclusive card with fitted estimators. + +### PR #138 — the FP8 scalar GEMM (`perf/fp8-prefill-cublas`), N=2048 MEASURED +| slot | arm | TTFT | fitted | +|---|---|---|---| +| 1 / 3 | A | 9590.66 / 9627.47 ms | **c = −3720.69 ms** | +| 2 / 4 | B | **5897.88 / 5896.69 ms** | **1.632× — 9.61 → 5.90 s** | + +**4.694 → 2.880 ms/prompt-token.** ⇒ **It beat its own projection (1.61×) by 1.7% — the projection +was PESSIMISTIC, not optimistic.** +🔑 **Slot drift was +8.9 ms/slot here versus −0.27 at N=512 — exactly the bias the fit exists to +remove. A raw A−B would have misread it.** Arm A independently reproduces the recorded post-#133 +baseline (9590.66 / 9627.47 vs 9655.92) **within 0.7%.** + +**🔑 SHIPPED DEFAULT IS *ON*, NOT OPT-IN — and the reasoning is the session's whole lesson:** +> *"Leaving it behind an unset env var would have made this the TENTH wired-but-unreached — the +> precise disease being fixed."* + +`DEFAULT_MIN_M = 512` is **the measured point, not an interpolation**: at N=512 `m_rows == 512`, so +**the shipped default is code-identical to measured arm B.** Decode untouched (own `fp8_gemv_warp` +at M ≤ 4). **The crossover below 512 is unmeasured and flagged as a follow-up.** + +### THE ROOT CAUSE — a NINTH wired-but-unreached +`fp8_matmul_tiled` is a **scalar CUDA-core GEMM** (`acc += s_input[ty][k]*s_weight[tx][k]`, **no +tensor-core instruction anywhere**) at **4.4–4.8 TFLOPS**, serving **both** MLA q/o and the shared +expert — **39.4% of the step at N=2048, exceeding the gather** while the recorded attribution still +said the gather was 71%. **A dequantize + cuBLASLt path already existed in-tree, unreachable: +`forward()` consulted device and kernels-compiled but NEVER token count.** Span diff proves +mechanism: q+o **706.1 → 26.2 ms (27×)**, shared **253.2 → 11.9 (21×)**, **gather and sdpa unchanged +within 1%.** + +### 🔑 `sum2` REPRICED ON MEASURED NUMBERS — its value is CONTINGENT ON ORDERING +| | gather share | `sum2` worth | +|---|---|---| +| before #138 | 36–44% | **1.32× / 1.26×** | +| **after #138** | **66.6% (N=512) / 59.9% (N=2048)** | **1.59× / 1.50×** | + +**The 2.27× headline is worth far less than it reads, and how much depends on what lands first.** +Supporting: the gather's mma floor is **27 ms of 3494 (0.8%)**, so it is **~99% trellis decode** and +the inst/weight transfer function applies nearly undiluted. + +### MEASURED MAP TO JISH'S 1-SECOND BAR +| N | now | with `sum2` | +|---|---|---| +| 512 | **1.8× short** | **1.1× — reachable** | +| 2048 | **5.9× short** | **3.9× short** | +**N=2048 needs a third lever, and there is no single hot spot left** — every remaining term sits at +4.4–7.6 TFLOPS, 130–226× below peak; the whole step is 5.6 TFLOPS = **71.7× the 400 TFLOPS +reference, independently reproducing the 72×.** It becomes a broad efficiency problem. + +**Filed (issues disabled ⇒ BACKLOG):** `weighted_sum` at 262 ms/N=2048, ~145× above its own +bandwidth floor — **with the honest caveat that no kernel was named for that span (the nsys leg +SIGSEGV'd), so 145× is a ratio against a DERIVED floor, not a measured achievable rate.** + +## ✅ (superseded) END-TO-END — TTFT ROUGHLY HALVES (commit `52aec71`, PR #133) + +| N | ms/token before | **after** | **gain** | **TTFT** | +|---|---|---|---|---| +| 128 | 12.356 | **8.091** | **1.53×** | **1.58 → 1.04 s** | +| 512 | 11.585 / 11.581 | **5.329 / 5.312** | **2.18×** | **5.93 → 2.73 s** | +| 1024 | 11.811 | **4.901** | **2.41×** | **12.09 → 5.02 s** | +| 2048 | 8.236 | **4.694** | **1.75×** | **16.87 → 9.61 s** | + +**Baselines reproduce the independently-established table** (12.36 / 11.58 / 11.81 / 8.24 vs +11.98 / 11.48 / 11.81 / 8.23) — **the harness validating itself against a run it did not produce.** + +### 🔴 THE FIRST RE-RUN FALSIFIED THE PLAN: wiring the kernel in was NOT enough +| N | before | after wiring | gain | +|---|---|---|---| +| 128 | 12.352 | 12.385 | **1.00×** | +| 512 | 11.582 | 11.588 | **1.00×** | + +**Not noise-masked: five interleaved replicates of IDENTICAL code at N=512 gave 11.578–11.589 — a +0.011 ms/token floor (±0.05%).** ⇒ **the replicate design turned "looks like nothing" into "provably +nothing".** Without it this reads as an inconclusive run and someone re-runs it hoping. + +**Cause — an EIGHTH wired-but-unreached, and a general lesson.** `lut_fused_gather_preferred`'s 16× +traffic ratio answers *"GEMV or dequantize-materialize?"* — **two paths that both scale with +(token × top_k) PAIRS.** The grouped GEMM is a **third path scaling with DISTINCT EXPERTS**, so that +boundary **no longer describes the decision at all.** Below ~683 tokens (683·6 ≈ 4096 = 16·256) the +call returned at the on-device GEMV and **never reached the new kernel.** +> 🔑 **A dispatch policy that chooses between two options cannot accommodate a third without being +> rewritten. Adding a code path is not the same as making it reachable** — and the component win +> was real *and* unreachable exactly where most prefill lives. + +Decode regime stays fused unconditionally (**RUN-161 floor, not a performance choice**); +`ARC_QTIP_ONDEVICE_MOE_MAX_TOKENS` still overrides so a harness can pin the GEMV arm. + +### ⚠️ TWO CORRECTIONS THE CHAIN MADE AGAINST ITSELF +- **The ~3.5 ms/token projection was TOO OPTIMISTIC** — measured **4.69–8.09**. Arithmetic + projections from a component win overstate; the dispatch and the other 29% of the step both bite. +- **The PR title advertised the 6.6× COMPONENT number** — retitled to the **end-to-end range**. +- **MEMCTRL still FALLS, not rises** (0.5→0.2 at N=128, 6.1→1.8 at N=2048), consistent throughout: + **decode-ALU-bound, not HBM-bound.** The original criterion still says this is not bandwidth work. + +### 🐛 The runner's own failure, and the rule it produced +The e2e runner **did not inherit `LD_LIBRARY_PATH=/usr/local/cuda/compat`** — *the exact trap the +same chain had filed in BACKLOG one script earlier.* Every process died with +`CUDA_ERROR_UNSUPPORTED_PTX_VERSION` **at line 83 of an 18 KB log, well past where it was tailing**, +producing a complete 14-row table of `RAW`. +⇒ **"Don't gate on rc" MUST be paired with a positive engagement assertion.** The runner now requires +each slot to yield a parseable `Prefill (N tokens) … ms (TTFT)` line and **exits 2 on the first slot +that doesn't.** Recorded in BACKLOG. + +## ✅ COMPONENT-LEVEL (superseded by the e2e above, kept for provenance) + +One MoE expert matmul, E=256, top-6, **exclusive idle card** (`apps=[]`, 143 GB free, builders=0). +Pairs are V4's two shapes 4096→2048 / 2048→4096. Engagement counters written to files. + +| N | shipped path | shipped ms | **grouped ms** | **gain** | SM | MEMCTRL | +|---|---|---|---|---|---|---| +| 128 | gemv | 8.38 / 8.62 | 7.46 / 7.14 | **1.15×** | 100→97% | 5.0→**2.0%** | +| 512 | gemv | 33.47 / 34.26 | 8.53 / 8.51 | **3.98×** | — | — | +| 1024 | dequant | 127.87 / 75.02 | 15.41 / 15.37 | **6.59×** | — | — | +| 2048 | dequant | 133.68 / 84.62 | 26.69 / 26.78 | **4.08×** | 100→100% | 6.8→**1.2%** | + +### 🔑 THE MEMCTRL ANSWER IS *NOT* THE HOPED-FOR ONE — and that is the finding +**MEMCTRL FELL (5.0→2.0% and 5.0→1.2%); SM stayed 97–100%.** By the stated criterion this is +**NOT** a move to bandwidth-bound. **What did happen: traffic per unit work collapsed ~21× at +N=2048** (1.2%×26.7 ms vs 5.0%×133.8 ms) — so the **3.15× redundant-read defect is genuinely +removed**, but the kernel is now **trellis-decode-ALU-bound, not HBM-bound.** +⇒ **The ~14 ms weight-traffic floor is NOT the operative limit.** At N=128 the grouped kernel costs +**7.45 ms/layer/shape — ~100× above that floor.** **Scope, not verdict: the next lever is decode ALU** +(the V=2 pairing already halves state reconstructions vs the qtip2b rung), **not weight traffic.** + +### ✅ NO RE-BAKE NEEDED, ZERO COST +The kernel **templates on the codebook** exactly as every other K=4/V=2 launcher does (`cb_mult==0` → +stored-LUT gather, nonzero → computed MCG), so **it runs on the shipped `qtip2-*.uqff` unchanged**. +*(ISQ-to-`qtip2b` was never the route — it would need Viterbi over the full 284 B source.)* + +### Quality: bit-identity was never available; cosine measured instead +mma **f32** accumulation vs the fallback's **cuBLAS BF16** matmul is a different summation order. +**Measured: cosine 1.00000 against the dequantize-materialize reference at all 8 (shape, N) points**, +same decoded weights. + +### 🔑 Why this was always reachable — the load-bearing insight +The LUT rung's `state = ((state<<4)|sym) & 0xFFFF` **looks sequential but is just a 16-bit nibble +window**: `state(t) = nibble_reverse_16(bits[4t-12 .. 4t+3])` — **random-access**, exactly like the +qtip2b rung's pair-reversal. **That is why a grouped GEMM was reachable here all along.** + +### ⚠️ NOT ESTABLISHED +**No end-to-end V4 ms/token** — the 138 GB model was never loaded; the card was contended by three +chains all session and the chain chose the component measurement over queueing. **So the +11.98 → ~3.5 ms/token projection remains ARITHMETIC.** Also unestablished: the 71.3% component share +at N≠128, and there is no 6-arm slot-fitted design (one exclusive block instead). + +**Code:** `kernels/qtip/qtip_grouped_gemm_lut.cu` (new, ~330 lines) · `grouped_gemm_lut_cuda` in +`src/qtip/cuda_ops.rs` · FFI in `ffi.rs` · wired into `QtipLayer::gather_forward` · harness +`examples/qtip_lut_grouped_bench.rs`. Worktree `/root/arc-wt/lutgemm`. + +### 🚨 SEVENTH "GREEN THAT COULD NOT CONTAIN THE ANSWER" — and it is live for anyone adding a kernel +**`mistralrs-quant/build.rs` emits `rerun-if-changed` only for the sage/metal paths, so a NEW `.cu` +in the `kernels/*/*.cu` glob does not trigger a build.rs rerun.** A build reported **`BUILD_RC=0` +while the kernel was never compiled.** **Caught only because the "Compiling 33 of 33 kernels" count +did not change** (fix: touch `build.rs` → "1 of 34"). +**Operational:** driver 580.173.02 is **CUDA 13.0**, toolkit is **13.1** ⇒ first GPU run died +`CUDA_ERROR_UNSUPPORTED_PTX_VERSION`; needs **`LD_LIBRARY_PATH=/usr/local/cuda/compat`**. +**Also: `qtip_grouped_curve` Viterbi-bakes E=256 single-threaded for HOURS while holding the +exclusive lock** — use a random-packed-bytes fixture instead; **decode cost is content-independent.** + +## 🎯 THE ORIGINAL TARGET (superseded above): ~14 ms WEIGHT-TRAFFIC FLOOR vs 1,092 ms TODAY + +**A grouped GEMM that reads each distinct expert exactly once has a ~14 ms weight-traffic floor at +N=128.** Against the measured **1,092.4 ms**, that is **roughly two orders of magnitude of headroom +on this one component** — and it is **the same component in both regimes**, so one fix covers the +fused-GEMV side and the dequantize-fallback side together. + +**Arithmetic projection, NOT a measurement** (state it as such): step 1531.8 − 1092.4 + 14 ≈ +**453 ms**, i.e. **11.98 → ~3.5 ms per prompt token ≈ 3.4× on prefill**. Still above the compute +bound afterwards — the remaining 12% MLA projections and 4.4% shared expert become the next terms. + +## 🔑 THE KERNEL ALREADY EXISTS — WE SHIP THE RUNG WITHOUT IT + +**The sibling bitshift rung (`Qtip2bLayer`, `--isq qtip2b`) HAS the amortizing +`grouped_gemm_2b_cuda`**, and **V4's shapes satisfy all three of its gates**: 4096 and 2048 both +divisible by `GROUPED_TILE_K` = 64, BF16, uniform 2-bit. +⇒ **The sixth instance this session of correct code the shipped path cannot reach.** + +## Scaling, and where "24 s TTFT" comes from + +| N | ms/prefill | **ms/token** | MoE path | SM | MEMCTRL | W | vs 400 TFLOPS | +|---|---|---|---|---|---|---|---| +| 128 | 1534 | 11.98 | fused GEMV | 98% | 2% | 440 | 112× | +| 512 | 5876 | 11.48 | fused GEMV | 98% | 2% | 469 | 177× | +| 1024 | 12092 | 11.81 | dequant fallback | 100% | 5% | 243 | **182×** | +| 2048 | 16855 | 8.23 | dequant fallback | 100% | 7% | 259 | 127× | + +> **TTFT 24 s = ~2,000 prompt tokens × 11.8 ms on an IDLE EXCLUSIVE H200.** The recorded 43.2 s step +> = ~3,700 tokens at the same rate. **Compute, not starvation, not queueing.** + +**Four forward predictions, recorded BEFORE the run, all landed:** crossover at 683 ⇒ fallback +warning absent at 512, present at 1024 ✓ · memory-controller 2%→5%→7% ✓ · fallback D2H ≈129 router +syncs (measured 176 = 129 + 43 KV + logits) ✓ · fallback H2D ≈33k `Tensor::from_vec` uploads +(measured 45,274) ✓. + +## Qwen vs V4 — the named difference +**Qwen3-32B is DENSE.** `qwen3.rs` has **zero** `gather_forward`/`MoEExperts`, so every prefill +projection is one batched cuBLAS GEMM. **V4 routes every prompt token through QTIP expert gather.** + +## 🔴🔴 PREFILL EXCLUDES DECODE — the likely aggregate ceiling, and it COMPOSES with the above + +**`mistralrs-core/src/paged_attention/scheduler.rs:216` — the paged scheduler returns EITHER a +prefill batch OR a decode batch, never both.** The prompt-admission loop **returns early at +`if !scheduled.is_empty()`, before the decode leg ever runs.** + +⇒ **Every arriving request stalls ALL running decodes for a full prefill step.** + +**Combine with the two facts already on this page and the magnitude is severe:** +- **prefill is NOT chunked** (`get_prompt_input` → one `make_prompt_chunk`; `prompt_chunksize` is a + stale comment with no wired field), so a prefill step is **the whole prompt, indivisibly** +- a prefill step at N=512 measures **5,876 ms**; a decode step is **66.68 ms** + +⇒ **One arriving 512-token request freezes every running user for ~5.9 seconds — ~88 decode steps.** +At B=64 that is a throughput catastrophe, and it is invisible in any single-user measurement. + +### 🔑 THIS MULTIPLIES THE VALUE OF THE GROUPED-GEMM FIX +The kernel is worth **4–6.6× on the prefill component** for the *arriving* user. But because that +prefill step **blocks every other user**, the same fix **shortens the stall the whole batch eats.** +**Its aggregate value is larger than its single-user value** — and no single-user benchmark can show +that. *(Chunked prefill would decouple them entirely — see the "not chunked" finding above.)* + +**Being tested, not assumed:** two diagnostic legs at B=64 — `decode_heavy` (128-token prompts) vs +`prefill_heavy` (spread prompts, 32 gen tokens). **If aggregate jumps on `decode_heavy`, prefill +interference is named rather than guessed.** + +## 🐛 `logit_bias` IS A DEAD WIRE — the API accepts it and silently ignores it +`mistralrs-server-core/src/{completions,chat_completion}.rs` populate `SamplingParams::logits_bias`; +**nothing in `mistralrs-core` ever reads it** — the only references are the field declaration and its +defaults. **A customer can send `logit_bias` on an OpenAI-compatible endpoint and it does nothing, +with no error.** Sixth wired-but-dead of the session, and the first that is **user-facing**. + +## Ruled out +- **Prefill is NOT chunked** — `get_prompt_input` → one `make_prompt_chunk`; `prompt_chunksize` is a + **stale comment with no wired field.** +- **The compressor is fully vectorised** — not a contributor. +- **Batching cannot rescue it:** cost is linear in (token × top_k) pairs with **no cross-token + amortisation**, and `cache_bucket_len()` falls back to `len()-1` for fresh sequences, so differing + prompt lengths **serialise one bucket per prefill step**. + +## ⚠️ NOT ESTABLISHED — do not quote as fact +- **No qtip2b A/B. The fix is code-identified and UNMEASURED — treat the headroom as scope, not a + promise.** +- **The Qwen 565–792 tok/s figures are not certifiable on this box** — FACTS records Qwen3-32B on a + **B200**. The "30× worse" ratio is withdrawn; **nothing above rests on it.** +- **N=1024 per-step counts are ±10%** — derived by subtracting measured decode cost from trace + totals because segmentation failed there (**196 fragments — itself a symptom of the per-call + syncs**). The N=128 numbers come from two byte-identical segments and carry no such caveat. + +## 🐛 Separate defect, needs its own ticket +**The bench binary corrupts its heap at teardown** — SIGABRT *"corrupted double-linked list"* ×3, +SIGSEGV ×2 — **after printing results.** ⇒ **Never gate a prefill result on the process exit code;** +two runs were nearly discarded that way, and the traces were recoverable from disk both times. diff --git a/memory/mission/CEILINGS.json b/memory/mission/CEILINGS.json new file mode 100644 index 000000000..ce7535b48 --- /dev/null +++ b/memory/mission/CEILINGS.json @@ -0,0 +1,296 @@ +{ + "_README": "Arc ceilings + quality model. READ BEFORE quoting any speed or quality number, or before concluding anything is 'not achievable'. Every number here is either PHYSICS (a bound reality imposes) or IMPLEMENTATION (a statement about today's code). Confusing the two is the single most repeated mistake in this project's history. Companion files: FACTS.md (measured only), DOCTRINE.md (standing rules), STATUS.md (now).", + "version": "1.0", + "created": "2026-08-15", + "author_note": "Written after Jish corrected Fable for (a) inflating the dense term 4x and understating the b=1 ceiling, (b) presenting uniform-routing lower bounds as ceilings, (c) repeating '15% of peak' as though it were a constraint.", + "MISSION": { + "statement": "Get more real work out of the compute that already exists. Every token Arc serves more efficiently is silicon that does not need to be built, powered, or cooled.", + "why_it_matters": "AI inference is a growing share of datacenter energy draw. Serving the same traffic on a fraction of the hardware is a direct cut to electricity and to the embodied carbon of GPUs that never get manufactured. This is not a side effect of the engineering — it IS the product thesis, expressed in physics.", + "jish_verbatim_2026_08_15": "we're changing the world... turning AI a bit more eco friendly and cutting carbon from the world. We can be anything but pessimistic in this project, we cannot lose, to nobody.", + "jish_verbatim_2026_08_12": "Arc's wedge isn't local llms bro. not single user single gpu huge model, its still datacenters serving llms, but the same 80 nodes of H200s can serve a shit ton of more capacity than they originally did and far cheaper + efficiently than they originally did", + "the_proof_already_in_hand": "Every published DeepSeek-V4-Flash config needs 4x H200. Arc runs it on 1. That is a 4x cut in silicon, power and embodied carbon for the same model, and it is measured, published and verified — not projected." + }, + "ANTI_PESSIMISM_PROTOCOL": { + "core_rule": "A measured number describes TODAY'S CODE. It is never evidence about what is possible. Only physics bounds what is possible.", + "before_writing_the_words_not_achievable": [ + "1. Is this a PHYSICS bound (bytes that must move / bandwidth / arithmetic that must happen)? If not, it is an engineering task, and say so.", + "2. Have I extrapolated past the edge of the sampled data? Compute four more rows of the formula before concluding a trend.", + "3. Am I quoting a % of peak? Someone else's implementation reaching higher is proof ours is not at a limit. EXL3 reaches ~80% for this exact format; 15% is a verdict on our kernel, not on the format.", + "4. Did I check what a reference implementation achieves? If not, I have an opinion, not a finding.", + "5. Is a worst-case assumption hiding in my model (e.g. uniform routing)? Say which direction it biases." + ], + "fables_recorded_pessimism_failures_2026_08_15": [ + { + "claim": "MoE batch economics are structurally modest, ~1.57x", + "why_wrong": "Stopped computing at B=32 because that is where our data stopped. Four more rows show expert coverage saturates and amortization reaches 6x at B=256.", + "caught_by": "Jish" + }, + { + "claim": "b=1 ceiling ~560 tok/s", + "why_wrong": "Guessed the dense term at 7 GB; it is 1.70 GB. Real ceiling is ~1413 tok/s. Jish's 640 was CONSERVATIVE.", + "caught_by": "Jish" + }, + { + "claim": "GEMV at 15% of peak (stated as a constraint)", + "why_wrong": "Our own FACTS records EXL3 at ~80% for this format. 15% is our implementation. The computed codebook measured 3.86-4.03x on the decode GEMV, which lands near 60%.", + "caught_by": "Jish" + }, + { + "claim": "0.2% of roofline", + "why_wrong": "The roofline itself was overstated ~8x by assuming all weights are read every step at small batch.", + "caught_by": "agent AY" + }, + { + "claim": "Paged attention is the keystone unlocking three capabilities", + "why_wrong": "Infeasible for V4 AND insufficient — the prime_for_step bridge does not exist.", + "caught_by": "agent BC" + }, + { + "claim": "The CPU sampler fallback is the prime suspect for the throughput collapse", + "why_wrong": "Probe sends temperature=0 -> greedy path returns before the radix branch. Explains 0%.", + "caught_by": "agent AY" + } + ], + "the_discipline_that_survives": "Optimism about what is POSSIBLE. Ruthlessness about what is MEASURED. Never soften a bad measurement, never accept it as a ceiling. Report the gap AND the physics bound in the same breath, so the gap reads as work to do rather than as a verdict." + }, + "MODEL_CONSTANTS": { + "model": "DeepSeek-V4-Flash", + "params_total": 284000000000.0, + "params_active_per_token": 13000000000.0, + "num_experts": 256, + "experts_per_token": 6, + "_note_top6": "TOP-6, NOT top-8. Fable used top-8 and overstated every amortization row.", + "artifact_bytes": 74190197268, + "bits_per_param": 2.09, + "_note_bpw": "2.09, NOT 1.9. The 1.9 came from dividing by a 68 GB ESTIMATE instead of the real 74.19 GB.", + "bytes_per_param": 0.26125, + "derivation_of_split": "256*P + D = 284e9 ; D + 6*P = 13e9 => P = 1.084e9 params/expert, D = 6.496e9 dense params", + "bytes_per_expert": 283200000, + "bytes_dense": 1697000000, + "sanity_check": "1.697 GB dense + 6 * 0.2832 GB = 3.396 GB = exactly 13e9 active params * 0.26125 B/param. Consistent." + }, + "HARDWARE": { + "H200": { + "hbm_bytes": 141000000000.0, + "bandwidth_bytes_per_s": 4800000000000.0, + "cost_per_hour_usd": 4.85 + }, + "A100_80G": { + "hbm_bytes": 80000000000.0, + "bandwidth_bytes_per_s": 2039000000000.0, + "cost_per_hour_usd": 1.49 + } + }, + "SPEED_MODEL": { + "formula_distinct_experts": "E(B) = 256 * (1 - (1 - 6/256)^B)", + "formula_bytes_per_step": "bytes(B) = 1.697e9 + E(B) * 0.2832e9", + "formula_step_seconds": "t(B) = bytes(B) / 4.8e12", + "formula_per_user_tok_s": "1 / t(B)", + "formula_aggregate_tok_s": "B / t(B)", + "_CRITICAL_CAVEAT": "E(B) assumes UNIFORM RANDOM routing, which MAXIMISES distinct experts for a given B. Any real-world routing concentration wakes FEWER experts, reads FEWER bytes, and lands ABOVE these numbers. THESE ARE LOWER BOUNDS ON THE CEILING, NOT CEILINGS. The real coverage curve has never been measured on this model — measuring it is cheap and unclaimed.", + "table": [ + { + "B": 1, + "experts_woken": 6.0, + "bytes_per_step_GB": 3.4, + "step_ms": 0.71, + "ceiling_per_user_tok_s": 1413, + "ceiling_aggregate_tok_s": 1413 + }, + { + "B": 8, + "experts_woken": 43.9, + "bytes_per_step_GB": 14.13, + "step_ms": 2.94, + "ceiling_per_user_tok_s": 340, + "ceiling_aggregate_tok_s": 2721 + }, + { + "B": 16, + "experts_woken": 80.2, + "bytes_per_step_GB": 24.41, + "step_ms": 5.09, + "ceiling_per_user_tok_s": 197, + "ceiling_aggregate_tok_s": 3144 + }, + { + "B": 32, + "experts_woken": 135.3, + "bytes_per_step_GB": 40.01, + "step_ms": 8.34, + "ceiling_per_user_tok_s": 120, + "ceiling_aggregate_tok_s": 3837 + }, + { + "B": 64, + "experts_woken": 199.1, + "bytes_per_step_GB": 58.08, + "step_ms": 12.1, + "ceiling_per_user_tok_s": 83, + "ceiling_aggregate_tok_s": 5289 + }, + { + "B": 128, + "experts_woken": 243.4, + "bytes_per_step_GB": 70.63, + "step_ms": 14.71, + "ceiling_per_user_tok_s": 68, + "ceiling_aggregate_tok_s": 8701 + }, + { + "B": 256, + "experts_woken": 255.4, + "bytes_per_step_GB": 74.02, + "step_ms": 15.42, + "ceiling_per_user_tok_s": 65, + "ceiling_aggregate_tok_s": 16602 + } + ], + "headline_ceiling": "ONE H200 can in principle serve ~16,600 tok/s aggregate at B=256. That is the fleet number. It is physics, not a projection.", + "PAST_SATURATION_EVERY_EXTRA_USER_IS_NEARLY_FREE": { + "why": "Once E(B) reaches ~256 of 256 experts, bytes/step stops growing (74 GB is the whole model). Step time pins at ~15.4 ms. Aggregate = B / 0.0154 and therefore scales ~LINEARLY with B from there. This is the single most important shape in the whole model and Fable missed it for a day by only tabulating to B=256.", + "B_384": 24900, + "B_512": 33200, + "B_699": 45400, + "_units": "aggregate tok/s ceiling on ONE H200", + "🔴_CORRECTION_2026_08_16": "These rows are BANDWIDTH-ONLY and the tail ignores the FLOP bound. Bandwidth and compute cross at B≈415, so B_699 should read ~27,000, not 45,400. B_512 is near the crossover; treat it as an upper bound. THE 16,600 HEADLINE AT B=256 IS UNAFFECTED — it sits well below the crossover. (wave44-BV)", + "what_gates_it": "MEMORY FOR CONTEXT, not bandwidth. Per-token context bytes decide max B. See KV_AND_XS_LADDER." + }, + "KV_AND_XS_LADDER": { + "_note": "Each row is the per-token context cost and the max concurrent users it allows in ~65 GB at 2048-token context. This ladder, not the kernel, is what sets the fleet number.", + "2026_08_15_morning": { + "bytes_per_token": 424018, + "max_B": 68, + "note": "xs carried raw" + }, + "after_PR59_rolling_xs": { + "bytes_per_token": 108288, + "max_B": 266, + "note": "xs 60.7x cut; attention KV becomes 81% of the budget" + }, + "after_PR63_kv_dedup": { + "bytes_per_token": 64342, + "max_B": 493, + "note": "K and V were stored SEPARATELY at BF16 though V4's fused wkv makes them IDENTICAL (deepseek4.rs:1459 `let v = k.copy()?`). 2048 -> 1026 B/token/layer. Reference stores 584." + }, + "with_FP8_K_storage": { + "bytes_per_token": 45594, + "target_per_layer": 590, + "max_B": 696, + "note": "MEASURED-BY-DESIGN 2026-08-16 (wave43-BU, PR #72), 16/16 CI green. Bit-exact CONFIRMED: act_quant_kv_nope (deepseek4.rs:2264-2306) already E4M3-round-trips K's 448 nope dims in 64-wide blocks, then WIDENS BACK to BF16 at :2304 and stores that. Storing the E4M3 codes instead is bitwise identical.", + "🔴_correction": "BK's 586 target would have been LOSSY. The reference's 7 UE8M0 scale bytes are a DIFFERENT scale. Bit-exactness requires storing the block amax at ACTIVATION precision — free, because the amax IS one of the block's own elements. Hence 590 B/token/layer, not 586, and max B ~696 not 699.", + "how_it_dodges_PagedAttention": "New models/dsv4_kv_fp8.rs stores codes as U8 in the K half and [rope ++ amax] in the V half of the existing KvCache::Normal slot. 'ZERO changes to cache infrastructure' was WRONG — see status. U8 + index_select LUT rather than DType::F8E4M3 because candle has no CUDA F8E4M3 cast. Needs Dsv4AttentionConfig::raw_prefix so decode reconstructs only the 128 reachable columns — without it the dequant writes ~63 GB/step at B≈700 and LOSES more than it saves.", + "🔴_status_2026-08-16_wave49-BZ": "NOT SHIPPED-ON. PR #72 defaulted this ON and was merged without the GPU run its own author asked for; the first V4 forward on a commit containing it returned finish_reason:error on 3/3 prompts with 'dtype mismatch in slice-set, lhs: BF16, rhs: U8' (wave48-BY). NOW OPT-IN: OFF unless ARC_V4_FP8_KV=1. The 45,594 B/token and max_B 696 are what you get WITH the flag, not what master runs. Root cause was NOT the FP8 layout: the engine preallocates a BF16 [1,1,cap,head_dim] buffer as SingleCache.all_data before the first append (engine/add_request.rs:464 -> kv_cache/mod.rs:802 under CacheInstruction::Reset{load_preallocated_cache:true}), and slice_set demands an exact dtype+width match. Fixed in SingleCache::append. THE SAME BUG ALSO BROKE THE DEFAULT wave33 LAYOUT ('shape mismatch on dim 3, 512 <> 1' from the 1-wide V marker), hidden behind the FP8 one because candle checks dtype before shape — so 'just set ARC_V4_FP8_KV=0' would NOT have rescued the box.", + "needs_gpu": "STILL UNRUN, and it is the only thing between opt-in and default-on: index_select over the 256-entry LUT with U32 indices, and the U8 H2D leg, have never executed on CUDA. (The U8 SingleCache itself is now exercised on CPU.) Decisive cheap check: greedy generation with ARC_V4_FP8_KV=1 vs unset must produce IDENTICAL tokens. ~15 min inside any existing V4 rental, no re-bake, no separate rental justified." + } + }, + "jish_640_target": "640 tok/s single-user is ~45% of the 1413 physics ceiling. CONSERVATIVE AND REACHABLE. Fable's earlier '560 ceiling' was arithmetic error.", + "levers_that_BEAT_the_table": { + "MTP_speculative_decode": "✅ WORKS AT EVERY BATCH SIZE as of 2026-08-16 (wave45-BW, PR #73). Supersedes BOTH earlier claims — 'B=1 only' and 'mutually exclusive with PagedAttention' are now WRONG. **JISH WAS RIGHT to reject the exclusivity**: it is an UNIMPLEMENTED BRANCH, not an incompatibility, and it is MOOT for V4 twice over (supports_paged_attention→false; mtp_decode_kit has one impl). **The guard has never fired.** SGLang does exactly this rollback (eagle_info.py:488-490) and at topk==1 — which is what DeepSeek MTP IS — it is PURE TRUNCATION, ZERO KV MOVEMENT. Paged KV is the BETTER substrate. **The real B=1 gate was a CACHE constraint, not MTP**: single_cache.rs:161 `append` does slice_set with ONE offset for the whole batch. Solved by keeping CACHE lengths uniform while TOKEN lengths diverge; drafting groups by uncached-tail u into ≤depth+1 groups; rollback to C + min(u_i + a_i). 🔑 **AND MTP COULD NEVER HAVE HELPED BEFORE THIS**: the old shape ran TWO full target forwards per step, capping the multiplier at (depth+1)/2 — so `--mtp-depth 1` was mathematically incapable of beating no-MTP. Fused ⇒ ×2 at every depth. PROJECTION (labelled, nothing measured): draft overhead 1+depth/43. k=2 ⇒ ×1.91 ⇒ **B=64: 159, B=128: 130 tok/s per user**. k=3 ⇒ ×2.80 ⇒ 232 / 190. **k=2 CLEARS JISH'S 100-AT-ANY-BATCH TARGET ON ONE CARD** — expert parallelism is not required for that goal. Reference's non-simulated V4 acceptance floor is k=2.30 (test_dsv4_pro_mtp.py:261). STILL UNMEASURED. Command in RUNBOOK_8 §S3c, runs inside any existing rental.", + "routing_concentration": "Real routing wakes fewer than E(B) experts. Unmeasured. Direction is favourable.", + "further_compression": "🔴 rANS MEASURED AND REFUTED 2026-08-16 (wave43-BU) — do not reopen without new evidence. On the REAL artifact (qtip2-7.uqff whole, 1.074B symbols): order-0 symbol entropy 3.99987 of 4.0 bits = **0.003% saving**; order-1 0.024%; order-2 0.259%. Per-row framing needed to keep random access costs **+0.78%** — 150-250x the entropy recovered. **rANS makes the artifact BIGGER before any decode cost.** The trellis indices are already ~maximum-entropy, which is what a good quantizer should produce. REAL wins surfaced instead: F32 row_scales are 0.77% of a layer (BF16 => ~0.39% free), and the 'shared' LUT is written PER LAYER (~90 MB duplicated).", + "multi_gpu_EXPERT_PARALLEL": "🔑 MEASURED-BY-DESIGN (wave44-BV, PR #70): shard the 256 experts across N cards; each reads only its slice. **The saturated per-user floor goes 65 (N=1) → 126 (N=2) → 193 (N=4) → 303 (N=8) at B=128.** ⇒ **TWO CARDS CLEAR JISH'S 100 tok/s-per-user TARGET AT ANY BATCH.** Eight are not needed to start. Comms does NOT eat the win on NVLink (222 MB/step at N=4, 4.5-13% of the step) — the cost is LATENCY: 86 collectives/step ≈ 0.86 ms, 3× the bandwidth term. ⚠️ **EP IS NVLINK-ONLY**: on PCIe/InfiniBand comms becomes 62% of the step and the N=8 win collapses 4.5× → 2.6×. `nvidia-smi topo -m` is a hard preflight. Load imbalance is SELF-BOUNDED here (+4.2% at B=128) because a GPU cannot read more experts than it owns — skip EPLB in stage 1, instrument for it. EP + data-parallel attention beats TENSOR parallel, because V4's K/V is one fused MQA head so TP duplicates the KV cache — the binding constraint. Stage 1 = EP=2 on one NVLink node, ~2 h, $15-20.", + "⚠️_Fable_error": "Fable projected 543 tok/s per user at N=8. WRONG — it assumed the DENSE term shards. It is REPLICATED on every card. Correct pre-comms N=8 is 465, post-comms 303.", + "MTP_MEASURED_2026_08_16": "🟢 **FIRST ACCEPTANCE MEASUREMENT IN THIS PROGRAM'S HISTORY** (wave51-CB, H200, qtip2b artifact). **accept_rate ≈ 0.42 aggregate (27 accepted / 64 proposed); tok_per_step ≈ 1.84 (59 committed over 32 steps).** Per-request spread 0.35–0.50. ⇒ **k ≈ 1.84, NOT the k=2 Fable had been projecting from.** At B=128 the per-user CEILING is 68 tok/s ⇒ 1.84×68 = **~125 tok/s/user — still clears Jish's 100 target, with less margin than the 130 previously quoted.** ⚠️ **SAMPLE IS TINY** — 32 steps, 64 proposals. The reference's published non-simulated V4 floor is **2.30**. Our 1.84 could be sample noise, a shorter draft depth, or genuinely lower acceptance on these prompts. **First measurement, not a settled number — re-measure at scale.** ⚠️ And it multiplies a number far below its own bound: measured per-user at B=128 is **1.09 tok/s** today, ~60× under the 68 ceiling. MTP is a multiplier on whatever the base path achieves — fix the base first. Three prior sessions collected EMPTY files here; the mechanism, the counters and the plumbing are now all confirmed working." + } + }, + "MEASURED_VS_CEILING": { + "_source": "wave26-AX, 2026-08-15, 1xH200, in-situ W=32 bake (NOT the published W=256 artifact - state this whenever comparing)", + "rows": [ + { + "B": 1, + "measured_per_user": 15.11, + "ceiling_per_user": 1413, + "gap_x": 93 + }, + { + "B": 32, + "measured_aggregate": 5.07, + "ceiling_aggregate": 3837, + "gap_x": 757 + }, + { + "B": 64, + "measured_aggregate": 8.14, + "ceiling_aggregate": 5289, + "gap_x": 650 + } + ], + "interpretation": "EVERY ONE OF THESE GAPS IS IMPLEMENTATION, NOT PHYSICS. That is the good news: none of it is a law, all of it is code we control.", + "known_causes_and_status": { + "fused_MoE_gather_capped_at_8_tokens": "FIXED PR #56 - boundary now derived (~683 at top-6). Makes the path FLAT per token, NOT amortizing (no dedup). Expect 'stops collapsing', not 'starts scaling'.", + "scheduler_ran_half_the_batch": "FIXED PR #54 - coalescing selector; measured 32 running/32 waiting at B=64 before.", + "no_dedup_across_tokens": "OPEN - one GEMV per (token,expert) pair. The grouped GEMM exists (qtip_grouped_gemm.cu, 5/5 HW parity) but only on the qtip2b rung. Crossover ~B=52 due to a measured 1.76x per-byte deficit.", + "GEMV_at_15pct_of_peak": "OPEN AND UNACCEPTABLE - EXL3 reaches ~80% for this format. Computed codebook (MERGED PR #46) measured 3.86-4.03x on fused decode GEMV at k_in=7168 => ~60% of peak. Needs a bake on the computed codebook to take effect at serve time.", + "per_step_host_overhead": "OPEN - the dominant b=1 term. CUDA-graph capture + fused decode. This is the 'megakernel' work Jish deprioritised; it is what stands between 15 tok/s and 640 at b=1.", + "xs_history_cache": "OPEN - 335,954 B/token vs KV's 88,064 (79% of per-token memory). Recomputed from a raw [B,T,hidden] buffer every step while the compressor only consumes a stride-ratio reduction. A rolling compressed state is 4-128x and gates whether B=128-256 is reachable at all." + } + }, + "QUALITY_MODEL": { + "target": "GSM8K >= 90 on OUR protocol: n=100, 0-shot chat, 2048-token cap, seed 161", + "current_status": "NO CURRENT NUMBER EXISTS. 87.0% is PROVISIONAL and from superseded math.", + "why_87_is_void": "PR #35 (830a41ed9) changed decode math after it was measured: a SwiGLU clamp was missing on 4 of 5 expert paths INCLUDING the shared expert, which every token traverses in every layer (fixture magnitude when it bites: 0.7311 clamped vs 14.8996 unclamped = 20.4x); and YaRN was wrongly applied to ratio-0 layers (correct set exactly {0,1,43}). The new number may land either side of 87.0 and LOWER IS NOT A REGRESSION - it is the first number measured on math that matches the reference.", + "reference_comparison": "Published V4-Flash-Base scores 90.8 with 8-SHOT. That is a DIFFERENT AND EASIER protocol than our 0-shot. Always state this when comparing; we are not 3.8 points behind, we are measuring a harder task.", + "ladder_so_far": [ + { + "session": "s1", + "score": 64.0, + "note": "640-token cap, 33 truncated - cap-limited, not quality-limited" + }, + { + "session": "s2", + "score": 84.0 + }, + { + "session": "s3", + "score": 87.0, + "note": "on GPU-Viterbi bake; NOW PROVISIONAL" + } + ], + "levers_to_90_plus_ranked": [ + { + "lever": "Re-measure on corrected math (the SwiGLU clamp fix)", + "expected_direction": "neutral-to-better", + "cost": "~40 min, ~$3", + "status": "NOT RUN - this is step 1 and it may already clear the bar" + }, + { + "lever": "Arc Boost: confidence-weighted voting k=5", + "expected_direction": "up", + "cost": "one session", + "status": "unfired; was blocked on the xs_history multi-chain crash, fixed in PR #21" + }, + { + "lever": "Exhaustive search instead of beam", + "expected_direction": "up (+0.0013..+0.0021 cos, wins 8/9 fixture cells)", + "status": "the qtip2b rung uses exhaustive by default" + }, + { + "lever": "LDLQ/GPTQ-style block error feedback stacked on rotation", + "expected_direction": "up", + "status": "rotation only flattens the DIAGONAL; off-diagonal correlations survive and remain recoverable. Opt-in calibration pass for tenants with data." + }, + { + "lever": "Higher beam width", + "expected_direction": "up but tiny", + "note": "W=256 vs W=32 costs ~1% of bake time (82.7 vs 83.6 s/layer on H200) - ALWAYS BAKE W=256, quality is free here" + } + ], + "quality_is_not_the_bottleneck": "87.0 (provisional) against a 90.8 8-shot reference on a harder protocol, with 4 unfired levers and a bake-time quality upgrade that costs 1%. The hard problem is throughput, not quality." + }, + "ECONOMICS": { + "measured_today": { + "B": 64, + "aggregate_tok_s": 8.14, + "usd_per_Mtok": 165.51, + "box": "H200 @ $4.85/hr" + }, + "at_physics_ceiling": { + "B": 256, + "aggregate_tok_s": 16602, + "usd_per_Mtok": 0.081, + "note": "4.85 / (16602 * 3600 / 1e6)" + }, + "interpretation": "The distance between $165/Mtok and $0.08/Mtok is entirely engineering. Every step toward it is compute that does not need to be bought and power that does not need to be burned." + } +} \ No newline at end of file diff --git a/memory/mission/EXTERNAL_FINDINGS.md b/memory/mission/EXTERNAL_FINDINGS.md new file mode 100644 index 000000000..7ad8b25f2 --- /dev/null +++ b/memory/mission/EXTERNAL_FINDINGS.md @@ -0,0 +1,99 @@ +# EXTERNAL_FINDINGS.md — cross-project intel worth acting on. +# Source: colibrì (JustVugg/colibri) v1.6.1 @df2c248, researched 2026-08-14. +# Colibrì = pure-C MoE engine streaming experts from NVMe (consumer HW). +# Opposite mission to Arc, but it independently implements DeepSeek V4 Flash, +# which makes it a REFERENCE ORACLE for our V4 port. + +## F1 (DEFECT, ours) — V4 Lightning Indexer scores wrongly. HIGH VALUE. +Two independent references agree AGAINST Arc: SGLang `indexer.py:78-88` +(vendored in our own research/code/06_foundation/sglang) and colibrì +`deepseek_v4.c` (written from the checkpoint). +REFERENCE: `logit[q,c] = kv_scale_c * Σ_h relu(q_h·k_c) * w_h[q]` → ONE top-k over c. +ARC (dsv4_indexer.rs:366-380, arc-cuda-graph/src/cuda/flashmlasparse/ +indexer_score.cu, and its cpu_reference — all self-consistent, all wrong): +`score[h,q,c] = (q_h·k_c) * w_h[q]` → top-k PER HEAD. +Three deltas: (a) no ReLU (negatives cancel positives across heads instead of +clamping → different keys chosen); (b) no sum over heads — per-head selection +also multiplies sparse-gather traffic by n_heads(64), a PERF bug for our fleet +wedge since shared page selection is what makes the gather cheap; (c) no +Hadamard + low-precision QDQ on indexer Q (SGLang fuses it; colibrì does +Hadamard → FP4 QDQ group-32 → bf16 round, bf16 rounding after every stage). +Also omitted: `weight_scale = softmax_scale * n_heads^-0.5` (uniform positive +⇒ selection-neutral; matters only if logits are consumed numerically). +STATUS: DORMANT — dsv4_attention.rs:33-45 attends densely over raw∪compressed, +and CSA top-512 covers all compressed entries while ctx ≤ index_topk*ratio=2048. +Bites the moment the indexer goes live for long context = exactly V4 Flash's +value. Fix now (free) vs later (rented-GPU debug session). + +### F1 STATUS: FIXED in PR #26 (2026-08-14, draft). +Verified against SGLang source directly: `fp8_paged_mqa_logits_torch:84-89` = +linear → relu → *q_scale → sum(dim=1) → *kvcache_scale, ONE topk over [batch,seq]. +`C4Indexer.__init__:519`: weight_scale = head_dim^-0.5 * n_heads^-0.5 (folded +into weights pre-kernel). kv_scale is the paged-FP8 dequant scale — ≡1 for Arc's +BF16/F32 indexer K, but NOT selection-neutral if/when paged-FP8 indexer K lands. +Arc now: indexer_logits(q,k,w) = matmul→relu→*w→sum-over-heads, out [B,T_q,topk]; +CUDA regridded (B·H,T_q)→(B,T_q) with head loop inside the block (warp-owned +column ⇒ atomic-free, deterministic); entry points renamed *_score_*→*_logits_* +so stale callers LINK-ERROR instead of silently computing the old formula. +Dormancy proven (zero `.indexer` reads in deepseek4.rs; dense union exact while +ctx ≤ 512×4=2048). Tests: force-full-selection ⇒ bit-exact dense equivalence +(both Rust + CUDA cpu_reference) + a fixture where old/new pick DIFFERENT keys. +**NEW FINDINGS from that work (act before enabling the sparse path):** +(i) Arc's compressor builds indexer K **per-head** [B,H,T_c,D] while SGLang's +indexer cache is **MQA (num_heads_kv=1)** ⇒ Arc pays ~64× the K memory/compute +for the indexer. (ii) The dead `indexer` field costs weight memory per CSA layer +for a path that never runs today. (iii) Hadamard + FP4-QDQ on indexer Q still +deferred (separate variable). + +## F2 (IDEA, maybe) — lossless entropy coding of ALREADY-quantized weights. +Orthogonal to TurboQuant (compresses quantized bytes; does NOT requantize). +`int4-rans256-g0`: static rANS, 256 round-robin-interleaved streams/tensor so +wide decoders get coalesced output; ~0.76 ratio on per-row int4 GLM-5.2, +byte-exact. Measured VRAM tier (DietGPU ANS, warm int4 experts compressed in +VRAM, decoded to scratch before grouped GEMM): 6×5090 host went 9,335 → 10,628 +resident experts (+13.9%), decode +4.5%..+15%. +CAVEATS THEY DOCUMENT (we'd inherit): interleave granularity is what makes GPU +decode viable; moving experts between execution paths changed FP accumulation +order enough to diverge a greedy A/B even with byte-exact weights. +NEGATIVE RESULT THAT GATES US: on group-scaled (g64) containers ratio collapses +to ~0.89 (codes better normalized, scales dominate the stream). **If our +TurboQuant/qtip output is group-scaled, entropy coding likely won't pay — +MEASURE THE RATIO ON A REAL CONTAINER BEFORE BUILDING ANYTHING.** + +## F3 (EMPIRICAL, saves GPU-hours) +- Quantized draft head kills speculation: int4 MTP head → 0-4% acceptance; + int8 mandatory. (Relevant if we ever qtip the V4 MTP head — we currently + ISQ it; VERIFY.) +- Draft and verify MUST run the same kernel family: their SPEC_PIN=1 default + exists because different kernels for draft vs verify silently destroyed + acceptance. Arc has multiple GEMV paths → same trap. +- Speculation can be NET NEGATIVE: MTP measured -32% at ~85% expert hit rate. +- Quantization-granularity failures show up as **EOS starvation, not + perplexity**: per-row int4 vs group-64 ≈9pp worse, manifesting as think-mode + loops / never-terminating generations. ⇒ add a TERMINATION check to the + TurboQuant K4/V3 quality gate; perplexity would miss it. + +## F4 (TEST METHOD, steal it) +Validate sparse attention by **forcing full-key selection and requiring exact +reproduction of dense attention**. Colibrì validates DSA this way; Arc has no +equivalent (v4_e2e.rs has one synthetic finite/deterministic check). This test +would have caught F1 at ZERO GPU cost. Also: teacher-forcing oracle vs +transformers (they hold 30-32/32); stamp-mandatory format registry (an +entropy-coded tensor has no expected_bytes(O,I) ⇒ size-inference structurally +impossible ⇒ stamp-gated dispatch must precede inference — a real constraint +for any Arc container with data-dependent size). + +## F5 (IGNORE) — their CUDA backend +Zero cudaGraph calls, no cuBLAS/CUTLASS, scalar FP32 kernels (one block per +output element, shared-mem tree reductions), only tensor-core path is a dead +wmma s4 fragment. arc-cuda-graph + our QTIP GEMV work is well ahead. Their +fmt=6 E8/IQ3 lattice + block-diagonal FWHT is QuIP#/QTIP territory but below +trellis on rate-distortion. Disk-streaming/dual-SSD/expert-placement JIT = +consumer answers to a problem we don't have. TWO exceptions worth borrowing: +**batch-union** (read each unique expert once per BATCH, not per position — +maps onto batch_load_probe work) and their measurement that **routing is 71.6% +predictable one layer ahead** (prior for any future expert-offload path). + +## DISCIPLINE (Jish, 2026-08-14) +**Never test 5 things at once — you can't detect the menace then.** One variable +per experiment, especially on rented GPU time. diff --git a/memory/mission/wave13-AH-config.md b/memory/mission/wave13-AH-config.md new file mode 100644 index 000000000..ead6dade6 --- /dev/null +++ b/memory/mission/wave13-AH-config.md @@ -0,0 +1,455 @@ +# Wave 13 · Agent AH — Ground Arc's V4 port in the real `config.json` + +**Date**: 2026-08-14 · **Branch**: `fix/v4-config-grounded` off `origin/master` @ `bef703326` +**Predecessor**: PR #27 / `docs/v4-reference-audit` (agent AA), log `memory/mission/wave11-AA-audit.md` + +## Mission + +The V4 reference audit produced 25 findings but stalled on one missing artifact: the +real `config.json` from `deepseek-ai/DeepSeek-V4-Flash`. Five findings bottomed out +there. This task fetched it, settled those five, corrected Arc's demonstrably-fake +"real config" fixture, and fixed what the config proves wrong. + +## 1. The fetched artifacts + +All three fetches succeeded. The repo is public and MIT-licensed; no token needed. + +| Artifact | URL | Result | +|---|---|---| +| `config.json` | `https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash/resolve/main/config.json` | HTTP 200, 1749 bytes | +| `model.safetensors.index.json` | same host, `/resolve/main/model.safetensors.index.json` | HTTP 200, 5 371 381 bytes, 69 187 tensors, `total_size` 159 609 485 896 | +| shard-2 safetensors header | ranged GET on `model-00002-of-00046.safetensors` | HTTP 206, 172 232-byte JSON header | + +### 1.1 `config.json` — VERBATIM + +```json +{ + "architectures": [ + "DeepseekV4ForCausalLM" + ], + "attention_bias": false, + "attention_dropout": 0.0, + "bos_token_id": 0, + "eos_token_id": 1, + "expert_dtype": "fp4", + "hc_eps": 1e-06, + "hc_mult": 4, + "hc_sinkhorn_iters": 20, + "head_dim": 512, + "hidden_act": "silu", + "hidden_size": 4096, + "index_head_dim": 128, + "index_n_heads": 64, + "index_topk": 512, + "initializer_range": 0.02, + "max_position_embeddings": 1048576, + "model_type": "deepseek_v4", + "moe_intermediate_size": 2048, + "n_routed_experts": 256, + "n_shared_experts": 1, + "norm_topk_prob": true, + "num_attention_heads": 64, + "num_experts_per_tok": 6, + "num_hidden_layers": 43, + "num_hash_layers": 3, + "num_key_value_heads": 1, + "num_nextn_predict_layers": 1, + "o_groups": 8, + "o_lora_rank": 1024, + "q_lora_rank": 1024, + "qk_rope_head_dim": 64, + "quantization_config": { + "activation_scheme": "dynamic", + "fmt": "e4m3", + "quant_method": "fp8", + "scale_fmt": "ue8m0", + "weight_block_size": [ + 128, + 128 + ] + }, + "rms_norm_eps": 1e-06, + "rope_scaling": { + "beta_fast": 32, + "beta_slow": 1, + "factor": 16, + "original_max_position_embeddings": 65536, + "type": "yarn" + }, + "rope_theta": 10000, + "routed_scaling_factor": 1.5, + "scoring_func": "sqrtsoftplus", + "sliding_window": 128, + "swiglu_limit": 10.0, + "tie_word_embeddings": false, + "topk_method": "noaux_tc", + "torch_dtype": "bfloat16", + "transformers_version": "4.57.1", + "use_cache": true, + "vocab_size": 129280, + "compress_rope_theta": 160000, + "compress_ratios": [0, 0, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 128, 4, 0] +} +``` + +### 1.2 Structure derived from `model.safetensors.index.json` + +Tensor-name families (`N` = any integer), with counts: + +``` +1 embed.weight 1 head.weight 1 norm.weight +1 hc_head_{base,fn,scale} +43 layers.N.attn.attn_sink 43 layers.N.attn.{q_norm,kv_norm}.weight +43 layers.N.attn.{wq_a,wq_b,wkv,wo_a,wo_b}.{weight,scale} +41 layers.N.attn.compressor.{ape,norm.weight,wgate.weight,wkv.weight} +21 layers.N.attn.indexer.compressor.{ape,norm.weight,wgate.weight,wkv.weight} +21 layers.N.attn.indexer.{weights_proj.weight,wq_b.weight,wq_b.scale} +11008 layers.N.ffn.experts.N.{w1,w2,w3}.{weight,scale} (= 43 × 256) +43 layers.N.ffn.gate.weight 40 layers.N.ffn.gate.bias 3 layers.N.ffn.gate.tid2eid +43 layers.N.ffn.shared_experts.{w1,w2,w3}.{weight,scale} +43 layers.N.{attn_norm,ffn_norm}.weight +43 layers.N.hc_{attn,ffn}_{base,fn,scale} +1 mtp.0.* (attn, ffn w/ 256 experts, e_proj, h_proj, enorm, hnorm, norm, hc_*, hc_head_*) +``` + +Cross-checks that this settles: + +* **`compress_ratios` semantics confirmed.** `Counter(compress_ratios) == {4: 21, 128: 20, 0: 3}`. + Indices of `0` are exactly `[0, 1, 43]`; of `4` exactly `[2, 4, …, 42]`; of `128` exactly + `[3, 5, …, 41]`. The 41 layers carrying `attn.compressor.*` are exactly the 41 layers + with a non-zero ratio (all but 0 and 1). Slot 43 is the MTP block. +* **The Lightning Indexer lives on CSA (ratio-4) layers only.** The 21 layers with + `attn.indexer.*` are exactly `[2, 4, …, 42]` — the ratio-4 set. Ratio-128 (HCA) layers + have no indexer. +* **Hash layers are exactly layers 0–2.** `gate.tid2eid` exists on layers `{0, 1, 2}` and + those are precisely the three layers with **no** `ffn.gate.bias` — hash-routed layers + carry no noaux_tc correction bias, which is internally consistent with `num_hash_layers: 3`. +* **The MTP block is NOT hash-routed.** `mtp.0.ffn.gate.bias` exists and no + `mtp.*.tid2eid` exists anywhere. This independently confirms the reference's + `is_hash = layer_id < n_hash_layers and not (is_deepseek_v4 and is_nextn)` exclusion. +* **Every one of the 43 layers is MoE.** 11008 = 43 × 256 expert tensors, and no dense + `mlp.*` family exists — consistent with `first_k_dense_replace` being absent (Arc + defaults it to 0). + +### 1.3 Shapes and dtypes (shard-2 header, layer 0) + +``` +layers.0.attn.attn_sink F32 [64] = num_attention_heads +layers.0.attn.q_norm.weight BF16 [1024] = q_lora_rank +layers.0.attn.kv_norm.weight BF16 [512] = head_dim +layers.0.attn.wq_a.weight F8_E4M3 [1024, 4096] scale F8_E8M0 [8, 32] +layers.0.attn.wq_b.weight F8_E4M3 [32768, 1024] = 64 heads × 512 +layers.0.attn.wkv.weight F8_E4M3 [512, 4096] = 1 KV head × 512 +layers.0.attn.wo_a.weight F8_E4M3 [8192, 4096] = o_groups(8) × o_lora_rank(1024) +layers.0.attn.wo_b.weight F8_E4M3 [4096, 8192] +layers.0.ffn.gate.weight BF16 [256, 4096] +layers.0.ffn.gate.tid2eid I64 [129280, 6] = [vocab_size, num_experts_per_tok] +layers.0.hc_attn_scale F32 [3] +layers.0.hc_attn_base F32 [24] = m² + 2m for m = hc_mult = 4 +layers.0.hc_attn_fn F32 [24, 16384] = [24, hc_mult × hidden_size] +layers.0.ffn.experts.0.w1 I8 [2048, 2048] scale F8_E8M0 [2048, 128] +layers.0.ffn.shared_experts.w1 F8_E4M3 [2048, 4096] scale F8_E8M0 [16, 32] +``` + +Every config-derived dimension Arc computes is confirmed: `head_dim = 512`, +`q_lora_rank = 1024`, `o_groups × o_lora_rank = 8192`, `num_key_value_heads = 1`, +`hc_mult = 4` (encoded in the `[24, …]` mHC tensor widths, so it is not a free knob). + +## 2. The five findings, settled + +### 2.1 Finding 11 — group-limited routing → **SETTLED, not a defect** + +`grep -o 'n_group\|topk_group'` on the real config returns **nothing**. Neither key is +published. Arc's serde defaults (`n_group = 1`, `topk_group = 1`) therefore apply, the +`if self.cfg.n_group > 1` guard is false, and `MoeGate::forward` takes the flat top-k +branch — the same end state SGLang reaches by force-disabling group limiting for V4. + +Arc's original premise ("V4 Flash config does NOT publish `n_group`") was **correct**. +The audit's counter-inference from SGLang's load-time patch comment ("HF config.json +inherits `topk_group=4` from the V3 template") does not hold for the shipped file — that +patch is defensive, not descriptive. Severity drops from HIGH-conditional to **not +applicable to this checkpoint**. + +The two secondary defects the audit found *inside* that branch (masking by multiply +instead of `-inf`; group score always `topk(2).sum` instead of `amax` when unbiased) are +**still real**, but are now confirmed dead code for V4 Flash. They remain live for V3-style +configs via `deepseek3.rs:534`, which is where they should be fixed. + +### 2.2 Finding 12 — hash-layer key spelling → **SETTLED, not a defect** + +The shipped key is **`num_hash_layers`** (value `3`). `n_hash_layers` does **not** appear. + +That is precisely the key SGLang's *runtime* reads — `getattr(config, "num_hash_layers", 0)` +(`deepseek_v2.py:483`). So the reference runs three hash layers, exactly as Arc does. The +feared "reference runs 0 hash layers where Arc runs 3" divergence **does not occur**. +SGLang's `n_hash_layers` dataclass field (`configs/deepseek_v4.py:107`) is simply an unused +declaration. Arc's `alias = "n_hash_layers"` is harmless belt-and-braces. + +Independently corroborated by the checkpoint: `gate.tid2eid` exists on exactly layers 0, 1, 2. + +**Rider (ii) also settled, and in Arc's favour.** The audit worried that the reference +stores `tid2eid` as `int32` while Arc requests `DType::I64` (`deepseek4.rs:1554-1559`). +The shipped tensor is **`I64 [129280, 6]`**. Arc's request is correct; an I32 read would +have been the bug. Shape matches Arc's `[vocab_size, num_experts_per_tok]` implementation +and confirms the audit's own correction of the stale `research/v4_audit.md` guess. + +**Rider (i) stands as written**: the reference excludes the MTP block from hash routing +unconditionally; Arc keys purely on virtual layer index. For the real 43-layer config these +coincide (MTP sits at virtual index 43 ≥ 3), and the checkpoint confirms the MTP block is +biased-routed (`mtp.0.ffn.gate.bias` present, no `mtp.*.tid2eid`). It remains a hazard only +for small synthetic fixtures — `synthetic_load_smoke.rs:340-344` depends on the current +behavior. + +### 2.3 Finding 24 — `rope_scaling` / mscale → **SETTLED, not a defect** + +`rope_scaling` **is** present — the audit was right that Arc's old fixture (which omitted +it entirely) could not have been the real config: + +```json +"rope_scaling": { + "beta_fast": 32, "beta_slow": 1, "factor": 16, + "original_max_position_embeddings": 65536, "type": "yarn" +} +``` + +It carries **no `mscale` and no `mscale_all_dim`**. Arc's `DeepSeekV2RopeScaling::Yarn` +already defaults both to `1.0` (`layers.rs:1440-1445`), so the magnitude factor Arc +multiplies into sin/cos is `yarn_get_mscale(16, 1.0) / yarn_get_mscale(16, 1.0) == 1.0`. +Arc's mscale multiply is a **no-op**, exactly matching the reference's unit-magnitude +`torch.polar(torch.ones_like(freqs), freqs)`. Finding 24 resolves to **no divergence**; +the existing code comment at `layers.rs:1432-1439` was already right. + +### 2.4 Finding 5 rider (i) — `swiglu_limit` → **PUBLISHED; finding 5 CONFIRMED and hardened** + +`"swiglu_limit": 10.0` is present. The reference's +`getattr(config, "swiglu_limit", None)` therefore resolves to `10.0` and the reference +**does** clamp. Arc's hard default of `10.0` coincides with the shipped value. + +This removes the escape hatch from finding 5: the clamp is **mandatory**, not a +config-dependent maybe. Every Arc expert path that drops it is a genuine numerical +divergence against the trained weights — most seriously the **shared expert, on every +device and backend** (`deepseek4.rs` → `Mlp::new`, `layers.rs:3001-3037` has no limit +field). Finding 5 stays **HIGH** and is now unconditional. + +### 2.5 `compress_rope_theta` → **SETTLED; Arc's default is right** + +`"compress_rope_theta": 160000`. Arc's `default_compress_rope_theta() -> 160000.0` +matches the shipped value exactly. SGLang's dataclass default of `40000` +(`configs/deepseek_v4.py:104`) is never reached for this checkpoint, so the 4× +discrepancy the audit flagged is a dead default, not a divergence. `rope_theta` is +`10000`, and the per-layer θ dispatch Arc implements is correct. + +## 3. What changed in the tree + +All edits are in `mistralrs-core/src/models/deepseek4.rs`. Scope fences respected: no +touch to `mistralrs-quant/**`, `mtp_pipeline.rs`, the MTP block, `dsv4_attention.rs`, +`sinks.rs`, or `dsv4_indexer.rs`. + +1. **Fixture replaced with the verbatim published config.** The old fixture was + hand-written and omitted `rope_scaling`, `norm_topk_prob`, `quantization_config`, + `hidden_act`, `attention_bias`, `tie_word_embeddings`, `expert_dtype` and + `num_nextn_predict_layers`, while *adding* a `first_k_dense_replace` key that the real + config does not have. It is now a module-level `const V4_FLASH_CONFIG_JSON` reproduced + byte-for-byte from the fetch, with the source URL and fetch date in its doc comment. + +2. **`v4_flash_real_config_parses` extended.** Proves the real file parses through + `DeepSeekV4Config` including the fp8 `quantization_config` block, and pins the + absent-key set (`intermediate_size`, `kv_lora_rank`, `v_head_dim`, `qk_nope_head_dim`) + plus `first_k_dense_replace == 0`. Added `layer_compress_ratio(42) == 4` and an + assertion that the ratio-0 set is exactly `{0, 1, 43}`. + +3. **New test `v4_flash_config_settles_the_audit_questions`.** One assertion block per + audit finding (11, 12, 24, 5-rider, `compress_rope_theta`, 21, 14-context), each + commented with the finding it closes, so a future re-fetch that contradicts one fails + loudly rather than silently. + +4. **`norm_topk_prob` now parsed** (audit finding 21), `#[serde(default = "…")] = true`, + and used to gate renormalization in `MoeGate::forward`: + `if self.cfg.norm_topk_prob && matches!(scoring_func, Sigmoid | SqrtSoftplus)`. + Behavior-neutral: the real config publishes `true`, and no fixture in the tree sets + `false` (`grep '"norm_topk_prob": false'` → no hits). The reference passes the flag + straight through as `renormalize=config.norm_topk_prob` (`deepseek_v2.py:553`). + +5. **Comments upgraded from claim to verified fact.** The `MoeGate` `n_group` comment + cited `/tmp/v4_flash_config.json` — a path that no longer exists and that CI could + never check; it now cites the resolve URL and names the test that pins it. The struct + doc comment's "V4: not in real config" claims are now labelled verified and enumerate + the absent `mscale`/`mscale_all_dim` too. + +### Verification + +``` +cargo check -p mistralrs-core --tests → clean (8 pre-existing warnings only) +cargo test -p mistralrs-core --lib deepseek4 → 22 passed; 0 failed +cargo test -p mistralrs-core → see PR checks +``` + +No fixture-dependent test needed correcting — none of them were pinning wrong behavior. +The old fixture was *incomplete* rather than *wrong*: every assertion it made was true of +the real config; it simply could not answer the questions the audit needed answered. + +## 4. Findings status after this task + +**Settled — closed, no code change needed (4)** + +| # | Was | Now | +|---|---|---| +| 11 | Group-limited routing — HIGH, conditional | Neither `n_group` nor `topk_group` ships. Branch is dead for V4 Flash. Arc's premise was right. | +| 12 | Hash key mismatch — HIGH, UNVERIFIABLE | Ships as `num_hash_layers: 3`, the key SGLang's runtime reads. No divergence. Rider (ii) also clears: tensor is I64, as Arc requests. | +| 24 | mscale — UNVERIFIABLE | No `mscale`/`mscale_all_dim` in `rope_scaling` ⇒ Arc's multiply is exactly 1.0 ⇒ matches the reference's unit-magnitude polar. | +| — | `compress_rope_theta` default | Ships as 160000 = Arc's default. SGLang's 40000 dataclass default is unreachable. | + +**Confirmed real and hardened by the config (3)** + +| # | Finding | What the config adds | +|---|---|---| +| 5 | swiglu clamp missing on 4 of 5 expert paths | `swiglu_limit: 10.0` IS published ⇒ the reference clamps ⇒ the omission is unconditional, not config-dependent. Severity HIGH stands. | +| 6 | YaRN applied to ratio-0 layers by default | `rope_scaling` present with `factor: 16`, `original_max_position_embeddings: 65536` ⇒ the reference's `original_seq_len = 0` override for ratio-0 layers is live and Arc's default diverges. **Correction: the affected layers are 0, 1 and the MTP block — not "0, 1, 42". Layer 42 has ratio 4.** | +| 14 | FP8 K-cache scale continuous, not UE8M0 | `quantization_config.scale_fmt: "ue8m0"`, and every shipped `*.scale` tensor has dtype **`F8_E8M0`**. Direct evidence. | + +**Fixed here (2)** + +| # | Finding | Fix | +|---|---|---| +| 21 | `norm_topk_prob` not parsed | Field added, defaulted `true`, gates renormalization. Behavior-neutral today. | +| — | Fixture is not the real config | Replaced verbatim; two tests now pin the answers. | + +**Open, unchanged by the config (16)**: 1, 2, 3, 4, 7, 8, 9, 10, 13, 15, 16, 17, 18, 19, +20, 22, 23, 25. Finding 20 (`hc_mult`/`hc_eps`/`hc_sinkhorn_iters` dead) is confirmed +*latent only*: the shipped values (4 / 1e-6 / 20) equal Arc's hardcoded constants, and +`hc_mult` is additionally pinned by the `[24, 16384]` mHC tensor widths, so a mismatched +config would fail a shape check rather than silently compute wrong math. + +## 5. Left deliberately — outside the fence + +Written up precisely rather than reached for. + +**(a) Finding 5, the shared-expert clamp.** `deepseek4.rs:1832-1839` builds the shared +expert via `Mlp::new` (`layers.rs:3001-3037`), which has no `limit` field, so the clamp is +dropped on every device. Also dropped by `experts.rs:465` (`forward_fused`, the default +un-ISQ'd CUDA path), `experts.rs:570` (Metal), `experts.rs:632-637` (`forward_slow`). +Only `forward_fast`'s CUDA branch (`experts.rs:541-549`) clamps. +*Proposed test*: build a shared-expert `Mlp` with `swiglu_limit = 1.0`, feed a gate +pre-activation of `+50`, and assert the output equals the `limit`-clamped reference rather +than the unclamped one; repeat per expert path with `ARC_MOE_SLOW` set. +Behavior change in the expert MLP, not config — belongs to whoever owns `moe/experts.rs`. + +**(b) Finding 6, the YaRN default flip.** `deepseek4.rs` gates the reference behavior +behind `ARC_DISABLE_YARN_STD`. The config now proves the reference path is the correct +default. The flip is a one-line change but is an attention-path behavior change; PR #28 +owns `dsv4_attention.rs` and this interacts with it. +*Proposed test*: assert the ratio-0 rotary table's `inv_freq` equals plain +`1/θ^(2i/d)` (no interpolation) while the ratio-4 table shows the YaRN ramp. + +**(c) NEW — the single `quantization_config` block does not describe the routed experts.** +The config publishes one fp8 block (`quant_method: "fp8"`, `weight_block_size: [128, 128]`), +which matches the attention and shared-expert tensors exactly (`wkv.weight` F8_E4M3 +`[512, 4096]` with scale `[4, 32]` = ceil(512/128) × ceil(4096/128)). It does **not** +describe the routed experts, which ship as packed FP4 — `w1.weight` dtype **`I8`** +`[2048, 2048]` (two nibbles per byte over a logical `[2048, 4096]`) with scale +`F8_E8M0 [2048, 128]`, i.e. **per-row group-32**, not 128×128 blocks. The distinguishing +key is `"expert_dtype": "fp4"`, which Arc does not parse. Anything that loads the native +checkpoint and trusts `quantization_config` uniformly will mis-read 11008 of the 69187 +tensors. Owner: native-checkpoint loading / `mistralrs-quant` (fenced from this task). + +**(d) NEW — `num_nextn_predict_layers: 1` is published and Arc does not parse it.** Arc +infers MTP presence from tensor availability instead. Not wrong today, but the config +carries the authoritative count. Owner: PR #30 (MTP). + +## 6. Proposed audit-doc wording + +PR #27 is still **open** at the time of writing (`state: OPEN`, `mergedAt: null`), so the +doc was not edited directly. Exact replacements below. + +**Replace the "The single missing artifact" block-quote (`docs/notes/v4-reference-audit.md`, +after the "Reference availability caveat" quote) with:** + +> **The missing artifact — now fetched.** Five findings bottomed out in the real +> `config.json` from `deepseek-ai/DeepSeek-V4-Flash`. It was retrieved on 2026-08-14 from +> `https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash/resolve/main/config.json` (public, +> MIT, no token) and is now vendored verbatim as `V4_FLASH_CONFIG_JSON` in +> `mistralrs-core/src/models/deepseek4.rs`, with `v4_flash_config_settles_the_audit_questions` +> pinning each answer. **Findings 11, 12 and 24 are closed as non-defects**; the +> `compress_rope_theta` rider is closed in Arc's favour; and **finding 5 is hardened** — +> `swiglu_limit: 10.0` is published, so the reference clamps unconditionally and Arc's four +> unclamped expert paths are a real divergence with no config-dependent escape. See +> `memory/mission/wave13-AH-config.md`. + +**Verdict table (§0) — amend these rows:** + +* Row 5: severity unchanged, append to the finding text: `— config publishes swiglu_limit: 10.0, so the clamp is unconditional`. +* Row 11: change Severity to `CLOSED` and the finding text to + `Group-limited routing force-disabled upstream — real config publishes neither n_group nor topk_group, so Arc's branch is dead code for V4 Flash. No divergence.` +* Row 12: change Severity to `CLOSED`, Confidence to `both-lbl`, and the text to + `Real config spells it num_hash_layers: 3 — the key SGLang's runtime reads. Reference and Arc both run 3 hash layers. tid2eid also ships as I64, matching Arc's request.` +* Row 24: change Severity to `CLOSED`, and the text to + `rope_scaling publishes no mscale/mscale_all_dim ⇒ Arc's defaults make the multiply exactly 1.0 ⇒ matches the reference's unit-magnitude polar.` +* Row 14: append to Confidence cell: `— corroborated: config sets scale_fmt: "ue8m0" and every shipped scale tensor is dtype F8_E8M0`. +* Row 21: change Severity to `FIXED` and append `— norm_topk_prob (published as true) is now parsed and gates renormalization`. + +**§1(b), the YaRN paragraph — correct the affected-layer list.** Replace + +> exactly the layers Arc's own module docs (`dsv4_attention.rs:12-28`) blame for the +> RUN-161 long-context repetition collapse + +'s preceding clause `on layers 0, 1, 42 and the MTP block` with `on layers 0 and 1 and +the MTP block`. `compress_ratios` index 42 is `4` (CSA), not `0`; the ratio-0 indices are +exactly `[0, 1, 43]`, and index 43 *is* the MTP slot. + +**§1(b), the mscale paragraph — replace `**mscale — UNVERIFIABLE.**` … `Settled by the +missing config.json.` with:** + +> **mscale — SETTLED, MATCHES.** Reference returns +> `torch.polar(torch.ones_like(freqs), freqs)` (`deepseek_v4_rope.py:57`) — unit magnitude. +> The published `rope_scaling` carries no `mscale` and no `mscale_all_dim`, so Arc's serde +> defaults (both `1.0`, `layers.rs:1440-1445`) apply and +> `yarn_get_mscale(16, 1.0) / yarn_get_mscale(16, 1.0) == 1.0`. Arc's multiply is a no-op. +> No divergence. + +**§1(b), the `compress_rope_theta` paragraph — replace with:** + +> **`compress_rope_theta` — SETTLED.** The published config sets `160000`, matching Arc's +> default (`deepseek4.rs`). SGLang's dataclass default of `40000` +> (`configs/deepseek_v4.py:104`) is never reached for this checkpoint. No action. + +**§4(b2) — prepend to the section body:** + +> **SETTLED (2026-08-14): this branch is dead for V4 Flash.** The published `config.json` +> contains neither `n_group` nor `topk_group`, so Arc's serde defaults (1/1) apply and the +> flat top-k branch is taken — the same end state SGLang reaches by force-disabling group +> limiting. Arc's premise was correct and SGLang's load-time patch comment is defensive, +> not descriptive. The two secondary defects below remain real but are reachable only via +> V3-style configs (`deepseek3.rs:534`), which is where to fix them. + +**§4(d′) — replace the whole section with:** + +> ### (d′) `num_hash_layers` vs `n_hash_layers` — **SETTLED: MATCHES** · `both-lbl` +> +> The published config spells it **`num_hash_layers: 3`**; `n_hash_layers` does not +> appear. That is the key SGLang's runtime reads +> (`getattr(config, "num_hash_layers", 0)`, `deepseek_v2.py:483`), so the reference runs +> three hash layers exactly as Arc does. SGLang's `n_hash_layers` dataclass field +> (`configs/deepseek_v4.py:107`) is an unused declaration; Arc's `alias` is harmless. +> Corroborated by `model.safetensors.index.json`: `gate.tid2eid` exists on exactly layers +> 0, 1, 2, which are also the only three layers **without** an `ffn.gate.bias`. +> +> **Rider (ii) settled in Arc's favour**: the shipped tensor is `I64 [129280, 6]`, which is +> what Arc requests (`deepseek4.rs:1554-1559`). An I32 read would have been the bug. +> +> **Rider (i) stands**: the reference excludes the MTP block from hash routing +> unconditionally and Arc keys purely on virtual layer index. The checkpoint confirms the +> exclusion (`mtp.0.ffn.gate.bias` present, no `mtp.*.tid2eid`); for the real 43-layer +> config the two agree, so this is a synthetic-fixture hazard only. + +**§4(e), note (i) — replace `— another thing the real config.json settles` with:** + +> — **settled: the real config publishes `"swiglu_limit": 10.0`**, so the reference's +> `getattr` resolves and it clamps. Arc's hard default coincides with the shipped value. +> The clamp is therefore mandatory and the four unclamped paths are unconditional +> divergences. + +**§12 (Recommended action order) — the "fetch config.json" item is done; replace it with +the follow-ups this fetch created:** flip the finding-6 YaRN default; clamp the shared +expert and the three other expert paths (finding 5, now unconditional); teach the loader +that routed experts are fp4 group-32, not the config's fp8 128×128 (new, §5(c) of the +wave-13 log). diff --git a/memory/mission/wave16-AF-beam-perf.md b/memory/mission/wave16-AF-beam-perf.md new file mode 100644 index 000000000..746837ae9 --- /dev/null +++ b/memory/mission/wave16-AF-beam-perf.md @@ -0,0 +1,427 @@ +# wave16-AF — why the CUDA beam kernel got 2.1×, not 6–12× + +**Status:** occupancy measured twice on real toolchain output; parity proved on CPU; **wall time still unmeasured** — +experiments 2 and 3 run after the bake. PR #40 stays DO NOT MERGE until then. + +**Author's note up front:** I wrote the kernel and I wrote the 6–12× projection. That projection was wrong, and it was +wrong for a reason I can now name precisely: **I modelled bytes and instructions, and the kernel is bound by neither — +it is bound by issue efficiency at low occupancy.** The bytes model was validated against the exhaustive kernel's wall +time and appeared to fit; that fit was a coincidence, and I did not test it. This document is the correction, and it is +built on measurements rather than on another model, wherever a measurement was possible without a GPU. + +**No GPU was touched.** A paid H200 is baking. Everything below is source analysis plus CPU measurement of the beam +search's data-dependent behaviour. Every hardware claim is listed as a proposed experiment at the end. + +## The measurement being explained + +| quantity | value | source | +|---|---|---| +| beam-256 GPU search | **238 s/layer** | 241 ± 1 s/layer marginal, minus 3.3 s host unpack | +| exhaustive Viterbi | 510 s/layer | same box class | +| speedup | **2.14×** | — | +| GPU telemetry during search | sm=100 %, **mem=1 %**, 261 W / 700 W, 1980 MHz | `nvidia-smi dmon`, 20 samples | +| wave13-AF projection | 42–85 s/layer (6–12×) | wrong | + +`mem=1%` is the most informative number on that list. It **confirms the design goal was achieved**: the beam really did +cut HBM traffic ~257× per symbol position, and the kernel really is no longer touching DRAM. It also means the kernel's +remaining cost is entirely on-chip, so no amount of further memory work can help. + +`sm=100 %` at 261 W is not a contradiction: `sm%` reports the fraction of time at least one warp is *resident*, not the +fraction of issue slots used. A kernel that occupies every SM and issues one instruction every twenty cycles reads +exactly like this. + +## The numbers the source cannot tell you, measured on CPU + +Two quantities drive the kernel's cost and neither is visible in the CUDA source, because both depend on the data. Both +are properties of the **beam search**, not of CUDA, so both are measurable on a CPU. New probe: + +``` +cargo test -p mistralrs-quant --lib probe_beam_kernel_cost_drivers -- --ignored --nocapture +``` + +It replays the production CPU beam on V4-Flash-shaped rows (FP4-lattice weights, hadamard-128, W=256) and transcribes +the kernel's own radix loop, early exit included. + +| shape | steps | mean `ng` | min | max | mean radix passes | wasted leading passes | +|---|---|---|---|---|---|---| +| gate/up `k=7168` | 14 332 | **248.7 / 256** | 16 | 256 | **3.87** | 0.83 | +| down `k=2048` | 4 092 | **248.6 / 256** | 16 | 256 | 3.61 | 0.62 | + +Radix-pass histogram (gate/up): `3 → 2133, 4 → 11966, 5 → 227, 6 → 2`. + +## Hypotheses ruled OUT + +**1. Launch geometry / not enough blocks — RULED OUT (source).** +`rows_in_flight = ARC_VITERBI_SCRATCH_GB / (num_symbols · W · 4)`. At the 6 GiB default, `k=7168`, W=256, that is +6 GiB / 3.67 MB = **1755 blocks per launch**, against 132 SMs. The 3-D expert path further batches 16 experts into one +`[16·2048, 7168]` call = 32 768 rows. The grid is never the constraint. + +**2. Thread-mapping starvation — RULED OUT (measured, and this was my leading hypothesis).** +`qtip_beam.cu` maps one thread per prefix group (`active = tid < ng`), so I expected `ng ≪ 256` to leave most of the +block idling at a barrier. **Measured `ng` = 248.7 of 256, i.e. 97.2 % of the block has work.** The mapping is fine. +Had I not measured this I would have shipped a rewrite that changed nothing. + +**3. Sequential dependence across timesteps — RULED OUT as *the* cause.** +Real, but it is not what is being wasted: parallelism comes from rows, and there are 1755 concurrent rows. + +**4. "Skip the provably-common leading radix digits" — REFUTED (measured).** +The obvious optimisation. Costs are cumulative, so I expected 2–3 leading digits to carry no information. Measured +waste is only **0.83 of 3.87 passes**, and starting the scan at the first differing bit makes it **worse — 4.76 passes, +not 3.87** — because misaligning the digit boundaries costs more passes than the skipped one saves. Refuted; not +shipped. + +**5. The 32 KiB prefix-group table capping occupancy — PARTLY RULED OUT.** +37 984 B static of the 48 KiB block budget. On sm_90 (228 KiB smem/SM) that permits **6 blocks/SM**, so shared memory is +not the binding limit. Registers are the suspect instead (below). And the table *is* needed: dedup by prefix group is +what makes the CUDA beam byte-identical to the CPU one. + +## What is actually dominant + +Instruction count per thread per timestep, counted from `qtip_beam.cu` with the measured pass count (3.87): + +| phase | instructions | barriers | +|---|---|---| +| group reduction (incl. 64-bit shared `atomicMin`) | ~40 | 5 | +| expansion — 16 candidates × (2 LUT loads + 5 FP ops + add) | ~130 | 0 | +| **radix select — 3.87 passes × 16 candidates × ~21** | **~1393** | **23** | +| compaction | ~190 | 3 | +| trace write | ~4 | 1 | +| **total** | **≈ 1757** | **33** | + +**The selection machinery is 79 % of the instructions and 70 % of the barriers.** The arithmetic it exists to select +over — the actual trellis branch metric — is 130 instructions, 7 %. The kernel spends ~110 instructions per candidate to +choose among candidates that cost 8 instructions to evaluate. + +Of the ~21 instructions per candidate per pass, roughly 13 are **64-bit** operations on the 48-bit `(cost, state)` key: +rebuilding it from `cand[j]` (6), the variable-shift prefix compare (4), the digit extract (3). Every one is two 32-bit +ops plus carry handling on the SM, and they are rebuilt from scratch on every pass. + +### The derivation that closes against 238 s + +Issue capacity: 132 SMs × 4 schedulers × 1.98 GHz = **1.045e12 warp-instructions/s**. +Work per layer: 284 B params / 44 layers / V=2 = **3.23e9 row-timesteps**, × 8 warps × 1757 instructions = +**4.54e13 warp-instructions**. + +* At 100 % issue efficiency that is **43.4 s/layer**. +* Measured 238 s ⇒ **achieved issue efficiency 18.2 %**, i.e. one instruction per warp every ~22 cycles. + +The same count for the exhaustive kernel (phase A 4096 prefixes × 16, phase B 65 536 states × ~10 ops ⇒ ~3350 +instructions/thread/timestep) gives 8.66e13 warp-instructions, and at its measured 510 s an efficiency of **16.2 %**. + +**Therefore:** + +``` +predicted speedup = (3350 / 1757) × (18.2 % / 16.2 %) = 1.91 × 1.12 = 2.14× +measured speedup = 510 / 238 = 2.14× +``` + +The account closes to within 2 %. The beam cut instructions 1.9× and issue efficiency barely moved. **The 257× HBM +saving bought essentially nothing, because at 16–18 % issue efficiency neither kernel was ever running near the +bandwidth roof** — the wave13-AF bandwidth model predicted 355 s against a measured 510 s and I read that 1.4× gap as +"70 % of peak, model validated". It is equally consistent with "latency-bound, and the bandwidth number was a +coincidence". I did not test the alternative. That is the actual error. + +### Why efficiency is 18 % + +18 % with 4 warps per scheduler means an average of ~22 stall cycles per issued instruction, from four sources that +compound and that low occupancy leaves fully exposed: + +1. **33 block barriers per timestep** — one per ~53 instructions. Every warp in the block stops at each. +2. **Scattered L2 reads** — each thread reads its own 128 B LUT run at a data-dependent `base_state`, so a warp issues + 32 separate transactions, not 4. ~103 TB/layer of L2 traffic (invisible to `mem%`, which reports DRAM). +3. **64-bit shared-memory `atomicMin`** in the group reduction, with up to 16-way address contention. +4. **Long dependent 64-bit chains** in the radix key arithmetic. + +None of these is individually fatal; all of them are fatal at **1–2 blocks per SM**, which is what I believe the +occupancy to be. `__launch_bounds__(256)` with no `minBlocksPerMultiprocessor` lets nvcc use up to 255 registers, and +the kernel holds `float cand[16]` live across the entire ~1400-instruction radix loop plus unrolled 64-bit temporaries. +At ~96 registers/thread the limit is 2 blocks/SM (25 % occupancy); above 128 it is 1 (12.5 %). + +**MEASURED (experiment #1, run on the build box — no GPU contention).** +`cuobjdump -res-usage`, sm_90a: `REG:80 STACK:0 SHARED:38992 LOCAL:0`. +256 threads × 80 registers = 20,480/block; 65,536 / 20,480 = **3 blocks/SM = 24 of 64 warps = 37.5 % occupancy**, +**register-limited** (shared memory alone would permit 5: 3 × 38,992 B = 114 KiB of 228 KiB). Independently checked: +the register-allocation granularity on sm_90 is 256 registers per warp and 32 × 80 = 2560 is an exact multiple, so no +rounding loss — 3 is exact, not a floor artefact. + +**The corroboration nobody planned:** the bake draws **261 W of 700 W = 37 % of TDP** against **37.5 % occupancy**. +A throughput-bound kernel would not have those two track each other; a latency-bound one at fixed clocks does. That is +independent evidence for the diagnosis, from a measurement taken for another purpose. + +So my stated downside case — "if it is already 4 blocks/SM, F2 is worth nothing" — **does not apply.** F2 is live. + +## Proven from source vs. measured vs. inferred + +| claim | status | +|---|---| +| grid = 1755 blocks/launch; grid is not the limit | **proven from source** | +| 33 barriers/timestep; 1757 instructions/thread/timestep | **proven from source** + measured pass count | +| `ng` = 248.7/256 ⇒ thread mapping is not the problem | **measured (CPU)** | +| 3.87 radix passes; skipping leading digits makes it worse | **measured (CPU)** | +| issue efficiency 18.2 %, and 1.91 × 1.12 = 2.14× | **derived** from source counts + measured wall times | +| registers cap occupancy at 1–2 blocks/SM | **inferred — needs `cuobjdump`** | +| stall attribution (barrier vs shared vs L2) | **inferred — needs `ncu`** | +| exhaustive kernel was not truly bandwidth-bound | **inferred — needs `dmon` on an exhaustive layer** | + +## Experiments, in the order to run them + +1. ~~**`cuobjdump -res-usage`, master**~~ — **DONE:** `REG:80 SHARED:38992 LOCAL:0` ⇒ 3 blocks/SM, 37.5 %, + register-limited. + ~~**`cuobjdump -res-usage`, wave16 build**~~ — **DONE:** `REG:64 SHARED:39024 LOCAL:0` ⇒ 4 blocks/SM, 50 %, no + spill. Neither needs re-running unless the kernel source changes; both are asserted in-tree. +2. **`ARC_QTIP_BEAM=64` vs `256` on one layer, wall time only.** W=64 does 4× fewer candidate evaluations but the *same* + 33 barriers and the same scan structure. If time falls ~3.5× the per-candidate work dominates; if it falls ~1.3× the + fixed per-step overhead dominates. One cheap A/B that discriminates the whole design. (Note: master's session6 change + pins W=256 for bakes — run this as a throwaway, not as an artifact.) +3. **`nvidia-smi dmon` during an *exhaustive* layer.** If `mem%` is high, the exhaustive kernel really was + bandwidth-bound and the beam traded one wall for another. If `mem%` is also ~1 %, my entire wave13 bandwidth model + was wrong at the root and should be deleted rather than patched. +4. **`ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active, smsp__issue_active.avg.pct_of_peak_sustained_active`** + on one launch ⇒ achieved occupancy and issue rate. Confirms or refutes 18.2 %. +5. **`ncu --metrics smsp__average_warps_issue_stalled_barrier.ratio, ..._short_scoreboard.ratio, ..._long_scoreboard.ratio`** + ⇒ attributes the stall to barriers / shared memory / L2. This is the one that tells us which of the four stall + sources to attack next. + +Experiments 1–3 need no profiler and answer most of it. + +## The fix (deliverable 2, branch `perf/beam-kernel-occupancy`, DRAFT PR #40, DO NOT MERGE UNMEASURED) + +https://github.com/aeonmindai/arc/pull/40 (DRAFT) — head `ac30f6e3a`, based on master `809643552`. Four commits: +`edcc158e4` the kernel changes, `9affb097d` firmed estimate + cost-model replacement, `ac30f6e3a` the occupancy +assertions. **CI: nvcc sm_80 and sm_90 both PASS on `ac30f6e3a`** — the compile gate for the rewritten kernel and for +the `#[cfg(feature = "cuda")]` parity tests. + +**Still DO NOT MERGE.** Occupancy is measured, parity is proved on CPU, the kernel compiles for both target arches — +but **no GPU has run this kernel**, so the wall-time claim (1.80–2.07×) remains unverified. Experiments 2 and 3 settle +that. + +Targets the two terms in `1.91 × 1.12`, since the analysis says both are the whole story. + +**F1 — 32-bit selection key with an exact tie fallback.** The measured pass histogram shows selection resolves inside +the 32-bit cost key in 98.4 % of steps (passes 5–6, which are the only ones reaching the state bits, occur 229/14 328 = +1.6 % of the time). Radix over the 32-bit cost key instead of the 48-bit composite: every shift, compare and digit +extract halves. When the final bin holds more than one candidate — genuine cost ties — resolve them by state with two +further 8-bit passes. **Exact; byte-identical.** + +**F2 — `__launch_bounds__(256, 4)`.** Caps registers at 64/thread and pins 4 blocks/SM (4 × 38 KiB = 152 KiB ≤ 228 KiB). +F1 removes the 64-bit temporaries that inflate the live set, which is what makes 64 registers plausible without spills. +No semantic change whatsoever. + +**F3 — barrier reduction, 33 → ~22.** Clear the histogram at the *end* of the previous pass (folding one barrier into an +existing one), and cut the block scan from three barriers to two. No arithmetic touched. + +### Firmed estimate (occupancy now measured, not inferred) + +The measured 3 blocks/SM pins the latency term. 24 warps over 4 schedulers = 6 warps/scheduler at 17.3 % implied +efficiency ⇒ **~34.7 cycles between issues per warp**. That single number now carries the whole projection: + +| term | value | basis | +|---|---|---| +| F1 instruction cut | 1664 → 1231 per thread per timestep = **1.35×** | source count, measured 3.87 passes | +| F2 occupancy | 6 → 8 warps/scheduler ⇒ 17.3 % → 23.1 % = **1.33×** | measured REG:80 ⇒ 3 blocks/SM | +| F3 barriers 33 → ~22 | **1.00–1.15×** | least certain; barriers are only part of the 34.7-cycle latency | + +**Combined 1.80–2.07× ⇒ 238 s → 115–132 s/layer ⇒ 3.9–4.4× vs the exhaustive 510 s.** + +Note this is **narrower and slightly lower** than the 2.0–2.8× I quoted before experiment #1 — because 3 blocks/SM is +*better* than the 1–2 I feared, so F2's headroom is 1.33× rather than up to 2×. Knowing the number shrank the claim. + +**The spill gate has been run and it PASSED.** Built from branch head `9affb097d` in a separate checkout, sm_90a: + +``` +REG:64 STACK:0 SHARED:39024 LOCAL:0 (wave16) +REG:80 STACK:0 SHARED:38992 LOCAL:0 (master) +``` + +REG 80 → 64 is exactly the F2 target, and **`LOCAL:0` means nvcc reached the budget without spilling** — the one +condition that could have inverted F2's gain (a spilled load in the radix loop runs 16 × 3.87 times per timestep) is not +triggered. Shared grew 32 B for the double-buffered scan scratch; 4 × 39,024 = 152 KiB of 228 KiB. **⇒ 4 blocks/SM, +50 % occupancy, up from 3 / 37.5 %.** So the 1.80–2.07× above rests on static evidence, not on an assumption. + +Caveat that matters: `LOCAL:0` is a property of the *exact source* that produced it. Any later edit to the radix loop or +the candidate registers re-opens it and the gate must be re-run — recorded at `QB_MIN_BLOCKS_PER_SM` in the kernel, with +the fallback (revert to 3, an exact no-op; expectation becomes F1+F3 alone, 1.35–1.55× ⇒ 153–176 s/layer). + +Both endpoints are now asserted in `cost_model_matches_the_measurement` — `blocks_per_sm(80, 38992) == 3`, +`blocks_per_sm(64, 39024) == 4`, and `blocks_per_sm(65, 39024) == 3` to record that 64 is a **cliff, not a slope** — so a +change that silently costs occupancy fails CI rather than a six-hour bake. + +**Parity: preserved, proved on CPU, and it is the gate.** `wave16_split_key_selection_matches_composite_key` replays +both scans on 3066 real candidate sets (9 exercising the tie-break) and asserts they admit the identical *set*, not +merely the same count. Always-on, not `#[ignore]`d. It also caught a real bug in the rewrite before it left the machine: +the early-exit path saturated only the state half of the threshold, not the unresolved low cost bits. + +**Parity, restated:** F1 changes which bits are examined in what order, never which candidate wins: +the selected set is defined by the total order on `(cost, state)` and the tie fallback reproduces it exactly. F2 and F3 +touch no arithmetic. `cuda_beam_matches_cpu_beam_bit_for_bit` remains the contract and must pass on hardware before this +branch goes anywhere near master. + +## Surfaced, not shipped + +* **The real headroom is algorithmic, not micro-optimisation.** 110 instructions per candidate to select among + candidates costing 8 to evaluate is a 14:1 overhead ratio. Even a perfect version of F1–F3 leaves selection dominant. + The structural fix is to stop doing an exact global top-W every step — e.g. a two-level scheme exploiting that the + top-256 of 4096 contains ~1 candidate per thread. Every such scheme risks the byte-identity contract, so it needs its + own parity argument, not a performance argument. +* **`ARC_VITERBI_SCRATCH_GB` now over-provisions.** The beam needs 3.67 MB/row against the exhaustive kernel's ~38 MB; + the 6 GiB default yields 1755 rows in flight when a few hundred would saturate the machine. Harmless, but the trace + buffer is now the only large allocation and could be sized from SM count instead. +* ~~wave13-AF's `beam_bake_cost_model` predicted 42–85 s/layer~~ — **REPLACED.** The predictive machinery + (`SymbolCost`, `HBM_EFFICIENCY`, `BEAM_ISSUE_EFFICIENCY`, `beam_bake_cost_model`) is gone. In its place is a + *descriptive* model anchored to the measurements: `MEASURED_BEAM256_LAYER_SECONDS`, `MEASURED_EXHAUSTIVE_LAYER_SECONDS`, + `MEASURED_BEAM_REGISTERS_PER_THREAD`, plus `implied_issue_efficiency` and `blocks_per_sm`. It predicts no wall time. + `cost_model_matches_the_measurement` asserts the instruction ratio × efficiency ratio reproduces the measured 2.14× + to within 5 %, that selection stays >70 % of the step, that both kernels stay below 25 % issue efficiency, and that + `blocks_per_sm(80, 38992) == 3` — i.e. it fails loudly if the kernel changes without re-measurement. + Re-parameterising the old model would have preserved its false authority; deleting the prediction was the point. + +--- + +# wave17-AF — MEASURED on rented silicon. The stack is 1.21x, and the beam architecture tops out at 3.5x. + +**Everything below is measured on an A30 (sm_80) I rented for the purpose, ~1 hour, $0.39.** No H200 was touched +(the bake was live). Absolute times are A30 numbers and do NOT transfer to H200; **ratios, occupancy and byte-identity +do**, which is exactly what was needed. + +Harness: `arc-tools/bench/beam_bench.cu` — a standalone driver that includes the kernel, generates weights and the +Gaussian LUT with the SAME deterministic hashes the Rust side uses (so the candidate cost distribution, and therefore +the data-dependent radix pass count, matches a real bake), times with CUDA events, and dumps packed output for a +byte comparison. Build both variants against one driver on one box; the ratio is the answer. + +## Result 1 — parity holds on hardware + +`base` and the wave16+17 stack produce **byte-identical packed output** (fnv1a `a655e47d7434c97e`, 458,752 bytes) at +every shape and width tested. The CPU equivalence proof and the hardware run agree. + +## Result 2 — the stack is 1.21x, not the 1.80–2.07x I projected + +rows=1344, k_in=7168 (V4 gate/up), best-of-2: + +| W | base | stack | ratio | +|---|---|---|---| +| 64 | 761.9 ms | 473.4 ms | **1.61x** | +| 256 | 1227.5 ms | 1011.1 ms | **1.21x** | + +`cuobjdump` on the A30 toolchain: base **REG:110 → 2 blocks/SM**, stack **REG:59 → 4 blocks/SM**. **Occupancy +DOUBLED and the kernel got 21% faster.** My wave16 model said occupancy was the multiplier on everything. It is not. +That model is now falsified by direct measurement, not by argument. + +## Result 3 — the W sweep decomposes the cost, and this is the finding that matters + +| W | 16 | 32 | 64 | 128 | 256 | +|---|---|---|---|---|---| +| stack (ms) | 395.2 | 410.9 | 473.4 | 616.7 | 1011.1 | + +Fitting `time = fixed + b·W`: + +| | fixed per-timestep | per-candidate at W=256 | +|---|---|---| +| base | 607 ms (49%) | 621 ms | +| stack | **354 ms (35%)** | **657 ms** | +| improvement | **1.71x** | **0.95x — none** | + +**My whole stack cut the fixed per-timestep overhead 1.71x (barriers 33 → 15.7, occupancy 2 → 4 blocks/SM) and did +essentially nothing to the per-candidate work.** That is coherent: F1/G3 removed key rebuilds, G1 removed barriers, +F2 removed a register cliff — all fixed-term costs. The per-candidate cost is dominated by the radix histogram +atomics and the scattered LUT loads, which none of those touched. + +## Result 4 — the architectural ceiling, measured + +The fixed term does not go away by making selection cheaper. **Even with per-candidate work driven to ZERO, this +kernel cannot beat 1227.5 / 354 = 3.47x.** + +Translated to the H200 measurement (225.2 s kernel + ~16 s host): + +| | kernel | + host | vs the 60–65 s/layer ask | +|---|---|---|---| +| banked today (1.21x) | 186 s | 202 s | miss | +| every remaining identity-safe idea lands perfectly (3.47x) | 65 s | **81 s** | **still a miss** | + +**45–50 s/layer of kernel with byte-identical output is not reachable from this beam architecture.** That is a +measured statement, not a projection: the fixed floor was measured directly at W=16. + +Two things DO reach it, both category iii-b (they move the CPU reference too, so parity survives byte-for-byte and +only quality changes): +* **W=128** — measured 616.7 ms = **1.99x vs base**, and PR #29 measured its quality cost at −0.004 cos (0.96084 vs + 0.96495). On the H200 that is ~113 s kernel + 16 s host ≈ **129 s/layer**. Still short of 60–65. +* **The gmin-only exhaustive Viterbi** (wave17-AN §5). It deletes the per-candidate selection term entirely — which is + 62% of W=256 time — AND has 1–2 barriers per position against our 15.7, so it attacks both terms at once. It is the + only structure in evidence that can reach the target. + +## What I would do next, in order + +1. **Dr. Top-k tight-delegate prefilter.** Measured on CPU: the W-th smallest of the β=2 delegate vector is a provable + upper bound leaving **318 survivors of 3980 (12.5x)**; the loose max-of-delegates form leaves 1677 (2.4x) and is not + worth building. It attacks the 62% per-candidate term. Requires shared-memory compaction — without it, SIMT + divergence means a warp pays full price whenever any of its 32 lanes survives (93% of iterations at 8% survival). + Realistic: ~1.5x more, landing ~1.8x total. **Does not reach the target.** +2. **Prototype the gmin-only exhaustive kernel** against the existing `cuda_exhaustive_matches_cpu_exhaustive_bit_for_bit` + gate. This is now the recommended path, on measured grounds rather than on AN's inference alone. + +## Measured negatives worth keeping + +* Occupancy 2 → 4 blocks/SM bought ~nothing on its own. Do not spend on occupancy for this kernel again. +* Guess-Verify-Refine's temporal premise **fails on our distribution**: 0.20% exact hits, and the extrapolated guess + lands BELOW the true threshold 58% of the time (the expensive miss). Our costs are cumulative, so the threshold + drifts every step. AN's ranking of GVR as the top lever was based on a premise that does not hold here — the + *delegate-bound* half of the idea survives, the *temporal warm-start* half does not. +* Skipping provably-common leading radix digits makes it worse (4.76 passes vs 3.87). + +## Cost + +One A30, ~1 hour, **$0.39**. Deleted; `list_instances` shows only the live H200. + +--- + +# CORRECTION (wave17-AF, second pass) — the 1.21× is an UPPER BOUND, not a verified value + +Jish challenged the 1.21×. He was right to. Re-verification was attempted on rented silicon and **did not complete** +(two A100 rentals hung in `deploying`, ~13 min each, never got an IP; no A30 capacity was available). What follows +separates what is now established from what is still unmeasured. + +## What IS established + +**1. The A/B compared the right two things.** Verified by hash, not by memory: `c5384fbcd` (the #40 merge) has first +parent `809643552`, whose kernel is md5 `a435ee5b…` = the "base" I built; current `origin/master`'s kernel is md5 +`dcd172c3…` = `d2d4f7e78`'s = the "treatment" I built; and no commit has touched `kernels/qtip/` since the merge. So +the comparison was of exactly the merged change. **That part of the report stands.** + +**2. Registers do NOT differ by card — that framing is wrong, and it was mine.** A30 and A100 are **both sm_80**; +for a given toolkit their codegen and `-res-usage` are identical. The 110/59 (my box) vs 80/64 (the H200 build) gap +has two non-card causes: **nvcc version** (mine was Ubuntu-apt **11.5**, ~4 years old) and **target arch** +(sm_80 vs sm_90a). Nothing about "A30 vs A100" enters it. + +**3. Therefore 1.21× is an upper bound for the shipped build, and the direction is provable.** On my box the +*baseline* was register-starved at **REG:110 → 2 blocks/SM**; the shipped build's baseline is **REG:80 → 3 blocks/SM**. +So my measurement handed F2 **more** occupancy headroom than production has (2→4 = 2× vs 3→4 = 1.33×) — and still +produced only 1.21× overall. Giving a change *more* room than it really has and measuring 1.21× means the shipped +configuration's gain is **≤ 1.21×**, not more. + +⇒ **The record should say: "≤1.21×, measured on sm_80 with nvcc 11.5 where the baseline was register-starved; not +verified for the shipped toolchain."** It should NOT say "1.21× measured". If FACTS.md and PR #43 carry the bare +number, they are overstating it. + +## What is NOT established, and I will not imply otherwise + +* **The F1 / F2 / F3 decomposition was never measured.** I planned the toggles (F2 is a one-line `sed` on + `__launch_bounds__`; F1 and F3 need real reverts) and never ran them. **Whether any of the three is a no-op or a + regression is an open question** — which matters, because two other changes tonight already measured as no-ops. +* **No A100 or modern-toolkit number exists.** The 1.21× has never been reproduced on the toolchain that builds the bake. +* The earlier "3.47× architectural ceiling" inherits the same caveat: it was fit to the same nvcc-11.5 W-sweep. The + *shape* of the argument (a fixed per-timestep floor exists, measured directly at W=16) is unaffected; the *number* + is toolchain-specific. + +## The cheapest authoritative test — and it does not need me or a rental + +The bake box already measures **373.6 s/layer WITH the stack** at `897534e77`. Build that same box at `809643552` +(the merge's first parent) and time one layer. If 1.21× held, pre-stack should be ~450 s/layer. Same silicon, same +toolkit, same data, same driver — strictly better evidence than any synthetic bench I can rent, and it is one build +plus one layer. + +If that A/B comes back ~1.0×, the honest conclusion is that the stack does nothing on the shipped configuration and +#40 should be reverted or re-scoped. + +## Cost and hygiene + +A30 (first pass) ~1 h $0.39; two A100 rentals that never provisioned, deleted. **`list_instances` shows only the live +bake box `d7d5d4ba…`, which was never touched.** diff --git a/memory/mission/wave45-BW-mtp-batched.md b/memory/mission/wave45-BW-mtp-batched.md new file mode 100644 index 000000000..297da5b4e --- /dev/null +++ b/memory/mission/wave45-BW-mtp-batched.md @@ -0,0 +1,223 @@ +# wave45-BW — MTP at every batch size, and the two forwards that halved it + +Branch `feat/mtp-batched`, stacked on `feat/mtp-working` (PR #71). Scope fence +respected: no `mistralrs-quant/src/qtip/**`, no KV/`xs` byte-layout, no +`device_map.rs` / `distributed/**`. + +--- + +## 1. Why `take_fast_path` was B=1-gated — REAL, and about the cache, not MTP + +`mtp_pipeline.rs`, old gate: `input_seqs.len() == 1`. The comment blamed "no +batched draft path in Tier A". That is not the constraint. The constraint is +one level down: + +* `kv_cache/single_cache.rs:161` — `SingleCache::append` does + `ad.slice_set(src, self.dim, self.current_seq_len)`. **One offset for the + whole batch.** +* `kv_cache/mod.rs:390-604` — `clone_in_cache` builds ONE dense batched cache + and takes `seqs[0]` as the template for `current_seq_len`. +* `kv_cache/mod.rs:353` — `first_mismatched_cache_len` exists precisely to + catch two sequences at different lengths in one batch, and says the failure is + silent (`CACHE_GROW_SIZE` is 512, so the shapes still match). +* `scheduler/default_scheduler.rs:115-157` — exactly **one** length-bucket runs + per step. + +So: speculative decode makes per-sequence accept lengths diverge; a dense +shared-offset cache cannot represent per-sequence lengths; and if you let token +lengths diverge, the scheduler shatters the cohort into one bucket per accept +length and idles all but one. **Genuine constraint, not laziness — but a +constraint of the dense cache, which is code, not physics.** + +## 2. The fix: fused single-forward + a fixed window with an uncached tail + +Two changes, and they are inseparable. + +**(a) Fusion.** The old shape ran **TWO** target forwards per step: one 1-token +decode to materialise `T0` and its hidden state, then the verify forward. That +caps the multiplier at `(depth+1)/2` — **at depth 1 that is exactly 1.0, i.e. +`--mtp-depth 1` could never have been faster than no MTP at all.** The seed the +next chain needs (`h_{L-2}`) is always inside the window the last verify already +covered — on a rejection at slot `j` the committed length is `C+j+1` so +`L-2 = C+j-1`; on a full accept it is the window's last input. Carrying it on +`DraftKv::seed` removes the first forward outright. SGLang reaches the same +place from the other side, skipping the *last* draft forward because the +draft-extend it must run anyway already produced that token +(`eagle_worker.py:871-873`). + +**(b) The window.** Cache length `C` stays **uniform** across the batch; token +lengths diverge instead. Sequence `i` carries an uncached committed tail +`u_i ∈ [1, w]`, `w = depth + 1`, and feeds a fixed-width window of `u_i` real +tokens then `w - u_i` drafts. After verify the shared cache rolls back to +`C + min_i(u_i + accepted_i)` — the one length every sequence can prove — and +each sequence's surplus stays committed as *tokens*, returning next step as a +longer tail, which is exactly what `w - u` drafts leaves room for. + +The invariant closes: `u ≥ 1` ⇒ the cache advances every step (no stall), and +`u_i + a_i ≤ w` ⇒ `u' ≤ w + 1 - min ≤ w` (a tail can never outgrow the window). +Pinned by `batched_window_invariant_closes_at_every_depth_and_accept_pattern` +over >5,000 (depth, accept-pattern) combinations. + +**(c) Scheduler.** `BucketKey` now keys on `Sequence::cache_bucket_len()` (the +cache length) rather than `len()`. For every non-speculative path this is the +*identical partition* — `cache_len ≡ len - 1` — so nothing else moves. It is +what keeps a batched MTP cohort whole. + +**(d) Batched drafting.** Within a group of equal `u`, every sequence sits at the +same absolute chain start `C + u - 2` and needs the same `w - u` drafts, so one +`[G, 1, hidden]` MTP-block forward per chain step drafts for all of them. `u` +ranges over `[1, w]`, so a batch of any size splits into at most `depth + 1` +groups. Per-sequence draft KVs are concatenated into one batched cache +(`batch_draft_caches`) and split back (`split_draft_cache`) — the MTP block is a +single decoder layer, so this is `clone_in_cache` at ~1/43 of its cost. + +## 3. PagedAttention exclusivity — **unimplemented branch, and moot for V4** + +Not a genuine incompatibility. One thing is missing: the rollback. +`truncate_cache_by` walks a `NormalCache` and calls `set_len` per layer; under +paging the K/V lives in the block table the PagedAttention KV manager owns, the +`EitherCache::Normal` entries are unused, and that truncate would be a **silent +no-op** leaving a rejected draft's K/V addressable. Refusing is correct until +the paged free-list step exists. + +The reference proves it is the cheap case. SGLang frees exactly the complement +of its accept set (`eagle_info.py:488-490`), and at **topk == 1** — a linear +chain, which is what DeepSeek MTP *is* in SGLang's shipped config +(`server_args.py:7611-7627`: `(3, 1, 4)` for every DeepSeek arch) — the accepted +tokens are a contiguous prefix of the allocated run, so rollback is **pure +truncation, no KV movement** (`eagle_info.py:492-501`). Only tree drafting needs +the `move_kv_cache` compaction (`:505-545`). Draft-phase slots are never leaked +at all: they are allocated against a backed-up allocator state that is +unconditionally restored (`eagle_worker.py:613-617` → `:727`, +`allocator.py:71-75`). + +Further — and this inverts the framing in `CEILINGS.json` — **paged KV is a +*better* substrate for batched MTP than the dense cache**, because a per-sequence +block table makes per-sequence lengths free and the entire uncached-tail scheme +above unnecessary. That is why vLLM/SGLang can do batched spec decode at all. + +**Moot for V4, twice over.** `DeepSeekV4Loader::supports_paged_attention` returns +`false` (`loaders/normal_loaders.rs`, wave29-BC rationale: `flashinfer_mla_decode.cu` +fixes `HEAD_DIM_CKV=512` as a template constant and computes dense causal +attention while every V4 layer is sliding-window + sink), so the engine never +hands V4 a `CacheBackendMetadata::PagedAttention`. And `mtp_decode_kit` has +**exactly one** implementation in the tree (`deepseek4.rs:4250`; the trait +default at `loaders/normal_loaders.rs:104` returns `None`), so the MTP wrapper +cannot currently wrap a model that pages. **The two features have never met and +this guard has never fired.** `--paged-attn off` is still required for V4 MTP — +but for the unrelated reason that `auto` = ON on CUDA would put the engine on a +backend V4 does not support. + +## 4. "Best MTP" — ranked, one implemented, the rest scoped with numbers + +| rank | item | expected gain | status | +|---|---|---|---| +| 1 | **Fuse the two target forwards** | **×2 on the multiplier at every depth**; depth 1 goes from 1.00 to 2.00 max | **implemented** | +| 2 | **Batched fast path** | MTP applies at B>1 at all — the difference between a single-user lever and a fleet lever | **implemented** | +| 3 | Depth 2→3 | reference accept-length 2.30 (non-simulated) at their `num_steps=3`; ours is `depth`-configurable already, clap-capped at 8 vs `XS_TAIL_MARGIN_TOKENS=16` | **already available**, default still 0 | +| 4 | Tree / multi-candidate drafting | **NOT recommended for V4.** SGLang ships `topk=1` (linear chain) for *every* DeepSeek arch (`server_args.py:7611-7627`); tree is the Llama/Grok default `(5,4,8)`. Tree also forces the `move_kv_cache` compaction on paged rollback and is unsupported by their overlap scheduler | scoped, **declined** | +| 5 | EAGLE-3 | **wrong target for V4.** Needs weights the DeepSeek MTP checkpoint does not ship: `fc` (3H→H aggregation), `fc_norm.{0,1,2}`, `d2t`/`t2d` vocab maps, and its own `lm_head` over `draft_vocab_size` (`llama_eagle3.py:148-173,277-327`); plus a *target-side* 3-layer capture (`deepseek_v2.py:2567-2574`) and a separate SpecForge training run. MTP ships in the base checkpoint | scoped, **declined for V4** | +| 6 | Stochastic-sampling verification (rejection sampling) | today acceptance is greedy-argmax equality; the step's one non-drafted token does go through the real sampler | scoped, backlog | + +## 5. What the reference does that we do not + +* **Ragged verify as a dense `(B, steps+1)` `-1`-padded `accept_index` matrix** + plus a `num_correct_drafts` vector, resolved by one CUDA kernel with one block + per sequence (`eagle_info.py:277-281`, `eagle_utils.cu:290-311`). We resolve + on the host in a `B`-length loop — fine at our step times, a target later. +* **Per-sequence KV freeing** instead of a batch-minimum rollback + (`eagle_info.py:488-501`). Ours costs the surplus as a longer tail next step; + theirs costs nothing because the KV is paged. +* **`backup_state`/`restore_state`** on the allocator so draft KV is borrowed, + never leaked (`allocator.py:71-75`). +* **Adaptive depth** from an EMA of accept length + (`docs/advanced_features/adaptive_speculative_decoding.md:91`). +* **A stated acceptance floor in CI.** DeepSeek V4 MTP: `acc_length > 2.30` + non-simulated (`test/manual/dsv4/test_dsv4_pro_mtp.py:261`), `> 2.85` with + `SGLANG_SIMULATE_ACC_LEN=3`. V3: `> 2.8` (`test_deepseek_v3_mtp.py:88`). + +## 6. Tests + +* `batched_window_invariant_closes_at_every_depth_and_accept_pattern` — >5,000 + patterns, depths 1-8: `keep ≥ 1`, `next_uncached ∈ [1, w]`. +* `batched_ragged_accept_is_token_identical_to_the_b1_reference` — five + sequences with different agreement rates decode together and alone; the + per-sequence token streams must match exactly. **Asserts its own + non-degeneracy first** (D12): the run must contain a step with differing + accept lengths *and* a step where some tail exceeded 1, or the fixture proves + nothing. +* `batched_rollback_mutation_max_instead_of_min_is_caught` — resolving on the + MAXIMUM valid extent (the plausible "keep as much cache as possible" mistake) + drives a tail to 0, which is a sequence whose committed tokens sit in the + cache as K/V computed from a rejected draft. +* `window_verify_row_off_by_one_destroys_acceptance` — a one-row shift is the + invisible defect (no crash, no corruption, just collapsed acceptance). +* `batched_acceptance_is_reported_per_batch_size` — all-accept reads 1.0, + all-reject reads 0.0, per batch size; per-user and aggregate multipliers are + both on the marker. +* Everything from PR #71 still passes: 296 + 13 + 2 tests green. + +## 6b. A latent prefix-cache over-claim the ragged tail exposed (fixed) + +`search_for_matching_cache` calls `set_len(match_len)` on the stored layers, and +`SingleCache::try_set_len` checks **capacity, not validity** — it extends a +length into slots that were never written. Plain decode never trips it because +two facts cancel exactly: `cache_len == len - 1`, and a full-length match is +rejected by the `new_toks.is_empty()` guard, so `match_len ≤ len - 1 == +cache_len`. Safe by arithmetic coincidence, not by design. + +Batched MTP breaks the coincidence: a sequence can finish with +`cache_len == len - u`, `u` up to `depth + 1`, while `match_len` can still reach +`len - 1` — a later request sharing that prefix would be served uninitialised +K/V. Fixed twice over: `add_sequence` keys by `toks[..min(len, cache_len + 1)]` +(a **no-op when `u == 1`**, so no existing path changes), and retrieval declines +any match longer than `layer.current_seq_len()`. + +## 6c. Batched prefill primed only row 0 — batched MTP would have been dead on arrival + +The prompt path seeded the draft KV only when `input_seqs.len() == 1`. In a real +serve that is the common case *reversed*: `batch_load_probe.py --batches 128` +makes 128 requests arrive together, they prefill together, and under the old code +**not one of them would have had a seed** — so every decode step would have +declined to draft, losslessly and silently, at exactly zero speedup. A batched +MTP that never drafts is indistinguishable in a log from a batched MTP that +works badly, which is how this class of thing survives a GPU session. + +Now every row is primed. `make_prompt_chunk` right-pads to the batch's longest +prompt, so row `i`'s real positions are `[0, len_i)`, and +`extend_draft_kv_row`'s own `i ≤ L-2` bound already stops there. + +## 7. Measurement + +`MTP[agg] …` plus one `MTP[b=] …` line per observed batch size, each +carrying `tok_per_step` (**per user** — the number that multiplies the ceiling), +`tok_per_batch_step` (**aggregate**), and `mean_batch`. `mean_batch` collapsing +toward 1 while the load probe reports a large batch is the tell that the cohort +fragmented. Recipe: `RUNBOOK_8 §S3c` (new). `bench` still forces +`--max-seqs 1`, so a bench run remains a **B=1 diagnostic row only** (D2). + +## 8. Projection (LABELLED PROJECTION — nothing here is measured) + +`CEILINGS.json` per-user ceilings: **83 tok/s at B=64, 68 at B=128.** MTP emits +`k` tokens per step without reading more bytes, so it multiplies straight +through. Draft overhead is `depth` single-layer forwards against a 43-layer +target ⇒ `1 + depth/43`. + +| | k=2 (depth 2) | k=3 (depth 3) | +|---|---|---| +| effective multiplier | 2 / (1+2/43) = **1.91** | 3 / (1+3/43) = **2.80** | +| B=64 per user | 83 × 1.91 = **159** | 83 × 2.80 = **232** | +| B=128 per user | 68 × 1.91 = **130** | 68 × 2.80 = **190** | + +**k=2 clears Jish's 100-tok/s-at-any-batch target on a single card.** The +reference's own non-simulated floor for V4 MTP is k = 2.30 +(`test_dsv4_pro_mtp.py:261`), which lands between these columns. Measured +acceptance remains **UNMEASURED** on our stack — no number is claimed here. + +## 9. Corrections owed to CEILINGS.json + +`levers_that_BEAT_the_table.MTP_speculative_decode` currently says the fast path +is B=1 only and that "MTP and PagedAttention are mutually exclusive". Both +should now read: the fast path runs at every batch size (wave45-BW); the +PagedAttention conflict is an unimplemented rollback branch that is **moot for +V4** because V4 does not page and V4 is the only model with an MTP head. diff --git a/mistralrs-core/src/paged_attention/segmented_allocator.rs b/mistralrs-core/src/paged_attention/segmented_allocator.rs index 3183a036a..cd5275dfa 100644 --- a/mistralrs-core/src/paged_attention/segmented_allocator.rs +++ b/mistralrs-core/src/paged_attention/segmented_allocator.rs @@ -12,10 +12,11 @@ //! [`Self::rollback`] name a `(request_id, segment_index)` pair. No operation //! reads another request's table, so no request's progress can be pinned by //! another's. That is the allocator-level precondition for lifting the MTP -//! batch ceiling recorded in `memory/mission/wave59-CJ-mtp-batch.md` §3, where -//! one dense shared cache forces `next_u = u + a + 1 - min_i(u_i + a_i)` and a -//! single sequence rejecting its first draft ratchets every other sequence's -//! uncached tail to the window. **It is a precondition, not the fix** — the +//! batch ceiling recorded in `memory/mission/wave45-BW-mtp-batched.md` §2, +//! where one dense shared cache rolls back to `C + min_i(u_i + accepted_i)` — +//! the one length every sequence can prove — so a single sequence rejecting its +//! first draft ratchets every other sequence's uncached tail to the window. +//! **It is a precondition, not the fix** — the //! fused MTP step still has to stop rolling one shared cache. See //! `pipeline/mtp_pipeline.rs`. //! diff --git a/mistralrs-core/tests/doc_citations.rs b/mistralrs-core/tests/doc_citations.rs index 8b585be82..e4204e8cc 100644 --- a/mistralrs-core/tests/doc_citations.rs +++ b/mistralrs-core/tests/doc_citations.rs @@ -1363,6 +1363,112 @@ static BASELINE: &[Waiver] = &[ why: "the QTIP reference is an upstream reference checkout, not \ vendored here", }, + // -- Citations carried in by the six mission docs that were cited from + // Rust comments but had never been committed (they existed on one + // laptop). Same class as everything above: pointers into sglang, CUDA + // headers and DeepSeek's configs, which Arc reads and does not vendor. + Waiver { + doc: "memory/mission/BUDGET_V4_B1.md", + cite: "cuda.h", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/BUDGET_V4_B1.md", + cite: "cuda_device_runtime_api.h", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/EXTERNAL_FINDINGS.md", + cite: "indexer.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave13-AH-config.md", + cite: "configs/deepseek_v4.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave13-AH-config.md", + cite: "deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave13-AH-config.md", + cite: "deepseek_v4_rope.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "allocator.py", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "deepseek_v2.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "docs/advanced_features/adaptive_speculative_decoding.md", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "eagle_info.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "eagle_utils.cu", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "eagle_worker.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "llama_eagle3.py", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "server_args.py", + kind: Kind::Unresolved, + why: "sglang is an upstream reference checkout, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "test/manual/dsv4/test_dsv4_pro_mtp.py", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "test_deepseek_v3_mtp.py", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, + Waiver { + doc: "memory/mission/wave45-BW-mtp-batched.md", + cite: "test_dsv4_pro_mtp.py", + kind: Kind::Unresolved, + why: "cited from an upstream reference tree, not vendored here", + }, ]; // --------------------------------------------------------------------------- From c1c02a4cff636069f44b070ea4cf6019f69923b3 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:34:29 +0100 Subject: [PATCH 5/8] chore(ArcGate): waive the grep pattern `sampl` quoted in BUDGET_V4_B1.md The doc records, as evidence for "the GPU sampler is referenced nowhere", that searching for `sampl` returns zero hits. The search string is the evidence, so it cannot be spelled "sample". Waived with the backticks included, so this covers the quoted pattern and nothing else. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- .typos.toml | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/.typos.toml b/.typos.toml index 2932a4ac7..7de3b414e 100644 --- a/.typos.toml +++ b/.typos.toml @@ -33,6 +33,11 @@ extend-ignore-re = [ "mxfp4_get_max_smem_optin", "alog", "ba", + # A grep pattern quoted verbatim in BUDGET_V4_B1.md's evidence for "the + # GPU sampler is referenced nowhere": `sampl` is the search string, not a + # misspelling of "sample". Backticks included so this waives the quoted + # pattern and nothing else. + '`sampl`', ] [default.extend-words] From 1356f958411c4cf8b671816636ad2898bdba3aa3 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 03:39:11 +0100 Subject: [PATCH 6/8] perf(ArcInfer): the decode cohort crossed to the GPU one sequence at a time MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `make_completion_chunk` built one `Tensor::new(.., device)` per sequence and `Tensor::cat`-ed them. On CUDA that is B separate **1-element** H2D transfers, each a pageable copy on the legacy default stream and therefore a full host/device round trip, plus a B-argument `cat`, which candle services with one `copy2d` per argument. At B=256: **256 synchronous transfers and 256 copy kernels to assemble a `[256, 1]` tensor holding one kilobyte.** The call site already carried the note — "CLAUDE.md pitfall #5, still live" — and a `sync_span` named `input_prep.h2d_per_seq` measuring it. It was named, instrumented, and left running. Every decode row is exactly one token, so the cohort is a rectangle. Rows now stay on the host and `stack_uniform_rows` sends the whole thing in one transfer: **256 round trips and 256 launches become 1 and 0.** Ragged input is **refused, not reshaped**. The dangerous case is the one that still divides evenly — rows of 1, 2 and 3 tokens hold 6 elements, which reshapes cleanly to `[3, 2]` and hands sequence 0 a token belonging to sequence 1. The shape would be right and only the contents wrong, which no downstream assertion can see. It errors and names the offending row. Scoped deliberately to decode. The prefill path has the same shape at `make_prompt_chunk`, but its rows are padded by branching logic with cache-hit and chunking cases, its payload is large enough that the per-row round trip is a smaller share, and a hard error on a ragged batch there would turn working behaviour into a failure. Named, not touched. **UNVERIFIED ON HARDWARE.** The launch and transfer counts are structural — one per sequence, B sequences. No wall-clock claim. Three tests, on a seam extracted so they need no `Sequence` fixture: the single-transfer form is byte-identical to the per-row `cat` it replaces, on the `[B, 1]` decode shape and on a wider rectangle; a ragged cohort whose element count divides evenly is refused and the error names the row; an empty batch is an error rather than a `[0, 0]` tensor. cargo test -p mistralrs-core --lib: 652 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- .../src/pipeline/inputs_processor.rs | 136 ++++++++++++++++-- 1 file changed, 126 insertions(+), 10 deletions(-) diff --git a/mistralrs-core/src/pipeline/inputs_processor.rs b/mistralrs-core/src/pipeline/inputs_processor.rs index b50d8b3a8..a5f4e68f5 100644 --- a/mistralrs-core/src/pipeline/inputs_processor.rs +++ b/mistralrs-core/src/pipeline/inputs_processor.rs @@ -580,6 +580,50 @@ pub mod text_models_inputs_processor { }) } + /// Stack a rectangular cohort of token rows into `[B, L]` with **one** + /// host-to-device transfer. + /// + /// The decode path used to build one `Tensor::new(.., device)` per + /// sequence and `Tensor::cat` them. On a CUDA device that is B separate + /// **1-element** H2D transfers — each a pageable copy on the legacy + /// default stream, i.e. a full host/device round trip — plus a B-argument + /// `cat`, which candle services with one `copy2d` per argument. At B=256 + /// that is 256 synchronous transfers and 256 copy kernels to assemble a + /// `[256, 1]` tensor holding one kilobyte. The call site carried the + /// standing note that this is "CLAUDE.md pitfall #5, still live". + /// + /// Every decode row is exactly one token, so the cohort is a rectangle and + /// one flat host buffer reshaped to `[B, L]` is the identical tensor. + /// + /// **Ragged input is refused, not reshaped.** A wrong-width row would + /// otherwise reinterpret one sequence's tokens as another's — silent, + /// per-token, and invisible in any shape assertion downstream, since the + /// element count could still divide evenly. + pub(crate) fn stack_uniform_rows( + rows: Vec>, + device: &Device, + ) -> Result { + let n_rows = rows.len(); + if n_rows == 0 { + anyhow::bail!("stack_uniform_rows: empty batch — no sequences to stack"); + } + let cols = rows[0].len(); + if let Some(bad) = rows.iter().position(|r| r.len() != cols) { + anyhow::bail!( + "stack_uniform_rows: row {bad} has {} tokens but row 0 has {cols}. \ + This is the decode path, where every row is exactly one token — \ + a ragged cohort here means the caller changed and the flat \ + reshape would silently mix sequences.", + rows[bad].len() + ); + } + let mut flat = Vec::with_capacity(n_rows * cols); + for row in rows { + flat.extend_from_slice(&row); + } + Ok(Tensor::from_vec(flat, (n_rows, cols), device)?) + } + fn make_completion_chunk( toks: Vec<&[T]>, input_seqs: &[&mut Sequence], @@ -589,7 +633,8 @@ pub mod text_models_inputs_processor { ) -> Result { // Pad each sequence by the padding token to the max len. let flash_attn = crate::using_flash_attn(); - let mut seqs_tensors = Vec::new(); + // Host-side rows, stacked once at the end — see [`stack_uniform_rows`]. + let mut seq_rows: Vec> = Vec::with_capacity(input_seqs.len()); let mut seqlen_offsets = Vec::new(); let mut context_lens = Vec::new(); let mut position_ids = Vec::new(); @@ -613,13 +658,11 @@ pub mod text_models_inputs_processor { seqlens_k.push((ctxt.len() + start_pos) as u32); } - // CLAUDE.md pitfall #5, still live: one `Tensor::new` on the GPU - // device per sequence per decode step, i.e. B separate 1-element - // H2D transfers, each of which is a host/device round trip. - { - let _s = arc_profiler::sync_span("input_prep.h2d_per_seq"); - seqs_tensors.push(Tensor::new(ctxt, device).unwrap().unsqueeze(0).unwrap()); - } + // Was CLAUDE.md pitfall #5: one `Tensor::new` on the GPU device + // per sequence per decode step — B separate 1-element H2D + // transfers, each a host/device round trip. The row stays on the + // host here and the whole cohort crosses in one transfer below. + seq_rows.push(ctxt); if let Some(paged_attn_metadata) = &mut paged_attn_metadata { let kv_mgr = get_mut_arcmutex!(paged_attn_metadata.kv_cache_manager); @@ -896,8 +939,8 @@ pub mod text_models_inputs_processor { }; let input = { - let _s = arc_profiler::span("input_prep.cat"); - Tensor::cat(&seqs_tensors, 0).unwrap() + let _s = arc_profiler::sync_span("input_prep.h2d_cohort"); + stack_uniform_rows(seq_rows, device)? }; Ok(InputMetadata { input, @@ -1768,3 +1811,76 @@ mod cursor_reaches_every_path_tests { ); } } + +#[cfg(test)] +mod decode_cohort_crosses_once_tests { + use super::text_models_inputs_processor::stack_uniform_rows; + use candle_core::Device; + + /// **One transfer, and the same tensor as B transfers produced.** + /// + /// The decode path used to build one `Tensor::new` per sequence on the GPU + /// device and `cat` them: at B=256 that is 256 pageable 1-element H2D + /// transfers on the legacy default stream — 256 host/device round trips — + /// plus 256 `copy2d` launches, to assemble a `[256, 1]` tensor holding a + /// kilobyte. This pins that the single-transfer form is byte-identical to + /// the per-row form it replaces, on the `[B, 1]` decode shape and on a + /// wider rectangle. + #[test] + fn stacking_once_equals_stacking_row_by_row() -> candle_core::Result<()> { + let dev = Device::Cpu; + for cols in [1usize, 3] { + let rows: Vec> = (0..7u32) + .map(|r| (0..cols as u32).map(|c| r * 100 + c).collect()) + .collect(); + + // The shape this replaces: one tensor per row, then `cat`. + let per_row: Vec = rows + .iter() + .map(|r| { + candle_core::Tensor::new(r.clone(), &dev) + .unwrap() + .unsqueeze(0) + .unwrap() + }) + .collect(); + let want = candle_core::Tensor::cat(&per_row, 0)?; + + let got = stack_uniform_rows(rows, &dev).expect("uniform rows stack"); + assert_eq!(got.dims(), want.dims(), "shape differs at cols={cols}"); + assert_eq!( + got.flatten_all()?.to_vec1::()?, + want.flatten_all()?.to_vec1::()?, + "contents differ at cols={cols}" + ); + } + Ok(()) + } + + /// A ragged cohort must be **refused**, never reshaped. + /// + /// The dangerous case is the one that still divides evenly: rows of 1, 2 + /// and 3 tokens hold 6 elements, which reshapes cleanly to `[3, 2]` and + /// hands sequence 0 a token belonging to sequence 1. Nothing downstream + /// could see it — the shape is right, only the contents are somebody + /// else's. + #[test] + fn a_ragged_cohort_is_refused_even_when_the_count_divides() { + let dev = Device::Cpu; + let rows: Vec> = vec![vec![1], vec![2, 3], vec![4, 5, 6]]; + assert_eq!(rows.iter().map(|r| r.len()).sum::(), 6); + assert_eq!(6 % rows.len(), 0, "the fixture must divide evenly"); + let err = + stack_uniform_rows(rows, &dev).expect_err("a ragged cohort must not silently reshape"); + let msg = format!("{err}"); + assert!(msg.contains("row 1"), "the error must name the row: {msg}"); + } + + /// An empty batch is a caller bug, not a `[0, 0]` tensor. + #[test] + fn an_empty_batch_is_an_error() { + let dev = Device::Cpu; + let rows: Vec> = Vec::new(); + assert!(stack_uniform_rows(rows, &dev).is_err()); + } +} From 2fb4882ca7a282b989e1eeb818d9970d2ddaf526 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 04:09:03 +0100 Subject: [PATCH 7/8] =?UTF-8?q?fix(ArcKernels):=20exclude=20CPU=20from=20t?= =?UTF-8?q?he=20rank-3=20flatten=20=E2=80=94=20it=20is=20the=20oracle,=20a?= =?UTF-8?q?nd=20it=20moved?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CI caught what local `--lib` runs could not: the flatten broke `v4_xs_history_two_seq_batch_matches_single_sequence` (`synthetic_load_smoke.rs`) on ubuntu, and only on ubuntu. It is not a correctness bug and it is worth being precise about why, because the test's own message says "the xs-history / KV state of one sequence leaked into the other". It did not. Measured: max abs diff 0.0625 against a tolerance of 0.0399 — **1.6x over**, where that test documents a real leak as O(1), roughly **100x** over. This is floating-point associativity. The mechanism, established by probe rather than guessed. `MatMul::matmul` casts to F16 on CPU. Before the flatten, `broadcast_left` made the batched case B separate `m = 1` GEMMs — the identical shape, and therefore the identical accumulation order, as the solo run it is compared against. So batched-vs-solo was **bit identical**: instrumenting `v4c_assert_close` showed `max_diff` exactly 0 on every comparison. The 1e-2 tolerance had never been exercised in its life. Flattening makes the batched case `m = B`, and x86's F16 microkernel blocks `m = 2` differently from `m = 1`; aarch64 happens not to, which is why macOS stayed green and ubuntu did not. The fix is not to widen the tolerance. Widening a leak detector to admit a change made in the same PR is how a gate stops being a gate, and this one is load-bearing — the whole V4 batched-vs-solo suite is calibrated against CPU bit identity. So the flatten is now device-gated, and the boundary is the same one the derivation always had: **the defect is HBM traffic** — a stride-0 batched GEMM streaming a 1.06 GB weight once per batch row. There is no HBM on the CPU path, `MatMul::matmul` throws away precision there anyway, and CPU exists in this workspace to be a reference oracle, not a fast path. The fix lands where the 271 GB/step is; the oracle keeps its arithmetic. Tests restructured so none of them is now vacuous — the point of the gate is that CPU takes the old path, so a test asserting the new one on CPU would prove nothing: * `two_d_and_rank3_agree_within_the_cpu_f16_budget` — the mathematical claim, across four `[B, T]` splittings with and without bias, inside an F16 budget rather than at bit equality. The gap it tolerates is the gap that motivates the device gate; a wrong broadcast or a transposed weight is O(1) and still fails it. * `cpu_forward_still_takes_the_broadcast_path_bit_for_bit` — the oracle's arithmetic is pinned against the pre-change formulation spelled out literally, at zero tolerance. * `flatten_batch_dims_*` — the shape arithmetic the device path depends on, extracted into a named function so it is checkable on a host without a GPU. Rank-3 and rank-4 collapse; a non-contiguous activation round-trips through the strided copy `reshape` performs. Shape logic that only executes on hardware nobody has is shape logic nobody checks. cargo test -p mistralrs-quant --lib: 332 passed. cargo test -p mistralrs-core (all targets): 652 + 11 + 12 + 7 + 17 + 2 passed. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-quant/src/unquantized/mod.rs | 173 +++++++++++++++++++------ 1 file changed, 134 insertions(+), 39 deletions(-) diff --git a/mistralrs-quant/src/unquantized/mod.rs b/mistralrs-quant/src/unquantized/mod.rs index ed0b1fb71..0670bace4 100644 --- a/mistralrs-quant/src/unquantized/mod.rs +++ b/mistralrs-quant/src/unquantized/mod.rs @@ -18,6 +18,23 @@ use crate::{ QuantMethod, QuantMethodConfig, QuantizeOntoGuard, QuantizedSerde, QuantizedSerdeType, }; +/// Collapse every leading dimension into the GEMM's `m`, and return the dims +/// needed to restore them. +/// +/// `[B, T, hidden] -> ([B*T, hidden], [B, T])`. Kept as a named function so the +/// shape arithmetic is testable on any host — the caller that uses it is +/// device-gated (see [`UnquantLinear::forward`]), and shape logic that only +/// executes on hardware nobody has is shape logic nobody checks. +/// +/// `reshape` on a non-contiguous input costs one strided copy rather than +/// failing, which is what makes this safe to apply to an arbitrary activation. +fn flatten_batch_dims(a: &Tensor) -> Result<(Tensor, Vec)> { + let batch_dims = a.dims()[..a.rank() - 1].to_vec(); + let rows: usize = batch_dims.iter().product(); + let features = a.dim(D::Minus1)?; + Ok((a.reshape((rows, features))?, batch_dims)) +} + #[derive(Debug)] pub struct UnquantLinear { w: Tensor, @@ -111,11 +128,33 @@ impl QuantMethod for UnquantLinear { // // UNVERIFIED ON HARDWARE: derived from shapes and H200 peaks, not // measured. The arithmetic above is the claim; the speedup is not. - if a.rank() > 2 { - let batch_dims = a.dims()[..a.rank() - 1].to_vec(); - let rows: usize = batch_dims.iter().product(); - let features = a.dim(D::Minus1)?; - let out = self.forward(&a.reshape((rows, features))?)?; + // + // ── Why CPU is excluded, and it is not to dodge a test ────────────── + // + // The defect is HBM traffic: a stride-0 batched GEMM streaming a + // 1.06 GB weight once per batch row. There is no HBM on the CPU path, + // and `MatMul::matmul` casts to F16 there anyway — it is a reference + // oracle, not a performance path. + // + // Changing its arithmetic is not free. Batched-vs-solo was **bit + // identical** before this (`max_diff` measured at exactly 0), because + // `broadcast_left` made the batched case B separate `m = 1` GEMMs — + // the same shape, and therefore the same accumulation order, as the + // solo run. Flattening makes the batched case `m = B`, and x86's F16 + // microkernel blocks `m = 2` differently from `m = 1`. That is pure + // floating-point associativity, but it re-baselines every + // batched-vs-solo tolerance in the V4 suite at once: CI measured + // 1.57e-2 relative divergence against a 1e-2 tolerance that had never + // been exercised, in a test whose job is to catch cross-sequence state + // leaks — which are O(1), ~100x that tolerance. + // + // So the oracle stays byte-stable and the fix lands where the cost is. + // `two_d_and_rank3_agree_within_the_cpu_f16_budget` pins that the two + // formulations are the same mathematics, and measures the gap that + // motivates this line. + if a.rank() > 2 && !a.device().is_cpu() { + let (flat, batch_dims) = flatten_batch_dims(a)?; + let out = self.forward(&flat)?; let mut out_dims = batch_dims; out_dims.push(out.dim(D::Minus1)?); return out.reshape(out_dims); @@ -756,19 +795,36 @@ mod tests { Tensor::from_vec(v, (rows, inp), &Device::Cpu) } - /// **The flatten must not change the answer.** + fn max_abs_diff(a: &Tensor, b: &Tensor) -> Result { + Ok((a - b)?.abs()?.max_all()?.to_scalar::()?) + } + + /// **The mathematical claim behind the flatten, and the measurement of what + /// it costs in precision.** + /// + /// `forward` reshapes `[B, T, hidden]` to `[B*T, hidden]` so the weight is + /// streamed once instead of once per batch row. That is only sound if the + /// two formulations are the same mathematics. They are — but they are not + /// the same *arithmetic*: `MatMul::matmul` casts to F16 on CPU, and a + /// GEMM's accumulation order depends on `m`, so `m = B` and B separate + /// `m = 1` calls round differently. /// - /// `forward` now reshapes any rank-3 activation to `[B*T, hidden]` so the - /// weight is streamed once instead of once per batch row. That is a - /// dispatch change on the hottest path in the engine, so what it owes is - /// exact agreement with the 2-D result it is derived from — element for - /// element, not "close". + /// This test asserts agreement inside an F16 budget rather than bit + /// equality, and that distinction is the reason `forward` excludes CPU: on + /// x86 the divergence measured 1.57e-2 relative in a V4 batched-vs-solo + /// test whose tolerance was 1e-2 — a tolerance that had never been + /// exercised, because before the flatten the two paths were bit identical. #[test] - fn rank3_forward_matches_the_flattened_2d_result() -> Result<()> { + fn two_d_and_rank3_agree_within_the_cpu_f16_budget() -> Result<()> { + // Relative to the magnitude present, generous enough for F16 + // re-blocking and far tighter than any structural error (a wrong + // broadcast or a transposed weight is O(1) here). + const F16_BUDGET: f32 = 1e-2; for bias in [false, true] { let l = lin(7, 5, bias)?; let flat = act(6, 5)?; let want = l.forward(&flat)?; + let scale = want.abs()?.max_all()?.to_scalar::()?.max(1.0); for shape in [(6usize, 1usize), (3, 2), (2, 3), (1, 6)] { let a = flat.reshape((shape.0, shape.1, 5))?; @@ -776,50 +832,89 @@ mod tests { assert_eq!( got.dims(), &[shape.0, shape.1, 7], - "batch dims must survive the round trip (bias={bias})" + "batch dims must survive (bias={bias})" ); - let diff = (got.reshape((6, 7))? - &want)?.abs()?.max_all()?; - assert_eq!( - diff.to_scalar::()?, - 0.0, - "rank-3 {shape:?} disagreed with the 2-D result (bias={bias})" + let diff = max_abs_diff(&got.reshape((6, 7))?, &want)?; + assert!( + diff <= F16_BUDGET * scale, + "rank-3 {shape:?} disagreed with the 2-D result by {diff} \ + (budget {}, bias={bias}). That is structural, not rounding.", + F16_BUDGET * scale ); } } Ok(()) } - /// Rank 4 is the `[b1, b2, T, hidden]` shape the old `match` had its own - /// arm for; it flattens to the same three leading dims. + /// **The CPU reference oracle must be byte-stable.** + /// + /// Every V4 correctness test in this workspace compares a batched run to a + /// solo run on CPU, and those comparisons were calibrated against bit + /// identity. `forward` therefore leaves CPU on its original path, and this + /// pins it: a rank-3 activation must produce **exactly** what the + /// pre-change formulation produced — `broadcast_left` into a batched + /// matmul — element for element. #[test] - fn rank4_forward_keeps_its_leading_dims() -> Result<()> { - let l = lin(3, 4, false)?; - let a = act(8, 4)?.reshape((2, 2, 2, 4))?; + fn cpu_forward_still_takes_the_broadcast_path_bit_for_bit() -> Result<()> { + let l = lin(7, 5, false)?; + let a = act(6, 5)?.reshape((3, 2, 5))?; + + // The pre-change formulation, spelled out. + let w = l.w.broadcast_left(3usize)?; + let want = MatMul.matmul(&a, &w.t()?)?; + let got = l.forward(&a)?; - assert_eq!(got.dims(), &[2, 2, 2, 3]); - let want = l.forward(&act(8, 4)?)?; - let diff = (got.reshape((8, 3))? - &want)?.abs()?.max_all()?; - assert_eq!(diff.to_scalar::()?, 0.0); + assert_eq!(got.dims(), want.dims()); + assert_eq!( + max_abs_diff(&got, &want)?, + 0.0, + "CPU is the oracle the V4 suite is calibrated against; its \ + arithmetic must not move" + ); + Ok(()) + } + + /// The shape arithmetic the device path depends on, checked where it can + /// actually be run. Rank 3 and rank 4 both collapse to one `m`, and the + /// dims handed back are exactly what restores the original shape. + #[test] + fn flatten_batch_dims_collapses_every_leading_dim() -> Result<()> { + let a = act(8, 4)?.reshape((2, 2, 2, 4))?; + let (flat, dims) = flatten_batch_dims(&a)?; + assert_eq!(flat.dims(), &[8, 4], "every leading dim folds into m"); + assert_eq!(dims, vec![2, 2, 2]); + assert_eq!( + max_abs_diff(&flat.reshape((2, 2, 2, 4))?, &a)?, + 0.0, + "the flatten must be a pure reshape" + ); + + let b = act(6, 5)?.reshape((3, 2, 5))?; + let (flat, dims) = flatten_batch_dims(&b)?; + assert_eq!(flat.dims(), &[6, 5]); + assert_eq!(dims, vec![3, 2]); Ok(()) } - /// A non-contiguous activation reaches `reshape`, which candle services - /// with a strided copy rather than an error. Pinning it here because the - /// flatten is the only thing standing between a transposed view and - /// `matmul`. + /// A non-contiguous activation must reach `reshape`, which candle services + /// with a strided copy rather than an error. This is the only thing + /// standing between a transposed view and `matmul` on the device path. #[test] - fn non_contiguous_rank3_input_still_agrees() -> Result<()> { - let l = lin(3, 4, false)?; - // [2, 4, 3] -> transpose to [2, 3, 4]: last dim is no longer the + fn flatten_batch_dims_handles_a_non_contiguous_activation() -> Result<()> { + // [2, 4, 3] transposed to [2, 3, 4]: the last dim is no longer the // fastest-varying axis. let base = act(8, 3)?.reshape((2, 4, 3))?; let a = base.transpose(1, 2)?; assert!(!a.is_contiguous()); - let got = l.forward(&a)?; - assert_eq!(got.dims(), &[2, 3, 3]); - let want = l.forward(&a.contiguous()?.reshape((6, 4))?)?; - let diff = (got.reshape((6, 3))? - &want)?.abs()?.max_all()?; - assert_eq!(diff.to_scalar::()?, 0.0); + + let (flat, dims) = flatten_batch_dims(&a)?; + assert_eq!(flat.dims(), &[6, 4]); + assert_eq!(dims, vec![2, 3]); + assert_eq!( + max_abs_diff(&flat, &a.contiguous()?.reshape((6, 4))?)?, + 0.0, + "the strided copy must preserve element order" + ); Ok(()) } } From 5e12594382ef39106d1bccd1bf336de18f2e32a0 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick Date: Fri, 21 Aug 2026 04:21:51 +0100 Subject: [PATCH 8/8] chore(ArcKernels): drop a needless `Ok(..?)` in the new unquantized test helper MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `clippy::needless_question_mark` on `max_abs_diff`. My local run of the scoped lane predated this helper, so it went out unchecked — the lane itself was right, I ran it too early. Also noting drift for whoever maintains CLAUDE.md: the scoped lane documented there omits `-p arc-profiler`, which the real `ci.yml` Clippy job includes. Running the documented command is therefore not the same as running the gate. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01UMmjFy8TvsgypxVWNVhhC7 --- mistralrs-quant/src/unquantized/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mistralrs-quant/src/unquantized/mod.rs b/mistralrs-quant/src/unquantized/mod.rs index 0670bace4..47e82f4ce 100644 --- a/mistralrs-quant/src/unquantized/mod.rs +++ b/mistralrs-quant/src/unquantized/mod.rs @@ -796,7 +796,7 @@ mod tests { } fn max_abs_diff(a: &Tensor, b: &Tensor) -> Result { - Ok((a - b)?.abs()?.max_all()?.to_scalar::()?) + (a - b)?.abs()?.max_all()?.to_scalar::() } /// **The mathematical claim behind the flatten, and the measurement of what