minfer

A minimal local LLM inference engine built from scratch in Rust, modeled on llama.cpp's ggml_cgraph + backend scheduler, with zero ML framework dependencies.

• Declarative compute graph — inference builds a ComputeGraph (pure IR), then assigns backends, fuses ops, allocates and executes via a scheduler. Params-only graph reuse (decode steps skip reconstruction).
• Backends — CPU + Apple MPS (Metal); optional CUDA (feature-gated).
• Models — Qwen2 / Qwen2.5 / Qwen3 (dense), GGUF v3 format.
• Tooling — OpenAI-compatible server, CLI multi-turn conversation, an interactive web visualizer, and per-layer debug dumps.

Explore

• Architecture — module map, pipeline, backend layering, adding an arch.
• Compute Graph — the declarative graph design + rewrite record.
• Architecture Roadmap — system-layer gaps and the prioritized backlog.
• Backends: Metal · CPU · GPU Safety.
• Qwen3 Support and the Model Support Roadmap.
• Tooling: Debug Dump, web visualizer (viz/).

Build & run

cargo build --release
./target/release/minfer <model.gguf> "hello"

Features

This page expands the feature list from the README with the full detail. Performance numbers refer to Qwen2.5-7B-Instruct q4_K_m prefill on GB10 (sm_121) unless noted.

Inference Core

Declarative compute graph

Inference builds a ComputeGraph (pure IR) then assigns backends, fuses ops, allocates and executes via a scheduler — inspired by llama.cpp's ggml_cgraph + backend scheduler. Graph reuse is params-only (decode steps skip reconstruction), backend assignment is per-op, and the whole design is documented in COMPUTE-GRAPH-DESIGN.md.

Interactive graph visualization (viz/)

A zero-dependency browser page for the compute graph. minfer viz <model> serves the page, live SSE inference, per-node tensor stats/heatmaps and logits top-5 in one process; --dump-graph-json / MINFER_TRACE export graphs and real traces. See the viz README for the full user guide.

GGUF loader

Parses GGUF v3 files (metadata + quantized tensors) with split multi-part support; weights are mmap'd and shared zero-copy with the GPU.

Self-contained BPE tokenizer

Loaded directly from GGUF metadata — no external dependency on tiktoken. tokenizer.ggml.pre selects the pre-tokenization rule (qwen2, alias deepseek-r1-qwen; qwen35), implemented as hand-written splitters because the regex crate has no lookahead; an unknown or missing rule, an empty merge table, a non-gpt2 model or an incomplete byte vocabulary refuses the load. Special tokens (GGUF type 3/4 table plus <|im_start|>/EOS fallbacks) match as single IDs before BPE, so special-token templates (DeepSeek-R1's <｜User｜>/<think>, etc.) tokenize exactly like llama.cpp, and an unmatched piece falls back byte by byte instead of silently emitting id 0. Token ids are gated byte-for-byte against transformers / llama.cpp over tests/fixtures/tokenizer/.

Backends

CPU — AVX2 / NEON+SDOT

All 8 quantized dot products as SIMD kernels (AVX2 on x86, NEON+SDOT via inline asm on Apple Silicon), plus a persistent row-parallel thread pool (-t/--threads). Qwen3-4B CPU decode runs ~52–58 tok/s on M4 Pro (vs 1.1 before the pool).

GPU — Metal (Apple Silicon)

Flash attention (single fused kernel for decode + prefill), simdgroup GEMM prefill for every quant type, SIMD-parallel RMSNorm, float4-vectorized kernels, a build-time precompiled .metallib (no per-run shader compile), and auto-selected f16 KV cache for 7B-class models. Tracked in METAL_OPTIMIZATIONS.md.

GPU — CUDA (NVIDIA, feature-gated --features cuda)

The performance headline of the project. The int8 tensor-core MMQ path is default-on in CUDA builds (opt-out per gate with "0"; MINFER_MMQ=0 restores the legacy f16 path):

• Default prefill: ~3581 tok/s (7B q4_K_m @3314-token prompt) = 1.080× llama.cpp (llama-bench 3323.3 same shape) — from 441 tok/s when the path first landed, an 8.1× campaign documented step-by-step in CUDA_OPTIMIZATION.md (75-step history table).
• Raw-nibble int8 mma.m16n8k32 GEMMs for q4_K and q6_K with producer-fused activation quantization (rms-norm/swiglu emit the transposed q8 plane directly, skipping intermediate writes), registration-time weight-expansion planes (W_exp / W_dsc) staged by cp.async, and flash attention with register-resident softmax (2.43× kernel).
• CUDA Graph capture/replay for repeated identical-length prefills; decode uses the MMVQ weight-streaming path.
• Memory/speed knobs: the weight-expansion planes cost ~3.3 GB device for ~+6% prefill; MINFER_MMQ_Q6K_EXP=0 / MINFER_MMQ_Q4K_DSC=0 return the memory.
• Device adaptation (doc 105): the CUDA banner reports the resolved device tier — CUDA: device tier <name> (<provenance>, mmq <bool>) — from a cc-keyed table (GB10 measured; consumer GPUs adopted from llama.cpp; unknown → GENERIC). Dispatch gates (MMQ prefill availability, future batch caps) read the tier; on foreign devices the smem/VRAM feasibility checks self-degrade to slower-but-correct paths. MINFER_DEVICE_TIER=<key> forces a row for soak testing.

Model Support

Qwen2 / Qwen3 architectures

GQA attention, SwiGLU FFN, RoPE (Neox style), RMSNorm. Qwen3 adds the decoupled head dim and per-head Q/K RMSNorm (attn_q_norm/attn_k_norm, Op::QkNorm). Supported models include Qwen2.5 0.5B/7B, Qwen3 0.6B/4B, and DeepSeek-R1-Distill-Qwen-1.5B — see the matrix in AGENTS.md.

User-Facing

Model download

Auto-download from the Hugging Face Hub or the Ollama registry, with resume and cache-name resolution.

Multi-turn conversation CLI (--cnv)

Append-only KV + incremental chat-template rendering: each turn only prefills the new message delta while the whole conversation accumulates in the KV cache. In-session commands (/clear, /regen, …), automatic overflow handling, --session persistence. On overflow the dropped turn's KV rows are removed in place and the tail is re-based/re-roped (Phase C / C2), so the turn prefills its own delta instead of the retained history (measured 185 → 14 tokens per overflowing turn on the 0.5B probe); MINFER_NO_CONTEXT_SHIFT=1 restores the exact drop-and-re-render path. Plan: CLI-CONVERSATION-PLAN.md.

OpenAI-compatible HTTP server (serve)

/v1/chat/completions (streaming + non-streaming), /v1/models, /health, and /metrics — a Prometheus text snapshot of request counts, queue depth, live KV/arena occupancy and (under MINFER_OP_TIMING) per-op seconds; multi-slot with queued serial execution (MINFER_BATCH=0); a request that finds every slot busy on the batched path is refused with 503 (minfer_jobs_dropped_total +1, an SSE error frame when streaming) rather than queued — queueing is #150; a decode step whose forward fails answers every run in that batch with 500 server_error, clears their cached prefix and frees the slots (#151), so a deterministic failure cannot spin the worker on the same forward; a worker wedged inside a step ends itself after a counted STALL_STEP_LIMIT (64) consecutive no-progress steps, answering every live and queued request once with 500 server_error and publishing minfer_worker_stalled_total (#196) instead of spinning at 100% CPU; SIGINT/SIGTERM drain bounded by MINFER_DRAIN_MS. Plan: OPENAI-CHAT-API-PLAN.md, F8 record: ARCHITECTURE-EXECUTION-PLAN.md.

Chat templates (F7)

The model's own tokenizer.chat_template is rendered by minijinja plus a Python-str-method hook, so the published Qwen2.5/Qwen3 templates (including Qwen3's <think>-block split and re-emission) render as published. Reference renderings for every supported model live in tests/fixtures/chat/ (transformers 5.17.0, generated from each model's tokenizer_config.json), and a rendered prompt must equal them byte for byte. A template the engine cannot render is a loud refusal naming the construct and the template line — checked at load, so the CLI exits and the server refuses to start; the generic ChatML renderer applies only to a GGUF with no template at all. Design + accepted and refused construct sets: CHAT-TEMPLATE-AND-TOKENIZER-DESIGN.md.

Constrained decoding — grammar and JSON Schema (F2)

A GBNF-style grammar or a JSON Schema is compiled once per request into a pushdown automaton that masks the logits inside the one sampler pipeline, so decoding cannot leave the accepted language.

• GBNF subset: rules, string literals, character classes with negation, ., grouping, alternation, */+/?, repetition ranges {m}/{m,}/{m,n}, and # comments.
• JSON Schema subset: type (string or array), enum, const, object properties/required/additionalProperties, array items/prefixItems/minItems/maxItems, strings, integer bounds (inclusive and exclusive), numbers, booleans, null, anyOf/oneOf, and $defs + local $ref (recursive schemas work).
• Anything outside the subset is a loud refusal naming the construct (CLI startup error or HTTP 400) — never a silent guess. The catalogues are in GRAMMAR-DESIGN.md.
• Token advancement is byte-level correct: a token whose piece is one byte of a multi-byte character is handled, a token that a rule only partially accepts is rejected with its longest accepted prefix named, end-of-generation is legal only at a complete state, and "no token is allowed" stops with a printed reason instead of emitting an arbitrary token.
• Surfaces: CLI --grammar/--grammar-str/--json-schema/--json-schema-str; the server's response_format (json_object / json_schema) and a grammar extension field. The mask is cached per automaton state and computed with a DFA-style transition memo (measured 5.4 ms per new state on a 151,936-token vocabulary).

Performance benchmark (bench)

minfer bench <model> runs llama-bench-style prefill (pp<P>) / decode (tg<T>) throughput tests on the active backend — mean ± stddev over reps after an untimed warmup, each rep from an empty KV context without a model reload — reported as a markdown/CSV/JSON table.

GGUF tooling — convert / quantize / split (F6)

minfer convert <hf-dir> out.gguf produces a GGUF v3 from a HuggingFace Qwen2 checkpoint with the metadata the strict tokenizer/template loader requires (tokenizer.ggml.model/pre, the 256 byte tokens, merges, special ids and tokenizer.chat_template); minfer quantize in.gguf out.gguf --type … re-encodes weights to q4_0/q4_1/q5_0/q5_1/q8_0 (byte-identical to llama-quantize on the same source) or f16/f32; minfer split in.gguf dir --max-size N writes split.no/split.count parts the loader merges back into one tensor index. Unsupported architectures, tensor names, dtypes and quant targets are refused by name. A converted f16 GGUF loads and runs (f16 weights are CPU-only). See GGUF-TOOLING.md.

Philosophy

No external ML framework — pure Rust; runtime deps are minimal (rand, regex, half, serde, serde_json, minijinja; axum/tokio only for the HTTP server). Attention, RMSNorm, RoPE, SiLU, Softmax and every quantized dot product are handwritten.

Build

Requirements: Rust (edition 2021, no ML-framework dependencies — runtime deps are minimal). A Nix flake devShell (nix develop) is available for a batteries-included dev environment.

Build commands

# CPU + Metal (macOS) — plain build, never touches nvcc
cargo build --release
./target/release/minfer <model.gguf> "hello"

# CUDA (NVIDIA GPU) — opt-in feature, requires the CUDA toolkit (nvcc)
cargo build --release --features cuda
# statically-linked cudart (no libcudart.so runtime dep)
cargo build --release --features cuda,cuda_static

# + per-node debug dumps (MINFER_DUMP_DIR)
cargo build --release --features debug_dump

GGUF tooling (F6) — no extra dependencies

minfer convert / quantize / split need no Python, no PyTorch and no gguf-py: safetensors is parsed as a length-prefixed JSON header plus raw bytes and tokenizer.json/config.json are plain JSON, both through the already-present serde_json. They are host-side file work — no GPU backend is initialized — and they build in every configuration (CPU, Metal, CUDA). The writer/encoders and their verification references are documented in GGUF-TOOLING.md.

Tests

cargo test --release                         # unit + integration, no model files needed
cargo fmt --all --check                      # formatting, with the pinned toolchain's rustfmt
scripts/real_model_gates.sh                  # the #[ignore]d real-model gate set (parallel on CPU)
PARALLEL=0 scripts/real_model_gates.sh       # CPU-only serial form
FEATURES=cuda scripts/real_model_gates.sh    # device gate set (serial, one GPU)
scripts/cuda_test.sh                         # the whole CUDA suite on a real GPU, serial

The #[ignore]d set needs the cached real models (the 0.5B for the default configuration, MINFER_BATCH_TEST_MODEL=…/Qwen3-0.6B-Q8_0.gguf for the f16-KV one) and writes temporary session files. scripts/real_model_gates.sh is the documented entry point: it runs serially when FEATURES includes cuda, which a device build requires (the CUDA state is a process-wide singleton, issue #64), and in parallel otherwise — on a CPU-only build that reason does not exist, and the parallel harness is where the server batching gate checks its own robustness. PARALLEL=1/0 overrides. Each platform's CI job compiles the test target as well as the crate: cargo test --release --no-run on the Linux CPU job (implicitly), cargo test --release --features cuda --no-run on the CUDA job and cargo test --release --no-run on build-macos (#303) — cargo build does not compile #[cfg(test)], and a macOS-only test module was invisible for eleven days because of it. Those three jobs are skipped when the change classifier finds no code change (ADR-0025): a docs-only change runs the doc gates instead, and the suite still runs whenever the suite ledger moves. Since the KV storage format became per engine (#99) the parallel harness no longer makes one gate size another gate's KV regions. The current counts live in TEST-BASELINES.md and its ledger scripts/test-baselines.toml, each record dated, box-labelled and command-labelled. Four fixes now make the parallel form trustworthy: the server batching gate (interleaved matched rounds, a median verdict), #158 (a work bound instead of absolute wall-clock deadlines), #160 (every while engine.busy() stepper in the server gates routed through one shared per-step WorkBound — progress plus a step budget, both counted, never a clock), and #196 (the production serve_loop itself carries the counted no-progress bound STALL_STEP_LIMIT = 64 consecutive steps, so the two gates that drive it are wedge-proof too, and a wedged server answers every live and queued request once and publishes minfer_worker_stalled_total instead of spinning). A wedge injected with MINFER_TEST_TICK=wedge now fails each of those gates in seconds instead of hanging the suite; =spin exercises the step-budget arm (the two serve_loop gates still have to be --skipped there: a step that keeps moving the work counter is what a counter cannot catch — see the #196 record). The wall-clock bounds that remain are named backstops, not verdicts: the two run_cli child-process ceilings (1800s for the real-model sessions, 60s for the no-model registry cases; MINFER_CLI_WATCHDOG_SECS overrides both) and the serve_loop feeder's FEEDER_POLL_BACKSTOP, now the last resort for a worker whose published metrics never settle rather than the only way out of a wedge. The rules behind all of this are GATE-CONTRACT.md §3/§4 and the full records are ARCHITECTURE-EXECUTION-PLAN.md §test-infrastructure.

Git hooks (git-hooks.nix)

The formatting gate is provided by git-hooks.nix (flake input) instead of a Cargo dev-dependency. Pinned in flake.lock; the rustfmt --check hook runs with the project's pinned toolchain (1.97.1).

• The hook is installed when entering the devShell (nix develop installs .git/hooks/pre-commit + the .pre-commit-config.yaml symlink into the nix store). Committing outside the shell has no hook — those contributors should run cargo fmt --all -- --check manually first.
• Format drift is rejected, not auto-fixed: run cargo fmt --all, re-stage the changed files, then commit again.
• The same check runs sandboxed as a derivation: nix flake check (checks.pre-commit-check), or nix develop -c pre-commit run --all-files.
• CI enforces it too (#211): the test-linux-cpu job runs cargo fmt --all --check before the suite, with the toolchain step installing the rustfmt component for the pinned 1.97.1 (dtolnay/rust-toolchain@1.97.1 + components: rustfmt — 1.97.1 ships no rustfmt by default). The pinned toolchain's rustfmt is used deliberately, never stable's: a formatter that disagrees with the toolchain is exactly how a "green CI, red hook" split appears.
• In a nested worktree (.worktrees/<scope>, the documented parallel-work shape) only tracked files exist, so the ignored .pre-commit-config.yaml is absent and pre-commit run cannot work. The substitute is the CI command itself, after installing the component once: rustup component add rustfmt --toolchain 1.97.1 && cargo fmt --all --check. cargo fmt reads rust-toolchain.toml, so it uses the pinned rustfmt rather than whichever one stable happens to carry.

macOS / Metal notes

On macOS the Metal backend is built in automatically: build.rs compiles src/metal/kernels/ → a precompiled .metallib via /usr/bin/xcrun at build time (no per-run shader compile). If Xcode tools are unavailable the build still succeeds and the shader is compiled from source at first run.

CUDA build details

• nvcc is located via PATH, CUDA_HOME/CUDA_PATH (e.g. /usr/local/cuda); with --features cuda a missing toolkit is a hard error (with a clear message), and plain builds never touch nvcc at all.
• The host compiler is auto-detected: nvcc's default works when compatible; otherwise the first accepted GCC is pinned via -ccbin (e.g. a nix devShell putting GCC 15 first while CUDA 13 accepts ≤ GCC 13). Force one with MINFER_CUDA_CCBIN=/path/to/g++.
• GPU architectures are auto-detected from what the toolkit accepts (SASS for sm_70…sm_121 as available, plus PTX for the highest and a backward-JIT compute_70/compute_72 PTX) — one binary covers older and newer GPUs. The candidate list includes sm_87/sm_88 (Jetson Orin) explicitly: the device-tier gate routes sm_87 to the int8 BT path, which is compiled out of any PTX below sm_80 — native SASS per target is mandatory because PTX JITs forward only (docs 105–106, plan §14 R9). The minimum is sm_70 (Volta): the kernels in src/cuda/kernels/*.cu use WMMA tensor cores (nvcuda::wmma), which require sm_70+, so Pascal (sm_61) is not a target. The Volta V100/Titan-V (sm_70/72) PTX is only emitted when nvcc supports it: CUDA 12.x does, CUDA 13 removed Volta, so keep Volta coverage by building with CUDA 12.8 (the only version supporting Volta + the Blackwell RTX 50 sm_120/121, which needs ≥ 12.8).
• cudart linking (mirrors llama.cpp's GGML_STATIC): by default -lcudart is a shared link, so the binary NEEDEDs libcudart.so.N and needs the CUDA toolkit runtime present at runtime (an rpath to <cuda_home>/lib64 is baked in). Adding cuda_static links libcudart_static.a instead: the binary has no libcudart.so NEEDED dependency and only needs the NVIDIA driver (libcuda.so.1, dlopen'd lazily at runtime) + libstdc++ — deployable without a CUDA toolkit. The driver is never a link-time dependency in either mode.
• In CUDA builds the int8 tensor-core MMQ prefill path is default-on (runtime gates MINFER_MMQ*, see the README's Performance section / CUDA_OPTIMIZATION.md).

Usage

Examples use the built binary (cargo build --release first — see the Build section of the project README); cargo run --release -- … works identically.

./target/release/minfer <model> [prompt] [OPTIONS]

<model> can be a local path, a download URI, or a cached model name:

If prompt is omitted, reads from stdin. Run minfer --help for the full option list; the subcommands are:

Sampling and runtime options:

• --temp <T> — sampling temperature (default 0.8; --greedy = greedy decoding)
• --top-k <K> / --top-p <P> — top-K / nucleus sampling (defaults 40 / 0.95)
• --repeat-penalty <N> — repeat penalty (default 1.1; 1.0 = off), plus --frequency-penalty / --presence-penalty
• F3 sampler set (#48) — every default below leaves the pre-F3 chain unchanged:
  • --min-p <P> — drop tokens whose probability is below P * max (0 = off; 1.0 = argmax only)
  • --typical <P> — locally typical sampling (1.0 = off; 0 keeps the single most typical token)
  • --xtc-probability <P> / --xtc-threshold <T> — XTC: with probability P, exclude the top choices whose probability is at least T (T <= 0.5; 0 = off)
  • --dry-multiplier <N> — DRY (Don't Repeat Yourself) penalty strength (0 = off), with --dry-base <N> (default 1.75), --dry-allowed-length <N> (default 2), --dry-penalty-last-n <N> (default 64) and --dry-sequence-breakers <L> — restart sequences as token ids, e.g. --dry-sequence-breakers 198;13,2 (; between sequences, , between ids)
  • --mirostat <0|1|2> — mirostat off / v1 / v2, with --mirostat-tau <N> (target surprise in bits, default 5.0), --mirostat-eta <N> (learning rate, default 0.1) and --mirostat-m <N> (v1 estimator window, default 100). In mirostat mode the temperature is ignored (mirostat's mu truncation subsumes it); --temp 0 still means greedy. Mirostat cannot be combined with --spec-draft.
  • --logit-bias <L> — add to raw logits: ID:BIAS pairs separated by ,, repeatable, e.g. --logit-bias 15043:-2.0,198:1.5. A token id outside the vocabulary, or a bias outside [-100, 100], is refused at startup. A nonsensical value for any of these is refused at startup (exit 1), never silently ignored.
• F2 constrained decoding (#47) — mutually exclusive; an unsupported construct is refused at startup (exit 1) with the offending token:
  • --grammar <FILE> / --grammar-str <GBNF> — constrain sampling to a GBNF grammar
  • --json-schema <FILE> / --json-schema-str <JSON> — constrain sampling to a JSON Schema (compiled to GBNF internally) The grammar is compiled once against the loaded vocabulary; the mask is applied inside the sampler, after the penalties/DRY and before the greedy shortcut, so --greedy respects it. --spec-draft cannot be combined with a grammar (a verify round samples several rows from one automaton state). The accepted GBNF and JSON-Schema subsets, and every construct that is refused, are catalogued in GRAMMAR-DESIGN.md; the honest narrowing to know about is that object properties are accepted in declaration order and oneOf is compiled as anyOf. Example:

  minfer model.gguf "Give me a person" --json-schema-str \
    '{"type":"object","properties":{"name":{"type":"string"},"age":{"type":"integer","minimum":0,
      "maximum":150}},"required":["name","age"],"additionalProperties":false}'
  # -> {"age": 25, "name": "John Doe"}
  

• --stop <STR> — stop generation at this string (repeatable)
• -n, --n-predict <N> — max tokens to generate (default 512)
• --seed <N> — RNG seed for sampling
• Speculative decoding (D5-R, ADR-0017) — a draft model proposes, the target verifies the proposals in one batched forward:
  • --spec-draft <model> — the draft model (any supported GGUF)
  • --spec-draft-n <N> — drafted tokens per round (default 2, so the verify batch is N + 1 rows)
  • --spec-draft-adaptive — pick d per round from the per-depth acceptance and cost EWMA, capped at 8 unless --spec-draft-n sets a smaller cap --spec-draft-n and --spec-draft-adaptive are no-ops without --spec-draft. Speculative decoding refuses mirostat and a grammar (both named above), because a verify round samples several rows from one shared RNG, or from one automaton state.
• --n-ctx <N> — sizes the KV cache (clamped to the model's max context)
• -t, --threads <N> — CPU worker threads
• --gpu <N> — CUDA device index; unset auto-selects the highest compute capability. Ignored on CPU and Metal, where there is nothing to index.
• --backend <name> — F4: fence the run to a set of backends. Repeatable, and both the comma-separated (--backend cpu,cuda) and the =-joined (--backend=cpu) forms work; the names cpu, metal and cuda are matched case-insensitively. cpu stays allowed whatever is named — it is the fallback — so --backend cpu is how a run is forced onto the CPU. Unset = every backend this build has; MINFER_BACKENDS is the environment form of the same surface, and the flag wins when both are set. It is extracted before subcommand dispatch, so serve, viz, bench and specverify honour it too. Three distinct refusals, each naming what it rejected: an unknown name (unknown backend 'nope'; known backends are: cpu, metal, cuda), a name this build does not contain (backend 'cuda' is known but not compiled into this build: the CUDA backend is compiled only with --features cuda), and a name this machine cannot use (backend 'cuda' is compiled in but not available on this machine: <why>).
• --gpu-layers <N|auto> — E5: run the first N transformer blocks on the device and the rest on the CPU (0 = CPU only; unset/MINFER_GPU_LAYERS = every block the device can hold, the pre-E5 behaviour; auto = as many as the budget allows). The placement is printed at load (offload: 4 of 24 blocks on cuda, 20 on cpu; embed/output on cpu (32.0 MiB of device weights; --gpu-layers 4)), and the tensors outside the blocks — token_embd, the final norm and lm_head — stay on the CPU unless every block is offloaded.
• MINFER_GPU_MEM <MiB> — the weight budget auto fits into; unset = three quarters of the device's free bytes (the same default the activation gate uses). auto also holds back a quarter of that budget for the KV arenas and the activation pool, and reports what it decided (offload: … auto: 5 of 24 blocks fit — weights budget 64 MiB, 16 MiB reserved for KV/activations; MINFER_GPU_MEM=64 MiB). Without a device that reports free memory (Metal today) auto needs MINFER_GPU_MEM.

Three flags describe the run instead of changing how it decodes:

• --meta — dump the GGUF's metadata KV pairs and the key tensors (token_embd, the norms, lm_head) — the same two dumps minfer info <model> prints. Unlike info it does not stop there: the model still loads and generation proceeds, so it is the flag form for "show me the file, then answer".
• --dump-graph <PATH> — export the compute graph the run just built (build → assign → fusion, with the real backend assignment and the same fusion environment toggles as the live path) as Graphviz DOT. It runs one prefill, writes the file, prints Graph DOT exported to <PATH> (<n> nodes) and exits before decoding.
• --dump-graph-json <PATH> — the same graph as the JSON the viz/ web visualizer reads. It prints Graph JSON exported to <PATH> (<n> nodes, <kind>), where <kind> names the phase the graph was built for, and also exits before decoding.

Chat templates and tokenizer

Chat rendering uses the model's own tokenizer.chat_template from the GGUF. The published Qwen2.5/Qwen3 templates use Python string methods (content.split('</think>'), .lstrip('\n')) that minijinja does not provide natively; minfer supplies them through minijinja's unknown-method hook with CPython semantics, so those templates render as published (this is what makes Qwen3's think-block handling and tool-call formatting reach the model). Reference renderings — multi-turn, a system message, a generation prompt, a <think>-reasoning turn — are committed under tests/fixtures/chat/ with their provenance.

A template the engine cannot compile or render is an error, never a generic prompt:

Error: chat template error — unsupported template construct: unsupported Python str
method `splitlines` (template line 41); minfer refuses to fall back to a generic
ChatML prompt. Supported Python str methods: capitalize, count, endswith, find,
join, lower, lstrip, replace, rfind, rsplit, rstrip, split, startswith, strip,
title, upper.

The CLI exits before inference, serve/viz refuse to start (checked before the worker thread is spawned), and a per-request refusal is an HTTP 400. The generic ChatML renderer survives only for a GGUF that has no tokenizer.chat_template at all, and the startup path prints a one-line notice when that happens. --no-template still bypasses template rendering entirely.

The tokenizer is byte-level BPE and is equally strict about what it does not implement. tokenizer.ggml.pre selects the pre-tokenization rule — qwen2 (alias deepseek-r1-qwen) or qwen35 — and an unknown or missing value, a tokenizer.ggml.model that is not gpt2, an empty merge table, or a vocabulary missing any of the 256 byte tokens refuses the load with the offending value named. Token ids match the reference (transformers AutoTokenizer, or llama.cpp llama-tokenize on the same GGUF) byte for byte over the corpus committed in tests/fixtures/tokenizer/.

Multi-turn conversation

--cnv (docs/CLI-CONVERSATION-PLAN.md): append-only KV + incremental template rendering — each turn only prefills the new message delta, the whole conversation accumulates in the KV cache:

./target/release/minfer --cnv qwen2.5-0.5b-instruct-q4_0           # interactive REPL
./target/release/minfer --cnv -st qwen2.5-0.5b-instruct-q4_0 "hi"  # single turn

In-conversation commands: /exit /quit, /clear, /regen (regenerate the last reply), /help; EOF (Ctrl+D) exits.

Conversation options:

• -st, --single-turn — run one turn, then exit
• --system <STR> — system prompt
• -mli, --multiline-input — submit input on an empty line
• --color on|off|auto — color output (default auto = tty)
• --session <FILE> (with --cnv) — save/load the conversation. FILE is the history as JSON; FILE.kv is a KV session companion written next to it (C5) that carries the rows those messages were rendered into, plus the host state they belong to. On start a matching companion is resumed and the history is not re-prefilled — the run prints resumed N message(s) and M KV row(s) … — 0 tokens prefilled; anything that does not match this run (another --n-ctx, another model's n_kv_embd, another MINFER_CACHE_TYPE, a history the user edited, an older file version) prints the reason and falls back to re-rendering the JSON, which is always correct if slower. On overflow the oldest turns are dropped automatically and generation continues: the dropped turn's KV rows are removed in place and the tail is re-based (Phase C / C2), so only the new turn's delta is prefilled; MINFER_NO_CONTEXT_SHIFT=1 forces the older, exact "drop the turns and re-prefill the rest" behaviour (an engine that cannot move rows — e.g. Metal, where it is Phase G — falls back to that path on its own and says so on stderr)

Qwen3-style <think>…</think> reasoning blocks are gray-highlighted (single-shot mode too, when stdout is a terminal or MINFER_COLOR=1).

OpenAI-compatible HTTP server

Continuous batching (the worker composes one decode batch across the active slots instead of one forward per slot — Phase E / E2) is on by default when the model's forwards run on CUDA or Metal, and off on CPU (E6; Metal joined in #44 part (b), 2026-10-06). The default was decided by measurement, not assumption: on this project's reference CPU batching measured slower than serving requests one at a time (0.49x on 7B Q4_K_M, 0.88x on 0.5B Q4_0, --n-slots 4), while on the GB10 it measured 1.97x faster (7B Q4_K_M, four identical prompts, equal work, --n-slots 4, default settings).

Both readings are dated, and both are historical. The CPU pair is the E2 record's step 3 (2026-09-17); the GPU 1.97x is E6's re-fetch with the default and no environment variable (2026-09-19, matching E2's 1.9x of 2026-09-18 within noise). All three predate the cross-request prefix sharing described below (C8a/C8b, 2026-09-21/22), so the CPU figure prices a per-request prefill that sharing has since made avoidable, and the reason the E2 record gave for it — concurrency forfeits the cross-request prefix reuse each slot otherwise keeps — no longer holds. They are history, not a current claim; the boxes, the workloads and the tables are in the E2 and E6 records of ARCHITECTURE-EXECUTION-PLAN.md §5.

--n-slots does not cap a request. The arena is divided among the slots as an initial, elastic share only: the server starts from n_ctx / n_slots per slot, but a request's real bound is the whole n_ctx. When a request needs more than its share, admission reclaims cells from idle slots (releasing their cached prefixes) and the allocator moves whatever runs are in the way, so a busy neighbour cannot block it (C7/C7b, the same plan §5). The serial path (MINFER_BATCH=0) is the exception: its per-slot graph region really is n_ctx / n_slots cells.

• --slots-file <PATH> (batched engine only) — resume the server's slot contexts from PATH at startup and rewrite it after every completed request (C5 S2). A request whose prompt matches a restored slot's tokens prefills only its own delta; the file also carries the KV rows, so a restart no longer re-prefills the conversations that had finished. The startup line prints the size it will write per request (about 12 MiB for a 0.5B/512-row arena), because that cost is a decision. A snapshot from another --n-slots or --n-ctx (or another model / KV element type) is refused loudly and the server starts empty.
• MINFER_BATCH=1 forces batching on (this is how to batch on CPU, for experiments or for a machine where your own measurement says it wins).
• MINFER_BATCH=0 forces it off.
• Any other value warns and uses the device default.
• Metal joined the default only once it could take the node the batched path needs: until #44 part (b) (2026-10-06) supports_attn_span() was false there, so batching would have failed loudly instead of serving and Metal stayed opt-in. The explicit switch is no longer needed.
• The server prints its choice at startup: [server] batching: on (device cuda; MINFER_BATCH=1 forces it on, =0 forces it off).

Two fuse-related switches are easy to confuse (D3):

• MINFER_NO_FUSE_QKV=1 / MINFER_NO_FUSE_FFN=1 disable the corresponding decode fusion (the decoder builds the plain matmul/rope/store — or gate+up+ silu+mul — path instead).
• MINFER_FFN_COMPOSITION=1 keeps the FFN fusion but builds it as the proven composition (concat matmul + gate/up windows + in-place SwiGLU) instead of the hand-written fused node. It is the reference the A/B is run against, and it is ignored with a warning on a backend without offset views (Metal until G5).

The KV cache type is its own switch (C4): MINFER_CACHE_TYPE=f32|f16|q8_0, strict — an unknown value fails the load on every device, f16 resolves to f32 on the CPU (no f16 KV kernel there) and q8_0 is refused where the attention kernel has no packed read — it is supported on the CPU, CUDA and Metal (Metal since #310). A packed q8_0 cache is 3.76× smaller and, since C4 S2, is read by a fused Q8_0 × Q8_0 K dot with V accumulated out of the cell; MINFER_NO_FUSED_Q8_KV=1 restores the older dequantize-into-a-scratch read for the A/B. Details, numbers and the named tolerance class: docs/BACKENDS.md, docs/ARCHITECTURE-EXECUTION-PLAN.md §5 (C4).

Two environment switches around the GPU are easy to get wrong:

• MINFER_DISABLE_CUDA is checked for presence, not value: setting it to 0 disables CUDA (and therefore also turns the batching default off, since the model then runs on CPU). To force the CPU path deliberately use MINFER_DISABLE_CUDA=1; to use the GPU, leave it unset.
• On CUDA, batched prefills stay per request by construction (fa_prefill tiles one query tile against one KV window — E1b), so --n-slots concurrency still pays one prefill per request there. The batched-decode win is unaffected.
• Prefix reuse across slots is a share, not a copy, wherever the attention kernel gathers a kv_map (CPU, CUDA, and Metal since #362): admission points the arriving request at the donor slot's rows and no byte is copied. A device that cannot gather, and any device at all under MINFER_NO_KV_SHARE=1 (the A/B switch), falls back to C8a's row copy — either way the arriving request prefills only its own suffix.

Slot saturation

A request that arrives while every engine slot is busy is rejected loudly, not queued (the decision #121 pinned down). The worker answers

• non-streaming: 503 Service Unavailable with {"error":{"code":503,"message":"no idle slot","type":"unavailable_error"}};
• streaming: the event stream has already started, so the status line is 200 and the signal is a data: error frame carrying the same error object (followed by [DONE]) — not an empty stream a client could mistake for a generation that produced nothing.

minfer_jobs_dropped_total moves by exactly one per rejected request. The alternative — holding the request until a slot frees, which minfer_queue_depth would then measure and which the pre-E2 plan assumed (see OPENAI-CHAT-API-PLAN.md §"Slot Lifecycle") — is a feature request, #150; the serial path (MINFER_BATCH=0) still queues, because its single worker pulls one job at a time from the same channel.

Structured output (F2)

/v1/chat/completions accepts OpenAI's response_format and a llama.cpp-style grammar extension:

grammar together with a non-text response_format is a 400 (one grammar per request, never a precedence rule), and so is any unsupported GBNF/schema construct — the schema is compiled on the handler side, before the request takes a slot, so the error is an HTTP 400 with the offending construct named, not a truncated generation:

curl -s localhost:8080/v1/chat/completions -H 'Content-Type: application/json' -d '{
  "messages": [{"role":"user","content":"Give me a person"}],
  "temperature": 0, "max_tokens": 40,
  "response_format": {"type":"json_schema","json_schema":{"name":"person","schema":
    {"type":"object","properties":{"name":{"type":"string"},"age":{"type":"integer","minimum":0,
     "maximum":150}},"required":["name","age"],"additionalProperties":false}}}}'
# {"choices":[{"message":{"content":"{\n  \"age\": 25,\n  \"name\": \"Qwen\"\n}"}}]}

A grammar is per request: the compiled automaton is shared by Arc and each request (serial or batch slot) carries its own position, so slots cannot perturb each other. A response body never ends mid-character: if a generation stops with a partial UTF-8 sequence pending, those bytes are dropped (with a printed note) rather than decoded to U+FFFD.

Metrics and observability (F8)

GET /metrics returns a Prometheus text snapshot (content type text/plain; version=0.0.4; charset=utf-8). It is served from its own router state, so a scrape needs neither the tokenizer nor the job channel and keeps answering while the model is busy; rendering only reads atomics, so scraping cannot perturb generation.

Family names and units (every family is minfer_-prefixed; *_bytes is bytes, the per-op family is seconds, everything else is a count):

• Request lifecycle: minfer_requests_total, minfer_requests_completed_total, minfer_requests_rejected_total (refused before queueing: the server is draining, or the worker is gone), minfer_requests_in_flight (accepted and not yet finished — the drain surface), minfer_jobs_dropped_total (the worker could not place the job), minfer_worker_stalled_total (the worker's counted no-progress bound tripped: every live and queued request was answered 500 the worker stalled and the worker stopped — #196).
• Depth: minfer_queue_depth (accepted - admitted: the channel backlog plus the worker's pending deque — the one number neither thread can see alone), minfer_worker_pending_jobs, minfer_requests_running (occupying an engine slot now).
• Throughput: minfer_prompt_tokens_total, minfer_completion_tokens_total, and minfer_completion_tokens_per_second — generated tokens/s over a trailing 16-second window (a lifetime average would keep reporting a startup burst on an idle server). Tokens are counted where the response is produced, so the batched and serial paths agree; a client that disconnects before its answer is complete is not counted, because those tokens were never delivered.
• Drain: minfer_draining (0/1), minfer_drain_abandoned_requests (still in flight when the deadline expired; 0 is a clean drain).
• Allocator (E4 MemoryReport, for the backend the server runs on): minfer_memory_{weights,pool,live,peak_live,budget,headroom}_bytes, minfer_memory_idle_slots, minfer_memory_reserved_classes. On an unbounded backend (CPU/Metal with no budget) the budget/headroom families are omitted, not reported as 0.
• KV arena: minfer_kv_{layers,rows,region_bytes}, minfer_kv_packed (1 for a packed Q8_0 cache), minfer_kv_{reserved,owned,shared,free}_cells, minfer_kv_{free_runs,sequences}, and the C3/C8b counters minfer_kv_{defrags,cells_moved,cows,cow_cells}_total.
• Per-op timing (present only when the flag below is set): minfer_op_seconds_total{op="matmul"} and minfer_op_calls_total{op=…}. The interval is the scheduler's per-node dispatch, so it includes the backend's own prologue and excludes split-level syncs and cross-backend copies; only ops that actually ran appear.

Occupancy is a live reading: the worker republishes the allocator's numbers after every step. On the batched path that is the one shared arena; the serial path has one arena per slot, so it reports the arena of the slot that served the last request.

Flags:

• MINFER_OP_TIMING — presence-checked (any value), off by default. Turns on the per-op timing above. Off, the scheduler never reads the clock and the timing family is absent from a scrape; on, the numbers reported change and the numbers computed do not (a greedy run is identical with and without it).
• MINFER_DRAIN_MS — how long a graceful shutdown may take, in milliseconds (default 30000). A value that is not a whole number of milliseconds is reported and the default is used; 0 means "stop now". See below.

MINFER_OP_TIMING=1 ./target/release/minfer serve --n-ctx 4096 --n-slots 1 qwen2.5-0.5b-instruct-q4_0
curl -s http://127.0.0.1:8080/metrics

Graceful shutdown (F8)

On SIGINT or SIGTERM the server stops accepting new work and lets the requests already accepted finish, up to MINFER_DRAIN_MS:

1. the listener is closed (a brand-new connection gets a connection error), and a request that arrives on an already-accepted connection gets 503 — it is counted in minfer_requests_rejected_total;
2. in-flight responses are allowed to complete;
3. at the deadline the server logs how many requests were still in flight, records that count in minfer_drain_abandoned_requests, and exits. It never waits on the worker indefinitely — an SSE client that never disconnects cannot keep the process alive.

With no signal the server runs forever exactly as before, and the --slots-file snapshot (written after every completed request) is unaffected.

./target/release/minfer serve --n-ctx 4096 --n-slots 1 qwen2.5-0.5b-instruct-q4_0
# POST /v1/chat/completions  (stream + non-stream)
# GET  /v1/models, GET /health
# GET  /metrics  (Prometheus text; see "Metrics and observability" above)

Performance testing (bench)

./target/release/minfer bench -r 3 <model>        # pp512 + tg128, markdown table
./target/release/minfer bench -p 3314 -n 128 <model>  # campaign-shape pp/tg

pp<P> ingests P prompt tokens (prefill-only, generate nothing); tg<T> prefills the context then decodes T tokens. -p 0 / -n 0 skip a test, -r sets the measured reps (1 untimed warmup each), -o csv|json emits the same fields machine-readable, --n-ctx only ever grows the auto KV sizing (P+T+16, clamped to the model's context length).

Verify-step gate bench (specverify)

./target/release/minfer specverify -p 512 -r 40 -o json <model>

Measures the batched verify-step cost C_T(nt) (nt = 1, 3, 5 by default) at a fixed deep KV depth and reports the per-token amortization nt·C_T(1)/C_T(nt) — the D5 speculative-decoding gate instrument (step doc 81). -p sets the depth, -r the timed reps (3 untimed warmups each), MINFER_SPECVERIFY_NTS=1,3,16 overrides the phase list, MINFER_SPECVERIFY_NOUT=1 forces prefill-style n_out=1. Exit code is 0 whenever the measurement completes; the PASS/FAIL verdict is in the JSON.

GGUF tooling — convert, quantize, split (F6)

# HuggingFace Qwen2 checkpoint -> GGUF (f16, or f32 for a lossless archive)
./target/release/minfer convert /path/to/Qwen2.5-0.5B-Instruct out.gguf --outtype f16
# quantize an existing single-file GGUF (q4_0, q4_1, q5_0, q5_1, q8_0, f16, f32)
./target/release/minfer quantize out.gguf out-q4_0.gguf --type q4_0
# split a single file into transport-sized parts (any type; tensor bytes copied verbatim)
./target/release/minfer split out-q4_0.gguf /tmp/parts --max-size 200M

convert reads config.json, tokenizer.json, tokenizer_config.json and model.safetensors (single file or an index.json shard map) and writes the GGUF v3 metadata the engine's strict loader requires: tokenizer.ggml.model = gpt2, tokenizer.ggml.pre = qwen2, the full 256-token byte vocabulary, the merge table, the special-token ids and tokenizer.chat_template. Supported architectures: Qwen2ForCausalLM only. 1-D tensors stay f32 under --outtype f16 (llama.cpp's rule); --outtype f32 is exact for bf16/f16 sources.

quantize re-encodes 2-D float weights; 1-D norms/biases keep their source type, and on a tied model a sub-8-bit target quantizes the shared token_embd.weight at q8_0 (both are printed). K-quants and I-quants are refused by name — minfer can read them but has no encoder that has been verified against llama.cpp.

split writes {stem}-NNNNN-of-MMMMM.gguf parts with split.no/split.count/split.tensors.count; the loader reads part 0 and merges every part into one tensor index. --split-max-size/--max-size accept bytes or a K/M/G suffix, and a tensor is never split across parts.

Full contract, supported/refused sets and verification references: GGUF-TOOLING.md.

Examples

# Local model
./target/release/minfer ~/models/qwen2-0.5b-q4_0.gguf "What is the capital of France?"

# Cached model by name (no full path needed)
./target/release/minfer qwen2.5-0.5b-instruct-q4_0 "Hello"

# Auto-download from Hugging Face + run (quant auto-detected, splits included)
./target/release/minfer hf:Qwen/Qwen2.5-0.5B-Instruct-GGUF:qwen2.5-0.5b-instruct-q4_0.gguf "Hello"

# Inspect GGUF metadata + key tensors
./target/release/minfer info qwen2.5-0.5b-instruct-q4_0

# List available GGUF files in a HF repo (without downloading)
./target/release/minfer download hf Qwen/Qwen2.5-0.5B-Instruct-GGUF

# Pull from Ollama and create a symlink
./target/release/minfer download ollama qwen2.5:0.5b

# List locally cached models
./target/release/minfer list

Sampler pipeline and invariants (moved from AGENTS.md)

sampler.rs (#48): one SamplerConfig drives one pipeline — logit bias → penalties (repeat / frequency / presence, last 64 tokens) → DRY → grammar mask (F2) → greedy shortcut (temp == 0) → top-k → typical → top-p → min-p → XTC → temperature or mirostat v1/v2, seeded StdRng. Every F3 knob defaults to a no-op, so the default path is bit-identical to the pre-#48 chain (pinned by test_default_pipeline_matches_the_pinned_pre_f3_sequence and, through the grammar-aware entry point, by test_default_pipeline_matches_the_pinned_pre_f2_sequence). SamplerConfig::validate refuses nonsensical values at CLI startup / HTTP 400 (never clamps silently), and logit-bias token ids are checked against the vocabulary. Mirostat's mu is caller-owned state (MirostatState: one per run / session / request / batch slot); speculative decoding refuses mirostat (--spec-draft), because a verify round samples several rows from one shared RNG. DRY sequence breakers are token-id sequences (--dry-sequence-breakers 198;13,2); llama.cpp's string form needs a tokenizer port (follow-up). CLI: --temp --greedy --top-k --top-p --repeat-penalty --frequency-penalty --presence-penalty --min-p --typical --xtc-probability --xtc-threshold --dry-multiplier --dry-base --dry-allowed-length --dry-penalty-last-n --dry-sequence-breakers --mirostat --mirostat-tau --mirostat-eta --mirostat-m --logit-bias --grammar --grammar-str --json-schema --json-schema-str -n --seed -t.

Constrained decoding (F2, #47). src/grammar.rs compiles a GBNF grammar (or a JSON Schema, through a generated GBNF) into one pushdown automaton — a flat program per rule, a set of {rule, pc} call stacks — and turns it into a per-state token bitset. The mask is applied inside sample_with_config_grammar, after DRY and before the greedy shortcut: every stage before it only shifts logits and every stage after it only removes candidates, so a forbidden token can never be chosen, and the mask consumes no RNG (mirostat/DRY are unperturbed). The compiled Arc<Grammar> is per request; the mutable GrammarState is per run, exactly like MirostatState. Token advancement is byte-level correct (a piece may be one byte of a multi-byte character); a token whose pending bytes can never complete to an accepted codepoint is rejected, EOG is legal only at an accepting state with no pending bytes, and no allowed token is a loud stop (SampleError::NoAllowedToken), never an arbitrary token. Unsupported GBNF or schema constructs are startup/400 refusals — never a silent guess; the accepted subset and every refusal are catalogued in docs/GRAMMAR-DESIGN.md. Server: response_format (json_object / json_schema) plus a grammar extension field; the two together are a 400. --spec-draft + a grammar is refused (a verify round samples several rows from one automaton state).

Decisions governing this document

This page is the current contract for the CLI surface; the decisions behind the flags and behaviours it documents are frozen in the ADR corpus:

• ADR-0006 — The KV storage format is a per-engine gate, not a process-wide global
• ADR-0014 — A KV session is a versioned, checksummed file — never a memory dump
• ADR-0015 — The offload auto fit takes a prefix, not a knapsack
• ADR-0017 — Speculative decoding refuses the features its identity contract cannot carry
• ADR-0018 — The grammar mask is one stage inside the single sampler pipeline
• ADR-0019 — A chat template that cannot be rendered refuses the load
• ADR-0021 — bf16 is a round-to-nearest-even cast, and 1-D tensors stay f32

Support Matrix

Supported quantization formats and model architectures. This page is the expanded version of the two support sections formerly in the README.

Supported Quantization Formats

minfer supports GGUF v3 files with the following quantized weight types. The CPU backend quantizes activations on-the-fly (Q8_0 for the simple weight types, Q8_K — llama.cpp's format with precomputed per-subblock sums — for Q4_K/Q5_K/Q6_K); the GPU backends read f32 activations directly for their non-MMQ kernels, matching llama.cpp's Metal backend.

Supported

¹ Metal prefill uses a simdgroup GEMM for every quant type (dispatched when nt ≥ 2 && (od ≥ 2048 || nt ≥ 9)); the scalar f32 multi kernels handle decode (nt==1) and tiny small-od batches. The compute-graph MetalBackend dispatches these kernels per op (quant_matmul_f32_on_gpu_buf), so every quant type above runs on the GPU. ² F32 weights are supported on both GPUs: 1-D norms/biases through the norm kernels and 2-D matmul weights through CUDA's launch_f32_f32_matmul and Metal's kernel_f32_f32_matmul (#317). Before #317 an f32 weight on Metal had no arm and silently ran the Q4_0 kernel. ³ F16 has no block: 2 B per element, so there is no integer dot to run. The CPU dot is vectorized — AVX2 uses F16C (_mm256_cvtph_ps) and aarch64 uses baseline NEON FCVTL (vcvt_f32_f16) — with an f64 scalar oracle/fallback (vec_ops::dot_f16_f32, f16_dot_path()), and the multi-token prefill decodes each weight row once and threads the row loop through the shared CPU pool (#141). The AVX2 column marks the hand-written x86 kernel; NEON is folded into CPU as in every other row. ⁴ CUDA decodes in-register (f16_f32_matmul_vec / _scalar, __half22float2) and the embedding gather has its own f16 kernel — the weights stay 2 B/element on the device, which is the point of the format. No MMQ route: MMQ streams quantized bytes and f16 is not one of its formats, so an f16 prefill runs the f32-activation kernel. Both supported architectures (Qwen2/Qwen2.5 and Qwen3) use it: #141 landed the registration branch and the graph type gate in the qwen2 loader/graph only, so until #167 an f16 Qwen3 model fell to the CPU on a CUDA build even though these kernels existed; the loaders now share one registration rule (models::weight_reg). The engine's f16 file contract is 2-D tensors f16 and 1-D norms/biases f32 (llama.cpp's rule; mat_mul_f16/the f16 embed decode have no f16-norm sibling) — what minfer convert --outtype f16 writes. ⁵ Metal registers the raw 2 B/element f16 weights and promotes in-register: kernel_f16_f32_matmul (src/metal/kernels/f16.metal) is the f32-activation matmul and kernel_get_rows_f16 the embedding gather, both selected by the TensorType::F16 arms of quant_matmul_f32_on_gpu_buf / embed_tokens_gpu (#164). The weights stay half width on the device — no registration-time f32 copy — and, like CUDA, an f16 prefill runs the f32-activation kernel, not a simdgroup GEMM. Both loaders admit the type, so weights_on_gpu's all-or-nothing check passes and the model is a Metal model; the file contract above means an f16 norm can never reach a d*2 kernel buffer. Measured on macbook (macOS 27.0.1, Apple M4 Pro) (2026-10-06) against the same file's CPU logits: max |Δlogit| 2.4e-3 on the 0.5B and 7.9e-3 on Qwen3-0.6B (bar 0.05), with an identical greedy continuation ([12095, 11, 323, 432] for Qwen2, [12095, 13, 576, 6722] for Qwen3). ⁶ BF16 weights (#142): the CPU decodes one row at a time (vec_ops::mat_mul_bf16, exact f32::from_bits(bits << 16), then the same vec_dot_f32 the f16 row path uses) and the embedding rows in Op::GetRows. The decode is a left shift, so there is no separate SIMD kernel to mark in the AVX2 column (the dot itself is the vectorized vec_dot_f32). minfer convert --outtype bf16 writes 2-D bf16 / 1-D f32 and is byte-identical to llama-quantize --pure <f32>.gguf … BF16 (docs/GGUF-TOOLING.md §4.1.1). ⁷ CUDA registers the raw 2 B/element bf16 words and promotes in-register — bf16_f32_matmul_vec / _scalar (the uint4 word load split by a bits << 16 shift, the f16 pair's exact sibling) and embed_rows_bf16 — selected by the TensorType::BF16 arms of matmul_f32_ptr_layout / embed_rows_on_gpu (#208; #141 is the f16 template). The decode is exact (f32::from_bits(bits << 16)), so unlike the quantized types there is no rounding at all; the weights stay half width on the device — no registration-time f32 copy — and, like f16, a bf16 prefill runs the f32-activation kernel, not the int8 MMQ GEMM (MMQ streams quantized bytes and bf16 is not one of its formats). Both loaders admit the type through the shared models::weight_reg::cuda_weight_reg rule, so weights_on_cuda's all-or-nothing check passes for both supported architectures and the graph's BF16 matmul / embed nodes are assigned Backend::CUDA; the 1-D side is the file contract above. bf16 does not fuse: cuda::concat_rows has no 2 B/element arm, so the attn_qkv / ffn_gu concat copies are not registered and the unfused matmul chain runs. Measured on a GB10 (2026-10-06, dgxspark): a 0.5B bf16 GGUF registers the device-weight figure recorded in §"f16 and bf16 weights" (the same number as its f16 twin, i.e. the 2 B/element claim is real), 169 bf16 matmul + 1 embed nodes on CUDA, device-vs-CPU max |Δlogit| 7.82e-5 absolute / 4.24e-6 relative (bar 0.01 / 1e-3) with an identical greedy continuation [12095, 13, 1084, 374]. ⁸ Metal registers the raw 2 B/element bf16 words and promotes in-register: kernel_bf16_f32_matmul (src/metal/kernels/bf16.metal) is the f32-activation matmul and kernel_get_rows_bf16 the embedding gather, both selected by the TensorType::BF16 arms of quant_matmul_f32_on_gpu_buf / embed_tokens_gpu (#208, the Metal half; the CUDA half is footnote 7). Its own kernel, not a dtype flag on the f16 one — bf16 and f16 are different 2 B/element layouts, so a shared kernel would branch per element in the hottest device kernel. The weights stay half width on the device — no registration-time f32 copy — and, like CUDA/f16, a bf16 prefill runs the f32-activation kernel, not a simdgroup GEMM. Both loaders' Metal arm (matches!(ttype, F32 | F16 | BF16)) admits the type, so weights_on_gpu's all-or-nothing check passes and the model is a Metal model; the file contract above covers the 1-D side. bf16 does not fuse (the fused device forms are CUDA-only). Measured on a Mac (2026-10-06, macbook (macOS 27.0.1, Apple M4 Pro)) against the same file's CPU logits: 169 bf16 matmul + 1 embed nodes all on Backend::METAL, the same device-weight figure, max |Δlogit| 1.889e-3 absolute / 1.025e-4 relative (bar 0.05 / 5e-3), with an identical greedy continuation [12095, 13, 1084, 374].

⁹ The K-quant dots (#56, 2026-10-08) have AVX2+FMA kernels (src/quants/avx2.rs) and AVX-512/VNNI variants (src/quants/avx512.rs), dispatched AVX-512 → AVX2 → scalar at runtime (MINFER_NO_AVX512=1 drops to AVX2, MINFER_NO_AVX2=1 to scalar — the x86 counterparts of MINFER_NO_NEON) and gated bitwise against the scalar reference (quants::avx2_correctness). The AVX2 column marks the hand-written x86 kernel; NEON is folded into CPU as in every other row.

CUDA notes: prefill (nt ≥ 16) runs the default int8 tensor-core MMQ path for the common quants (Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q4_K via the f16-wmma GEMM, Q4_K/Q6_K via the raw-nibble int8 kernels — the promoted ~3581 tok/s path, see CUDA_OPTIMIZATION.md); the f32-activation kernels cover every type including Q5_1/Q5_K, and decode (nt==1) uses the dp4a MMVQ kernels (Q4_K/Q5_K/Q6_K with shape gates). Q5_K requires id % 32 == 0 (tail-masking granularity).

GPU grouping note: the old whole-layer layer_gpu path required all 7 weight matrices in a layer to share one quant group (all-Q4 or all-QK) and fell back to CPU otherwise. The compute-graph path (default) has no such restriction — backend assignment is per op, so mixed-group layers run fully on the GPU. Raw weights are not supported on GPU and select the CPU backend for those ops.

KV Cache Storage Type by Backend

MINFER_CACHE_TYPE picks the KV cache element type, which is a separate axis from the weight type above (graph/kvformat.rs is the single authority, and the answer for "can this backend read it" is the registry's reads_packed_kv). A value the backend has no kernel for is refused at load, never silently mapped to f32.

Notes:

• f16 on the CPU resolves to f32 — the CPU has no f16 KV kernel, and an env var set for a GPU run must not break a CPU one.
• The default on CUDA/Metal is the model's own auto policy (f16 when n_layers × n_kv_embd ≥ 8192, i.e. the 7B class, f32 for small models); the table's "default" row is the region shape, which f16 does not change.
• Q8_0 is the packed one: 3.76× smaller than f32 and 1.88× smaller than the f16 auto policy. On CUDA a Q8_0 decode runs the layout-tagged split-K kernel together with the packed fused QKV epilogue (attn_bias_rope_store_q8_0, #144), and a prefill at head dim 128 runs the packed FA prefill — its f16-tile staging dequantizes each packed block, so the tensor-core route is offered for a packed cell too (#144: Qwen3-0.6B pp2048 564.5 → 8231.1 tok/s). Still off their tuned route, and stated: the verify band (1 < nt ≤ 16) takes the general layout-tagged kernel and the hybrid 4-warp decode dispatch is f16-typed. The dp4a packed K dot is not a follow-up — it landed in #186 (cc19b4f, 2026-09-27) as the packed decode route, and docs/ARCHITECTURE-EXECUTION-PLAN.md §C4 #186 records it DONE with its tolerance class re-measured. The general layout-tagged kernel remains the fallback for every packed path. A speculative session refuses a packed cache outright (its greedy identity contract rests on the batched split kernel). See docs/ARCHITECTURE-EXECUTION-PLAN.md §5 C4 #144 and docs/cuda_optimization_steps/107-c4-packed-q8-kv-cuda.md.
• Metal (#310) — enabled. kernel_store_kv_q8_0 writes the same bytes as the CPU quantizer. Two read mechanisms cover every attention shape, selected by the pure crate::metal::packed_attn_route: mechanism A reads packed cells natively in the decode flash family (kernel_flash_attn_ext_q8_0 / _hd128_q8_0, nt == 1, hd ∈ {64,128}); mechanism B dequantizes the needed window into a transient f32 stage (kernel_dequant_kv_q8_0_to_f32) and runs the unchanged f32 prefill / windowed-flash family. The classic kernel_gqa_attn_q8_0 / _window_q8_0 / _map_q8_0 remain the fallback for a small/odd hd, an nt == 1 explicit window and any MINFER_NO_* opt-out. READS_PACKED_KV is now true, so MINFER_CACHE_TYPE=q8_0 loads and runs on Metal. The memory win is 3.76×; the measured speed (macbook (macOS 27.0.1, Apple M4 Pro), 2026-10-08, minfer bench -p 1024 -n 64 -r 3 --n-ctx 2048, 3 interleaved runs, medians; f16 baseline): Qwen3-0.6B pp1024 4844 → 4683 tok/s (0.967×) and tg64 193.5 → 176.6 (0.913×), Qwen2.5-0.5B pp1024 6175 → 6154 (0.997×) and tg64 295.7 → 245.5 (0.830×); region sizes 58 720 256 → 15 597 568 B and 6 291 456 → 1 671 168 B (both 3.76×). The stage is f32, not f16: f16 staging's second rounding was measured to amplify to 16.9 logit delta on Qwen3-0.6B over 8 decode steps, outside the inherited C4 class; f32 staging restores it (1.28 / 0.36) at parity speed. docs/METAL-BACKEND-DESIGN.md §4.4 records the mechanisms, the measurements and the gates.
• A Q8_0 cell width must be a whole number of 32-element blocks (so n_kv_embd % 32 == 0, which every supported architecture satisfies); ensure_kv refuses anything else.
• The KV write/move side (Backend::copy_cells for C3 compaction / C8a prefix copy / C8b S3 copy-on-write, and GraphAllocator::copy_kv_to_cpu for the C2 shift and C5 sessions) is implemented on all three backends since #44 part (b): Metal moves rows one at a time with MTLBlitCommandEncoder in the overlap-safe order and reads its regions back through the registry host_read hook. A physical shift of an f16 region refuses loudly and is pinned by a gate (#306: the host round trip has no dequantize → re-rope → requantize map, so the CLI re-renders the retained window instead); the per-engine kv_format is what makes a Metal session describe the width its region really uses.

Not Yet Supported

Q5_K and Q5_1 are fully supported on CPU and both GPU backends — Q5_K_M models run at full GPU speed.

f16 and bf16 weights — the file contract and the measured device deltas

Moved here from AGENTS.md (its Support section keeps the one-line summary).

f16 weights (F6/#49; device half by #141 on CUDA, #164 on Metal): an f16 GGUF runs on CPU, CUDA and Metal for both supported architectures (Qwen2/Qwen2.5 and Qwen3). Both device loaders register the raw 2 B/element weights and the kernels promote in-register; the CUDA side shares one registration rule with the CPU (models::weight_reg, which also carries the q4_K W_dsc plane gate of #165), while the Metal side is the loader's matches!(ttype, F32 | F16 | BF16) registration arm (the BF16 arm is #208's Metal half). Op::MatMul decodes one f16 weight row at a time (vec_ops::mat_mul_f16) and Op::GetRows decodes f16 embedding rows; 1-D norms/biases stay f32 — the file contract every producer writes (minfer convert --outtype f16, llama-quantize … F16, minfer quantize --type f16), and the CUDA norm path checks the registered byte length so an f16 norm cannot reach a d*2 buffer (#169). CUDA registers the raw 2 B/element weights and converts in-register (f16_f32_matmul_vec / _scalar, embed_rows_f16); an f16 prefill does not enter the int8 MMQ GEMM — it runs the f32-activation kernel. Metal is the same shape: kernel_f16_f32_matmul + kernel_get_rows_f16 (src/metal/kernels/f16.metal) are its f32-activation matmul/embed, selected by the TensorType::F16 arms of quant_matmul_f32_on_gpu_buf / embed_tokens_gpu, and its prefill runs that same f32-activation kernel (no simdgroup GEMM). minfer quantize supports q4_0/q4_1/q5_0/q5_1/q8_0 and the K-quants q4_K/q5_K/q6_K (byte-identical to llama-quantize; the K-quant reference is llama-quantize --pure, and --type q4_K writes one uniform type, not the Q4_K_M mixture — #203), f16 and f32, and refuses every type without an encoder by name (#140); a K-quant row that is not a multiple of 256 is demoted the way llama.cpp's tensor_type_fallback does it; --type f16 keeps 1-D tensors f32, --type f32 writes every tensor f32. The gate inputs are recorded: tests/fixtures/f6-fixtures.json carries one entry per cached artifact content (path, bytes, sha256, the exact producer command, the producer identity — minfer commit, or llama.cpp commit + compiler + effective -ffp-contract — date and an absolute box label), scripts/check_f6_fixtures.py audits that manifest in CI, verifies the whole cache where it exists, and — since #345 — re-runs a recorded producer with --regenerate and re-records the content identity it produces (idempotently; a llama-quantize reference, or a minfer entry whose recorded commit is not the one that runs, is refused by name rather than re-recorded), and every F6 gate verifies the fixture it resolves (src/tooling/tests/f6_fixtures.rs) so a stale or replaced reference refuses the run by name and digest instead of being compared against — #205.

bf16 weights (#142): minfer convert --outtype bf16 writes 2-D bf16 / 1-D f32, f32→bf16 being round-to-nearest-even (ggml_compute_fp32_to_bf16, NaN forced quiet; general.file_type = 32 = LLAMA_FTYPE_MOSTLY_BF16). The CPU decodes one bf16 weight row at a time (vec_ops::mat_mul_bf16, exact f32::from_bits(bits << 16), then the same vec_dot_f32 the f16 row path uses) and Op::GetRows decodes bf16 embedding rows; TensorType::BF16 is the new type the loader maps GgmlType::BF16 to. CUDA registers bf16 since #208's CUDA half — models::weight_reg::cuda_weight_reg answers Raw for it (the f16 arm's sibling: 2 B/element, no f32 copy, clear_nb_bt_only), Op::MatMul takes the TensorType::BF16 arm of matmul_f32_ptr_layout (bf16_f32_matmul_vec when id % 8 == 0, else _scalar; both shift bits << 16 in-register, which is exact) and Op::GetRows the BF16 arm of embed_rows_on_gpu (embed_rows_bf16), and both loaders' CUDA branch admits the type through that one shared rule, so the all-or-nothing gate answers Cuda for both supported architectures. bf16 is not an MMQ format, so its prefill runs the f32-activation kernel; cuda::concat_rows has no 2 B/element arm, so the attn_qkv/ffn_gu concat copies are not registered and bf16 runs unfused. Measured on GB10 2026-10-06 (dgxspark, scripts/cuda_test.sh device gate f208_bf16_weights_run_on_the_cuda_device): 0.5B bf16 registers 942.4 MiB of device weights (the f16 twin's number), 169 bf16 matmul + 1 embed nodes all assigned Backend::CUDA, device-vs-CPU max |Δlogit| 7.82e-5 absolute / 4.24e-6 relative (bar 0.01 / 1e-3), greedy [12095, 13, 1084, 374] identical. Metal registers it too since #208's Metal half — kernel_bf16_f32_matmul + kernel_get_rows_bf16 (src/metal/kernels/bf16.metal, the pl_bf16_f32 / pl_get_rows_bf16 pipelines) are its f32-activation matmul/embed, selected by the TensorType::BF16 arms of quant_matmul_f32_on_gpu_buf / embed_tokens_gpu, and both loaders' Metal arm is now matches!(ttype, F32 | F16 | BF16), so weights_on_gpu passes and both supported architectures answer Device::Metal. Measured on a Mac (macbook (macOS 27.0.1, Apple M4 Pro), 2026-10-06, device gate f208_bf16_weights_run_on_the_metal_device): the 0.5B bf16 file's 169 bf16 matmul + 1 embed nodes are all assigned Backend::METAL, 942.4 MiB of device weights, device-vs-CPU max |Δlogit| 1.889e-3 absolute / 1.025e-4 relative (bar 0.05 / 5e-3), greedy [12095, 13, 1084, 374] identical. The writer is byte-identical to llama-quantize --pure <f32>.gguf … BF16, 290/290 tensor payloads (169 BF16 2-D, 121 F32 1-D; docs §4.1.1 — the reference is cast from the f32 conversion, because bf16→f16 is lossy below 2^-14 and the f16 file cannot carry those values). On the bf16-source 0.5B, bf16-vs-f16 CPU logits differ by max 2.29e-5 / 1.24e-6 relative (not bitwise — the f16 file's 123 024 subnormal weight values are the whole difference) with an identical greedy continuation [12095, 13, 1084, 374].

Operator Coverage by Backend

supports_op decides at graph-build time which backend runs each node (docs/ARCHITECTURE.md §5). This table is the contract, generated from the three implementations — keep it in step with them.

The table's concrete F32 Op rows are pinned by graph::op_matrix::support_table_matches_support_matrix_doc, which fails when a backend's supports_op disagrees (each backend column is checked wherever it is compiled in). Two kinds of row it cannot check, so they are prose plus their own gates: the composite rows (one line spelling several ops) and the capability nuances that are not an Op field — a partial View at offset 0 (the allocator backstops it; supports_op sees only the offset), the set-valued kv_map attention window (Device::gathers_attn_map), and a packed q8_0 KV region (BackendCaps::reads_packed_kv).

Notes on the asymmetries — these are the rows where a model's decode path differs by platform:

• FusedQkvNorm is Metal-only. Qwen3 decode on CUDA takes the unfused QkNorm path, which is numerically equivalent but issues more dispatches. Making CUDA fused is a Phase G / CUDA-verifiable ticket, not a correctness gap.
• QkvBiasRopeStore is CUDA-only — a recorded decision (#52), not a gap. It is the mixed-quant decode epilogue: q/k/v use different quant types (so they cannot share FusedQKV's single concat weight), so three separate matmuls (no bias) feed one bias×3 + RoPE×2 + store×2 pass. CUDA fuses it (10 dispatches → 4 per layer, −6); Metal keeps the unfused chain and the graph builder never emits the node there (metal_backend.rs's false arm is a design statement). Porting would save 6 dispatches per mixed-quant layer — 84 per decode token on Qwen2.5-7B-Q4_K_M, the realistic case, whose 14 of 28 layers carry attn_v as Q6_K against Q4_K q/k — with no numerical difference (supports_op is a build-time gate and the unfused chain is the reference). The whole forward's host-encode is ~0.2 ms against a ~20 ms/token decode, so those 84 dispatches are a sub-1% slice of decode time; the A/B of the concat-class fusion that removes more dispatches (MINFER_NO_FUSE_QKV=1, −8 on the same 14 layers) sits within run-to-run variance on macbook (macOS 27.0.1, Apple M4 Pro) (2026-10-06, five interleaved bench -p 0 -n 128 -r 4 pairs: 48.06 vs 46.49 t/s means, individual pairs crossing zero), so a second kernel path and its bitwise gate are not earned by a ~1% ceiling. Provenance and the rejected alternative, stated so the numbers are not misread. The ~0.2 ms host-encode figure is MINFER_OP_PROFILE=1 on that 7B — per-op host-encode totals, not per-label counts — and the 10 → 4 / 84-per-token dispatch counts are the CUDA D3-8 ledger applied to Metal's dispatch table; Metal has no per-op profiler, so they were not re-counted on the device. And the port is cheap: Metal already has the class-1 attn_bias_rope_store kernel, so the refused work is mainly a three-pointer binding — the decision rests on the measured ceiling, not on the size of the change.
• Interleaved RoPE is CPU-only. Both loaders hard-code NonInterleaved today, so no shipped model hits this; a family that needs interleaved RoPE needs a loader change plus a CUDA kernel.
• Scale/Softmax are CPU-only and unused. Attention kernels fuse the softmax and carry the scale in AttnMeta, so no supported architecture emits either node.
• BatchMatMul is deferred everywhere (single-output IR) and nothing emits it.

Supported Model Architectures

minfer currently supports two model architectures.

Qwen3 support: dense architecture only (no MoE / hybrid-SWA / VL variants yet). The dense models reuse the Qwen2 graph with two deltas — the head dim is read from qwen3.attention.key_length (decoupled from n_embd / n_head) and Q/K go through a per-head RMSNorm (blk.{i}.attn_q_norm / attn_k_norm) before RoPE. See QWEN3-SUPPORT-PLAN.md for the design + verification record.

How Architecture Detection Works

minfer reads the general.architecture string from the GGUF metadata header. Only the exact values "qwen2" and "qwen3" (case-sensitive) are accepted. Any other value produces a clear error:

Unsupported architecture: 'llama'

The loader will not silently misinterpret a non-Qwen2 model — it fails immediately with a descriptive message. All model-agnostic components (BPE tokenizer, Jinja2 chat template renderer, samplers) are ready for additional architectures once the graph construction (build_graph) is added.

Hyperparameter Keys

The Qwen2 loader reads GGUF keys from both qwen2.* and llama.* prefixes. The llama.* fallback exists for compatibility with older GGUF converters that used the llama. prefix as a de-facto standard for Llama-family hyperparameters. This does not mean Llama architecture is supported.

Adding a New Architecture

See AGENTS.md for a step-by-step guide. In brief:

1. Create src/models/<name>/ with mod.rs, graph.rs, loader.rs
2. Add a match branch in src/models/mod.rs::load_model()
3. Define HParams, LayerWeights, and implement the ModelDef trait (including build_graph(&self, params) -> ComputeGraph, which is deterministic in params — the graph-reuse invariant)

Architectures that share Qwen2's tensor naming convention (LLaMA, Mistral, Phi) should be relatively straightforward to port.

Decisions governing this document

This page is the current contract; the decisions behind it are frozen in the ADR corpus:

• ADR-0013 — The CPU quantizes activations to Q8_0; a device reads f32
• ADR-0021 — bf16 is a round-to-nearest-even cast, and 1-D tensors stay f32
• ADR-0005 — Metal becomes a first-class backend
• ADR-0006 — The KV storage format is a per-engine gate, not a process-wide global
• ADR-0020 — A quantized file is byte-identical to llama-quantize, or it is wrong
• ADR-0032 — The packed-KV staging window is f32, not f16

minfer Architecture

A pure-Rust LLM inference engine written from scratch, inspired by llama.cpp, with zero ML framework dependencies. Inference runs through a declarative compute graph (builder → scheduler → per-backend kernels), modeled on llama.cpp's ggml_cgraph + backend scheduler. This document describes the overall design: module responsibilities, the compute graph pipeline, the CPU / Metal backend layering, quantization, adding a new model architecture, and adding a new backend. The pre-graph imperative forward is preserved at the end as an appendix (Appendix A).

1. Design Principles

1. No ML framework — attention, RMSNorm, RoPE, SiLU, softmax are all handwritten. Only 5 external crates: rand, regex, half, serde/serde_json, minijinja.
2. Declarative compute graph, not an imperative loop — the forward pass is built as a pure ComputeGraph (no side effects at build time), then assigned to backends, fused, allocated, and executed by the scheduler. This mirrors llama.cpp (ggml_cgraph + ggml_backend_sched) and enables graph reuse, per-op backend assignment, DOT export, and a clean path to new backends. Design and implementation record: docs/COMPUTE-GRAPH-DESIGN.md.
3. Bytes-in / bytes-out tensors — weight tensors are raw &[u8]; SIMD dot-product kernels (AVX2 / Metal shaders) operate on byte slices matching the exact GGML quantized block layout (repr(C) in block.rs).
4. Activations stay f32 — CPU matmuls quantize activations to Q8_0 on-the-fly; the Metal backend reads f32 activations directly for all weight types (matching llama.cpp's Metal backend).
5. Backend assignment is a build-time decision, never a silent mid-execution fallback — supports_op decides which ops run where; cross-backend transfers happen at split boundaries; kernel-invariant violations abort (gpu_abort / Err), they never silently fall back to CPU.

2. Module Map

The backend layers — runtime, launch, kernels, executors

Each device backend is organised along four layers, and only the last one is polymorphic:

(The CUDA paths in that table have landed (#262, #263); the Metal ones are still the target of Step 4 (#265, Mac-local). docs/SOURCE-LAYOUT-PLAN.md §8 has the file-by-file tree and each step updates this table as it lands.)

Two rules follow, and they are why the crate keeps one flat interface instead of a directory per layer:

1. Device is the first axis, the layer is the second. Backend + registry.rs (src/graph/backend.rs) are the one device seam — llama.cpp keeps ggml-backend*.cpp/h flat beside its per-device directories for the same reason. L2 cannot leave L1: the CUDA launchers are inherent methods of CudaState, the Metal ones methods of MpsCommandBuffer; splitting them out would be a type refactor, not a file move.
2. A common needs a second real implementation. A shared abstraction is added only when at least two backends implement it and at least two callers use it with the same semantics. The one candidate today is allocplan::DeviceMemory (CUDA answers it; the Metal half is #53). CPU's quants.rs/vec_ops.rs are deliberately not a device-private layer: they are the crate's numeric kernel library, consumed by graph/kvformat.rs and graph/cuda_backend.rs.

src/graph/ — the compute graph core

3. Inference Pipeline

The top-level flow lives in main.rs. The whole engine is a single-pass prefill followed by an autoregressive decode loop; both call ModelDef::forward, which routes through the compute graph. Every forward is build → assign → fuse → alloc → execute (the subgraph below), but the built graph is cached per GraphParams and reused — only the first call with new params pays the build/assign/alloc cost.

flowchart TD
    A["CLI args: model, prompt, flags"] --> B{"resolve model<br/>download::resolve"}
    B -->|local path| C["load GGUF v3<br/>single or split parts (mmap)"]
    B -->|"hf:… / ollama:…"| D["auto-download → path"]
    B -->|cached name| C
    C --> E["parse metadata KV + tensor table"]
    E --> F["init GPU backend<br/>MPS / CUDA"]
    F --> G["load model<br/>dispatch on general.architecture"]
    G --> MODE{"mode?"}
    MODE -->|"--cnv"| CV["conversation REPL<br/>(run_conversation)"]
    MODE -->|"serve"| SV["OpenAI-compatible HTTP server<br/>multi-slot, streaming"]
    MODE -->|"viz"| VZ["viz server (page + live SSE)"]
    MODE -->|single shot| H["load BPE tokenizer from GGUF"]
    H --> J{"no-template?"}
    J -->|no| K["render chat template<br/>GGUF chat_template via minijinja<br/>(unrenderable → loud refusal;<br/>no template at all → ChatML)"]
    J -->|yes| L["raw prompt"]
    K --> M["tokenize prompt<br/>n_ctx = max(--n-ctx, prompt len)<br/>(sizes the KV regions once)"]
    L --> M
    M --> N["PREFILL<br/>graph forward, all prompt tokens at once"]
    N --> O["last-token logits (n_out = 1)"]
    O --> P{"DECODE loop<br/>while generated < n_predict"}
    P --> Q["sample next token<br/>penalties → top-k → top-p → temp"]
    Q --> R{"stop token or<br/>stop string match?"}
    R -->|yes| S["done"]
    R -->|no| T["append token, decode+print"]
    T --> U["graph forward, single token<br/>KV persists in the allocator"]
    U --> P

    subgraph GRAPH["every forward: build → assign → fuse → alloc → execute"]
        G1["GraphBuilder<br/>build_graph (pure IR)"] --> G2["assign backends<br/>priority Metal → CUDA → CPU"]
        G2 --> G3["fuse<br/>SwiGLU + decode fusions (gated)"]
        G3 --> G4["alloc<br/>liveness + persistent KV"]
        G4 --> G5["execute<br/>per split, cross-backend copies"]
    end
    N -.->|"GraphCache: params-only reuse"| GRAPH
    U -.->|"GraphCache: params-only reuse"| GRAPH

Generation parameters (defaults match llama.cpp): temp=0.8, top_k=40, top_p=0.95, repeat_penalty=1.1 (last 64 tokens), frequency_penalty=0.0, presence_penalty=0.0, seed=42, n_ctx=4096, n_predict=512. Sampling applies the three penalties in one pass, then top-k → top-p → temperature.

Timing is dual-caliber: Prefill: = prompt tokens / prefill wall time; Generated: = generated tokens / decode wall time (pure decode, matches llama-bench "Generation" caliber); Total: = blended.

4. Compute Graph Architecture

Inference = build a ComputeGraph (pure, side-effect free) → assign backends → fuse → allocate → execute. The graph is built once per distinct GraphParams and reused (decode steps reuse the same graph; the model's forward() routes through Qwen2Graph::forward).

4.1 The IR

• ComputeGraph = topologically ordered nodes (builder appends sources before consumers), inputs, outputs, uid (for CUDA-Graph-style reuse).
• CNode = op + src dependencies + out_shape/out_dtype + backend (None until assigned) + meta (weight names, rope/attn params).
• Op carries full payloads (RmsNorm{eps}, MatMul{transpose_b}, KvcacheStore{layer} …) so graphs are structurally comparable.

4.2 Builder (GraphBuilder)

Per-architecture code calls builder methods (mirroring llama.cpp's llm_graph_context): embedding, rms_norm, qk_norm (Qwen3 per-head Q/K norm), matmul, get_rows, rope, silu, add, mul, swiglu, softmax, attn, kvcache_store/kvcache_load, plus the decode-fusion constructors fused_qkv/qkv_bias_rope_store/fused_ffn. Building is pure — no computation happens at build time.

4.3 Scheduler pipeline

assign_backends → fuse → alloc_graph → execute

1. assign_backends — capability-driven: each node gets the highest-priority backend whose supports_op returns true (priority Metal → CUDA → CPU). Weight registration decides GPU feasibility.
2. fuse — pattern matching (Mul(Silu(X),Y) → SwiGLU) gated per backend by supports_fused (no double-fusion with hand-written kernels). The plan's second rule (RoPE(Add(X,B)) → FusedBiasRope) was removed: no backend ever claimed the capability, and the bias+rope work ships as the builder's decode nodes (FusedQKV/QkvBiasRopeStore, whose attn_bias_rope_store kernel subsumes the bias+rope part). BatchMatMul stays deferred (single-output IR limitation, COMPUTE-GRAPH-DESIGN.md §5.4).
3. alloc_graph — per-backend liveness allocator: buffers shared between nodes whose live ranges don't overlap; persistent per-layer KV regions survive rebuilds; in-place ops alias their input buffer (see §4.5).
4. execute — per split (contiguous same-backend runs): sync the previous backend, copy split inputs across backends (copy_across, a host round trip via shared memory), run the nodes, then a final sync. Metal batches one MpsCommandBuffer per split, submitted at synchronize().

4.4 Reuse (GraphCache)

Params-only deterministic reuse (llama.cpp allow_reuse invariant): GraphParams = n_tokens / n_out (tail rows) / gtype / cparams (n_ctx, flash_attn, gpu, fuse_qkv, fuse_ffn, explicit_span) / weights_version deterministically determines the topology — equal params ⇒ identical graph. n_past is deliberately absent (it is execution data). CParams.gpu records backend participation so a backend toggle forces a rebuild. GraphCache owns the allocator (so the KV regions persist across rebuilds, e.g. the prefill→decode transition) and try_reuse compares params only; debug builds assert structural consistency (Op: PartialEq).

4.5 Core invariants (must not be violated)

1. KV positions are data, not structure. KvcacheStore/Load carry only the layer index; write positions come from the positions input node. Topology never depends on n_past.
2. Each layer owns TWO persistent KV regions (K and V), resolved by kv_pair(layer) (KvProvider). The store node's output buffer is the K region; backends write the V sibling via kv_pair.
3. GGUF weight layout: metadata [in, out] (ne[0] fastest), memory [out][in] row-major → matmul od = shape[1], id = shape[0]. Activations: shape metadata [d, nt, 1, 1], memory token-major [nt][d]. I32 inputs are stored as f32::from_bits bit patterns (fill_input_i32).
4. In-place ops (Silu, RoPE) alias their input buffer — the allocator maps the output to the input's BufRef (only when the input's sole consumer is this op AND it is on the same backend). Never host-copy a GPU-pending buffer: a host copy_in of a producer that is encoded but not submitted reads stale data (the Phase-3 KV-corruption bug). Cross-backend in-place inputs get a fresh buffer (the producer completed before the split boundary, so the copy is safe there).
5. Execution follows build order (a valid topological order by construction) — guarantees a KV store executes before the attention that reads it. Nodes with no allocated buffer (dead, e.g. fusion orphans) are skipped.
6. CPU vs GPU activation paths differ numerically: CPU matmuls are Q8_0×Q8_0 (activation-quantized), the GPU backends read f32 activations for all weight types (CUDA prefill additionally offers the default-on int8 MMQ path). Compare each path against its own reference, not against the other.

4.6 Per-layer computation (Qwen2) — as built by graph.rs

flowchart LR
    A["token_embd lookup (GetRows)"] --> B["hidden"]
    B --> C["RMSNorm attn_norm"]
    C --> D["WQ / WK / WV matmuls + bias"]
    D --> E["RoPE on Q and K (in-place alias)"]
    E --> F["kvcache_store: K/V → layer regions"]
    F --> G["GQA attention<br/>Q·K^T → softmax → ·V"]
    G --> H["WO matmul + bias"]
    H --> I["+ residual → hidden"]
    I --> J["RMSNorm ffn_norm"]
    J --> K["FFN gate + up matmuls"]
    K --> L["SiLU(gate) × up (fused SwiGLU)"]
    L --> M["FFN down matmul"]
    M --> N["+ residual → hidden"]
    N --> O["next layer / output_norm"]

GQA: each query head h maps to KV head hk = h / gqa. The KV head dimension is independent (n_kv_embd read from the K weight's ne[1]), so Qwen2.5-0.5B (n_embd=896, n_head=14, hd=64, n_kv_embd=128) strides correctly.

4.7 Prefill vs decode

The graph path keeps K/V on the executing backend (attention and KV are on the same backend by construction), so there is no per-token KV drain.

5. Backend Layering

5.1 The Backend trait (graph/backend.rs)

Backend (the thing the IR stores on a node) is not this trait: since F4 it is a Copy handle — an id into src/graph/registry.rs — and the trait below is implemented by the pools (CPU, Metal, CUDA) and held by the registry entry. The name, id, priority and capability record live on the handle/entry, not on the trait (Backend::name() is registry.rs:105):

#![allow(unused)]
fn main() {
pub trait Backend: Send + Sync {
    fn kv_pair(&self, layer: usize) -> Option<(usize, usize)>;   // persistent per-layer K/V regions
    fn supports_op(&self, op: &Op, dtype: DType) -> bool;
    fn supports_fused(&self, fused: &FusedOp) -> bool;
    fn supports_attn_span(&self) -> bool { false }       // E1: explicit attention windows (CPU/CUDA/Metal)
    fn weights_bytes(&self) -> usize { 0 }               // E4: bytes the memory budget is charged for
    fn alloc_buffer(&mut self, size: usize) -> usize;    // backend's own pool
    fn free_buffer(&mut self, id: usize);
    fn pool_len(&self) -> usize;
    fn alloc_fresh(&mut self, size: usize) -> usize;     // bypasses the recycle free list (split-boundary staging)
    fn execute_node(&mut self, node: &CNode, in_bufs: &[usize],
                    out_buf: usize, kv_pair: Option<(usize, usize)>) -> Result<(), String>;
    fn copy_cells(&mut self, dst: BufRef, src: BufRef, dst_row: usize, src_row: usize,
                  rows: usize, elems_per_cell: usize) -> Result<(), String>;   // C3 compaction / C2 shift
    fn read_host(&self, id: usize) -> Option<&[f32]>;
    fn write_host(&mut self, id: usize, data: &[f32]) -> Result<(), String>;
    fn write_host_window(&mut self, id: usize, offset: usize, data: &[f32]) -> Result<(), String>;
    fn synchronize(&mut self);
    fn retire(&mut self) { self.synchronize(); }         // release device objects for a dropped pool
    // CUDA only: try to replay a captured graph for (uid, range); capture is
    // gated to decode-shaped graphs. Default impl returns false.
    #[cfg(feature = "cuda")]
    fn graph_replay(&mut self, uid: u64, range: (usize, usize), nt_hint: Option<usize>) -> bool;
}
}

• CPU (cpu_backend.rs): Vec<f32> pool; executes via kernel.rs + vec_ops.rs; F32 weights use a plain f32 matmul, quantized weights use cpu_quant_matmul_f32 (Q8_0-activation path).
• Metal (metal_backend.rs, macOS): shared-memory MTLBuffer pool; per-op dispatch to MpsState's kernels (rms_norm, quant_matmul_f32_on_gpu_buf, rope, silu, add, mul, swiglu, embed_tokens_gpu, store_kv, gqa_attn_f32); one command buffer per split. Weights resolve by name from MpsState's registry (weight_buf(name) -> (buffer, offset)).
• CUDA (cuda_backend.rs, feature-gated --features cuda): wraps the cuda.rs device layer — per-op dispatch with int8 MMQ prefill + MMVQ decode (default-on; MINFER_MMQ=0 reverts; the MMQ gate resolves from the device_tier.rs device-tier table), split-KV attention, and CUDA Graph capture/replay keyed on the graph uid (graph_replay, decode-shaped only). Implementation record: docs/CUDA-BACKEND-DESIGN.md; per-step optimization history in docs/CUDA_OPTIMIZATION.md.

The allocator owns every backend pool (single source of truth); the scheduler orchestrates assignment, cross-backend copies, and sync.

5.2 Selection rules

Assignment asks each node in the registry's priority order — Metal 300, CUDA 200, CPU 100 (src/graph/registry.rs:72-75) — for the first backend that has a pool, is not fenced off, and reports the op supported; the CPU is tried last and is always allowed, so the answer is total (GraphAllocator::supports_for, src/graph/alloc.rs:403). The fence is --backend <name> / MINFER_BACKENDS=<csv>, and an unknown name, a name this build does not contain and a name this machine cannot use are three distinct startup refusals. The two orders, the name surface and those refusals are docs/BACKEND-REGISTRY-DESIGN.md's contract; the user-facing flags are docs/USAGE.md.

• Metal: all graph weights must be GPU-registered (Qwen2Graph::weights_on_gpu mirrors the old per-layer check). MINFER_DISABLE_MPS=1 forces CPU.
• CUDA: feature-gated (--features cuda), and every graph weight must be GPU-registered with a device present. A node whose block the E5 offload plan left on the CPU never reaches the device (supports_for answers CPU for it, so a partially offloaded model cannot run a block whose weights were never registered there). MINFER_DISABLE_CUDA=1 is presence-checked and forces the CPU.
• CPU: always available; AVX2 dispatch via is_x86_feature_detected!("avx2") on x86, NEON+SDOT on aarch64, scalar fallback elsewhere (MINFER_NO_NEON=1 forces scalar; the K-quant dots also have an AVX-512/VNNI path since #56, gated by MINFER_NO_AVX512=1 and MINFER_NO_AVX2=1).

5.3 GPU safety

All Metal submits wait bounded (10 s) and check status; no early return past a threadgroup_barrier; device limits queried at runtime, never hardcoded; guard failures gpu_abort with actual values. In the graph architecture, kernel-invariant violations return Err from execute_node and must NOT be treated as a silent CPU fallback — backend assignment at build time decides where ops run; only genuine support limitations (e.g. Raw weights) select the CPU backend. See docs/GPU_SAFETY.md.

6. Quantization & Tensor Layout

• Weight layout (block.rs): repr(C) blocks matching ggml-common.h. Q4_0/Q4_1/Q5_0/Q5_1/Q8_0 = 32-value blocks; Q4_K/Q5_K/Q6_K = 256-value super-blocks. Supported: Q4_0, Q4_1, Q8_0, Q4_K, Q6_K, Q5_0, Q5_1, Q5_K (CPU + Metal), F32/F16 norms & biases.
• Activation quantization: CPU matmuls quantize f32 activations to Q8_0 on-the-fly (quants.rs); Metal reads f32 directly.
• GGUF parsing (gguf.rs): ggml_pad(x, n) = (x + n - 1) & !(n - 1) alignment; tensor strides computed from type block size exactly like ggml; split multi-part models are merged into one tensor index.

flowchart LR
    A["GGUF file"] --> B["metadata KV<br/>hparams + tokenizer + template"]
    A --> C["tensor table<br/>name / type / shape / offset"]
    C --> D["quantized data blob"]
    D --> E["Tensor: type, shape, strides, Vec&lt;u8&gt;"]
    B --> F["HParams"]
    B --> G["Tokenizer"]
    B --> H["Chat template"]

7. KV Cache

The graph path owns the KV cache: two persistent regions per layer (K and V), sized n_kv_embd × n_ctx, allocated in the backend pool the layer runs on. kv_pair(layer) resolves them; the KV store node writes K/V at the positions carried by the positions input; attention reads the written prefix. The regions live inside the GraphCache's allocator and survive graph rebuilds (the prefill→decode transition). MINFER_CACHE_TYPE=f16 selects an f16 GPU cache where the kernels support it. The pre-graph cache.rs KVCache type — and the vestigial &mut KVCache argument of ModelDef::forward — was deleted in #252; these regions are the only KV store.

8. Adding a New Architecture

1. Create src/models/<name>/ with mod.rs, graph.rs, loader.rs.
2. Add a match branch in src/models/mod.rs::load_model() for the new general.architecture value.
3. In loader.rs: define HParams (including n_kv_embd) and LayerWeights, parse them from GGUF metadata (both qwen2.* and llama.* prefixes are accepted).
4. In graph.rs: implement build_graph(&self, params: &GraphParams) -> ComputeGraph deterministically in params (the reuse invariant), using GraphBuilder — mirror llama.cpp's llm_graph_context builder methods.
5. In mod.rs: implement ModelDef (forward, build_graph, forward_graph, forward_graph_cached, as_any, special_tokens, n_layer/n_head_kv/n_embd_head/n_kv_embd/n_vocab, rope_style). models/qwen3/ is the worked example of a second architecture (decoupled head dim + per-head Q/K norm).
6. If needed, add a chat template format in template.rs.

Architectures that share Qwen2's tensor naming convention (LLaMA, Mistral, Phi) are the easiest ports. The RopeStyle enum defines both Qwen2 (non-interleaved) and Llama (interleaved) pairings, but only the non-interleaved form is wired up: both loaders hard-code it (models/qwen2/loader.rs:155, models/qwen3/loader.rs:173) and the CUDA backend refuses the interleaved style (graph/cuda_backend.rs:1522). A family that needs interleaved RoPE requires a loader change plus a CUDA kernel — see docs/MODEL-SUPPORT-ROADMAP.md, "Cost model: what a port actually costs".

9. Model Download

download/mod.rs resolves hf:<repo>[:<quant>] and ollama:<model>[:<tag>] URIs, downloads via curl with size-checked resume, quant-matches single or split files case-insensitively, and stores them under ~/.cache/minfer/models. Cached filenames can be used directly as the model argument.

10. Related Documentation

Appendix A — Legacy Imperative Architecture (removed in Phase 6)

Historical reference. The imperative per-layer forward (models/qwen2/forward.rs) was the engine's core until the compute graph replaced it (Phase 6); it is preserved here verbatim in structure so old notes, benchmarks, and kernel analyses remain interpretable. Do not treat this as the current design.

A.1 Design stance (then)

The engine ran a direct per-layer forward loop instead of a compute graph: "simpler, easier to trace, and the whole layer can be fused onto the GPU." The GPU fallback was per-layer and safe: a layer that could not run on the GPU (e.g. unsupported weight type) submitted partial GPU work, downloaded the hidden state, and continued on the CPU.

A.2 The old forward pass

ModelDef::forward(tokens, positions, kv) was implemented in models/qwen2/forward.rs as an imperative loop:

flowchart LR
    A["token_embd lookup"] --> B["hidden"]
    B --> C["RMSNorm attn_norm"]
    C --> D["WQ / WK / WV matmuls + bias"]
    D --> E["RoPE on Q and K"]
    E --> F["store K/V into KV cache"]
    F --> G["GQA attention<br/>Q·K^T → softmax → ·V"]
    G --> H["WO matmul + bias"]
    H --> I["+ residual → hidden"]
    I --> J["RMSNorm ffn_norm"]
    J --> K["FFN gate + up matmuls"]
    K --> L["SiLU(gate) × up"]
    L --> M["FFN down matmul"]
    M --> N["+ residual → hidden"]
    N --> O["next layer / output_norm"]

Decode-time optimizations: fused QKV (nt==1) via a concatenated blk.{il}.attn_qkv weight, fused bias+rope+store (attn_bias_rope_store), fused SwiGLU kernel, and the last layer computed only the tail n_out rows (an inp_out_ids-style partial-row optimization). The n_out tail-row optimization was subsequently carried into the graph path as the G3 work: the builder inserts GetRows(wo, tail_ids) (and the same for the residual) before the last layer's FFN, so the last FFN, the final norm and lm_head all run on n_out rows only (models/qwen2/graph.rs:64, :228-231; GraphParams.n_out is part of the reuse identity).

A.3 Old backend layering & fallback

flowchart TD
    A["forward nt tokens"] --> B{"embedding on GPU?"}
    B -->|yes| C["GPU embed lookup → buf_hidden"]
    B -->|no| D["CPU embed → upload hidden"]
    D --> E["upload positions"]
    C --> E
    E --> F{"per-layer: layer_gpu ok?"}
    F -->|"yes, all layers"| G["output_norm_gpu<br/>on GPU"]
    F -->|"no at layer i"| H["submit partial GPU work<br/>download hidden, sync KV to CPU"]
    H --> I["CPU loop from layer i"]
    G -->|"output on GPU"| J["download logits → return"]
    G -->|"output fell back"| I
    I --> K["output_norm + output matmul on CPU"]
    K --> L["return logits"]

Selection rules (then):

• Metal: layer 0 must have all 7 weight matrices + norms registered on the GPU. Within a layer all 7 matrices must be all Q4 group (Q4_0/Q4_1) or all QK group (Q4_K/Q5_0/Q6_K); Q5_1/Q5_K use the f32 path and are exempt. MINFER_DISABLE_MPS=1 forces CPU.
• CUDA (--features cuda): requires every layer's 7 matrices to be all Q4_0/Q4_1 or all Q4_K/Q6_K. Decode replays a captured CUDA Graph.
• CPU: always available; AVX2 dispatch via is_x86_feature_detected!("avx2"), scalar fallback elsewhere (plus the AVX-512/VNNI K-quant dots, #56).

The GPU path skipped the per-token CPU→GPU KV drain (no sync_kv_to_cpu) because GPU-layer failure is deterministic by weight type — the sync only happened in the fallback branch.

A.4 Old KV cache

cache.rs provided an architecture-agnostic per-layer cache: k/v were pre-allocated Vec<f32> of max_size × dim, size tracked the current sequence length. store_multi wrote K/V for many positions at once (prefill); decode wrote one. The GPU maintained its own buffers and sync_kv_to_cpu copied them back only on the CPU-fallback path. MINFER_CACHE_TYPE=f16 selected an f16 GPU cache (opt-in).

A.5 Old "Adding a New Architecture"

1. Create src/models/<name>/ with mod.rs, forward.rs, loader.rs.
2. Add a match branch in src/models/mod.rs::load_model().
3. Define HParams (including n_kv_embd) and LayerWeights in loader.rs.
4. Implement the per-layer forward pass in forward.rs using kernel::, vec_ops::, cache::.
5. Implement ModelDef in mod.rs.
6. Add a chat template format in template.rs if needed.

Decisions governing this document

This page is the current contract; the decisions behind it are frozen in the ADR corpus:

• ADR-0001 — Inference runs through one declarative compute graph
• ADR-0012 — Device is the first axis, the layer the second — and no premature common
• ADR-0002 — Topology is a function of GraphParams alone, so positions cannot be structure
• ADR-0007 — No ML frameworks: every operator is hand-written
• ADR-0008 — GPU safety: bounded waits, no early return past a barrier, runtime device limits
• ADR-0009 — A failure is an error, never a silent fallback
• ADR-0013 — The CPU quantizes activations to Q8_0; a device reads f32

minfer Architecture Roadmap

Status: analysis + prioritized backlog (no code changed). Baseline: minfer HEAD = 293eb19 (2026-09-15, working tree clean). Method: direct reading of src/ (44,105 LOC) and docs/ (~15.6k LOC). Every claim about minfer carries a file:line anchor; claims taken from a document rather than code are marked "per docs".

Scope. Architecture-level work on the system layers: IR, scheduler, allocator, KV/sequence state, batching, backend abstraction, and model/quant coverage. Kernel micro-optimization is out of scope — CUDA_OPTIMIZATION.md and METAL_OPTIMIZATIONS.md cover it, and both campaigns are formally converged. Model-family selection is a separate document: docs/MODEL-SUPPORT-ROADMAP.md.

0. Verdict

minfer's kernel work is complete and competitive for the architectures it supports: eight quantized weight types on CPU, Metal and CUDA; int8 MMQ (quantized matrix-multiply on tensor cores) prefill and dp4a MMVQ (quantized matrix-vector multiply) decode on CUDA; simdgroup GEMM (general matrix-multiply) and flash attention on Metal; speculative decoding with adaptive draft depth. On the measured CUDA target, 7B Q4_K_M reaches ~3581 tok/s prefill and ~51.2 tok/s decode; on Metal, 7B decode is at parity and prefill reaches ~0.82–0.85× of the achievable ceiling, with the residual measured as not source-addressable.

The remaining work is at the system layer, and three items dominate it:

1. No multi-sequence batching. The IR, the attention kernels, the KV (Key/Value) allocator and the server are all single-sequence. E1/E1b made the attention side sequence-aware and E2 landed sequence-addressable batches and a batching server worker; the CPU payoff measured negative while the GPU payoff was 1.9x, so the default follows the device (E6: batches on CUDA and, since #44 part (b) landed the one-range attn_span read and the matching write/move side on a Mac (2026-10-06), on Metal too; off on CPU).
2. KV cache is a fixed per-layer buffer, not a sequence-addressable cell store: no sequence ids, no eviction/context shift, no defragmentation, no state save/restore, no quantized KV. (Phase C has since closed every one of those: C1/C2 the cell store and physical removal/shift, C3 the compaction — a pure planner, the counters, Backend::copy_cells on CPU and CUDA, the K re-rope a move needs while positions are cells, and a model-level continuation gate — C4 the packed Q8_0 cache (CPU, CUDA and Metal — the last in #310) and C5 the session container. What is still open is driving the compaction from the server's serving model, which reserves every slot once and packed, the CUDA packed-decode residual's stall attribution (#212 — #144/#186/#202 landed the fused epilogue, the dp4a dot and the FA prefill), and Metal's packed half (G5, #44; #310, enabled on Metal — mechanism A native packed decode + mechanism B f32 stage). See §2.4.)
3. The server has no persistent context. Every request builds a fresh GraphCache (KV regions + device pool re-allocated, CUDA Graph capture re-warmed) and re-prefills the whole prompt. Fixed in B2/B3 for the prefix-matched case: a slot keeps its cache and reuses the rows its prompt already covers (measured ≈11× TTFT on a second turn).

Everything else — IR expressiveness (views, multi-output), memory placement policy (VRAM budget, layer offload), backend pluggability, additional model families, quantizer tooling — is real but secondary, and mostly enabled by fixing (1) and (2) first.

Ranked backlog (full list with numbering in §3)

1. Current state

1.1 Scale

1.2 Landed surface

Compute graph (build → assign → fuse → alloc → execute) with params-only reuse (graph/cache.rs:47, graph/params.rs:51); per-op backend assignment (graph/scheduler.rs:60); liveness allocator with persistent per-layer KV regions (graph/alloc.rs:165, :385); eight quantized weight types (Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q4_K/Q5_K/Q6_K) on CPU + Metal + CUDA; two model families (qwen2, qwen3 dense); GGUF v3 with split parts and Metal zero-copy mmap; self-contained BPE tokenizer; chat templates via minijinja; greedy/penalty/top-k/top-p/temperature sampling; speculative decoding with adaptive depth; multi-turn conversation CLI; OpenAI-compatible server (chat completions, streaming, multi-slot); a runtime graph introspection stack (MINFER_TRACE per-node capture, live SSE visualizer, pipeline view, DOT/JSON export).

1.3 Documentation hygiene

The index-level documents (README.md, ARCHITECTURE.md, CPU_OPTIMIZATIONS.md, QWEN3-SUPPORT-PLAN.md, OPENAI-CHAT-API-PLAN.md, MODEL-SUPPORT-ROADMAP.md) are the entry points AGENTS.md routes readers through, so a stale statement there costs more than one buried in a campaign record.

A pass on 2026-09-15 fixed these stale statements:

• README.md and ARCHITECTURE.md — the removed BiasRope fusion in the two architecture diagrams, plus the fusion.rs row in the module table (the pass produces only SwiGLU; the decode fusions are built by the model code, not by FusionPass).
• ARCHITECTURE.md — the "n_out tail-row optimization is not yet ported" claim (it landed as the G3 work), and the RopeStyle "already covers Qwen2 and Llama" claim (only the non-interleaved form is wired up).
• CPU_OPTIMIZATIONS.md — the pre-compute-graph snapshot is now banner-marked as historical, with the still-open AVX2 gap pointed at §2.7.
• QWEN3-SUPPORT-PLAN.md — the self-contradictory status line.
• OPENAI-CHAT-API-PLAN.md — the "KV is f32" note (f16 is auto-selected for 7B-class models).
• MODEL-SUPPORT-ROADMAP.md — the misleading status line.

Keep the invariant: when a change lands, update the index rows in the same commit. No outstanding doc-truth work item remains.

2. Layer-by-layer analysis

The layer names follow the vocabulary of docs/GLOSSARY.md, which classifies the campaign's terms into seven layers; §2.8 adds one cross-cutting layer for ops, safety, testing and observability.

Legend for Gap: 🔴 structural (blocks a class of use cases) · 🟠 material (measurable capability or performance loss) · 🟡 hygiene.

2.1 L1 — Graph IR and build

Today. ComputeGraph is a flat, topologically ordered vector of CNode (graph/mod.rs:83-115). Each node has exactly one output, a static [usize; 4] shape, and an Op carrying its full payload. Op::View { offset, shape }, Reshape, Permute exist (ops.rs:105-116) but both backends execute them as identity copies (cpu_backend.rs:384-388, cuda_backend.rs:433-438, metal_backend.rs likewise) — the allocator has no aliasing for views, so a view costs a full tensor copy. The builder appends sources before consumers, so node-id order is the execution order (scheduler.rs:214).

Gap. 🔴 Strided, zero-copy views and multi-output nodes were missing — closed by D1 (2026-09-19): views are zero-copy with allocator-known aliasing, and GraphBuilder::split_parts exposes a node's single output as independent parts, so what remains under this header is model work (item 17), not IR work. Without them the IR cannot express slicing, concatenation, or a copy between overlapping regions as compositions of primitive ops — which is why minfer carries four hand-written decode-specific fused ops instead: FusedQKV, QkvBiasRopeStore, FusedFFN, FusedQkvNorm (ops.rs:125-154). Each one costs changes in five places: the builder constructor (builder.rs:193-281), the allocator's special cases (alloc.rs:232-273), the scheduler's kv_pair resolution (scheduler.rs:261-271), every backend's supports_op + execute_node, and both models' build_graph. BatchMatMul is deferred for the same reason (COMPUTE-GRAPH-DESIGN.md §5.4).

Consequence for the model work. MODEL-SUPPORT-ROADMAP.md ranks new model families by how much of the existing graph they reuse. That ranking is accurate for parameter-isomorphic architectures (Llama/Mistral) but not for anything needing new structure: MoE (Mixture-of-Experts) needs mul_mat_id (3-D expert indexing), MLA (Multi-head Latent Attention) needs a different KV layout and a view-based latent split, and sliding-window attention needs a mask parameter threaded into attention. Each of those currently pays the five-site cost.

Recommendation. Add strided views with allocator-known aliasing, and let a node carry a small output list. Fusions can then be produced by the fusion pass from compositions rather than hand-written as IR variants, and the four decode fusions become compiler output instead of bespoke ops.

2.2 L2 — Scheduler and execution

Today. assign_backends (scheduler.rs:68-83) walks nodes in build order and assigns the highest-priority backend whose supports_op returns true. split_graph partitions into contiguous same-backend runs derived from that positional assignment (scheduler.rs:73-120); split inputs/outputs are the crossing edges. Execution then walks splits, calling sync_backend(prev) and copy_across for each cross edge (scheduler.rs:176-189), and finally sync_backend.

Gap. 🟠 Two distinct issues.

1. No cost model. Splits are positional. A single unsupported op in the middle of an otherwise GPU graph produces three splits with two host round trips; the assignment never considers the cost of the resulting movement, or whether a neighbouring op could be moved so the op becomes supportable. In practice this is masked today because the supported models build a single GPU split on both GPU backends (CUDA-BACKEND-DESIGN.md:313-315), but it is exactly what breaks the first time an op is unsupported — e.g. interleaved RoPE on CUDA (cuda_backend.rs:1522), or FusedQkvNorm, which CUDA does not advertise (cuda_backend.rs:1905) while Metal does (metal_backend.rs:686).
2. Synchronous, host-mediated cross-backend movement. — partly closed by F5 (#58, 2026-09-24): the boundary is now two registered phases. Phase A (copy_across → BackendEntry::copy_cross) enqueues the transfer — for a CUDA source a device→host cudaMemcpyAsync into a pinned slab plus a cudaEventRecord, where the pre-F5 code did copy_to_cpu (a full stream sync + a blocking cudaMemcpy) — and phase B (await_cross) waits on the recorded event at the documented synchronization point, once per staged input. The CPU's hooks are a synchronous host round trip and a no-op (it has no device memory); Metal runs the same two phases — a MTLBlitCommandEncoder copy into a shared staging buffer plus encodeSignalEvent on a MTLSharedEvent, waited once by a bounded waitUntilSignaledValue — ported by #137 and verified on a Mac (2026-10-05: blocking_host_copies 21 → 0, copies == waits, bitwise identical). Counters (graph/copystats.rs), the per-backend table, the enumeration of the synchronization points and the measured before/after are in docs/BACKEND-REGISTRY-DESIGN.md §11 and the plan's F5 record. What remains: nothing filed for the boundary copy. The wait is already deferred to the consumer's first use (#138), so a boundary with several staged inputs holds them in flight at once; true cross-split overlap has no target on the reachable macOS topology, and #300 recorded that negative result — a single device split feeding a host consumer whose first node reads the dominant staged tensor (docs/BACKEND-REGISTRY-DESIGN.md §11.6). Multi-device execution is still the place that would need it.

Also: the cross-boundary staging map is keyed by node id alone (alloc.rs:34), so a node consumed by two different foreign backends can only have one staging buffer; the scheduler compensates by filtering on backend (scheduler.rs:252-255), which is correct only while at most two backends are enabled. Fixed in A5: the map is keyed by (node, destination backend), one node can feed two foreign consumers, and the consumer-side filter is gone.

Recommendation. Add an assignment pass that (a) propagates support backwards from unsupported ops and (b) scores a candidate assignment by crossing count; and replace the host round trip with an async copy plus a recorded event when the source and destination are different devices (F5 landed the CUDA half of this; the Metal half and the actual consumer-side waiting that would let a copy overlap independent work are the remaining follow-ups). Neither is urgent today; both are prerequisites for multi-device execution and for any heterogeneous split.

2.3 L3 — Allocator and memory

Today. GraphAllocator::alloc_graph (alloc.rs:513-521) rebuilds the whole node→buffer mapping on every graph rebuild: it frees every previously live buffer back to the backend pool (:165-177), recomputes last_use over build order, and re-allocates. Buffer pools are per backend, not unified (alloc.rs:315-335), and allocation is by exact element count — both the Metal (metal_backend.rs:322-332) and CUDA (cuda_backend.rs:1333-1353) pools scan a free list for an exact byte-length match and otherwise allocate fresh. free_buffer never returns memory to the device (cuda_backend.rs:1355-1362).

Gap. 🟠 Three consequences.

1. Per-rebuild teardown. Non-persistent buffers are freed and re-derived on every topology change. For a decode loop with a stable shape this happens once; for anything with a varying n_tokens (speculative decoding with adaptive depth, server requests of differing prompt lengths, future continuous batching) it happens every time, together with a full device re-allocation for any size not seen before.
2. No size classes / no rounding. Because matching is exact, a workload touching k distinct activation shapes ends up with k sets of live buffers. The pool is a per-GraphCache high-water mark that never shrinks (documented as accepted debt, CUDA-BACKEND-DESIGN.md:447) — but that was reasoned about for a fixed-shape CLI run, not for a server whose prompt lengths vary per request. E4 S2 (2026-09-23) rounds every pooled activation to its class, so shapes inside one class share a buffer across rebuilds; E4 S3 then split reservation from assignment (a slots table per backend and class: a released buffer is re-assigned rather than returned to the pool, so a rebuild re-maps and CUDA's pool_gen — the capture generation — stops moving) and gave GraphCache one graph per GraphParams, so a switch is a re-map, not a build (§14 row 3, closed by E4 S3).
3. No VRAM budget or feasibility check. cudaMemGetInfo is queried once at init and only printed (src/cuda/methods/init.rs:126); the sole consumer of free-memory information today is a valve guarding the optional f16 weight cache (src/cuda/methods.rs:57-62) — the activation/KV allocator has no accounting at all. Out-of-memory surfaces as a null pointer that fails at execute time (cuda_backend.rs:1345-1348). There is no "this graph will not fit, offload the last n layers" fallback because there is no layer-offload concept at all (§2.6).

Recommendation. Split allocation into a reserve phase (size the graph for the worst-case shape the session will use: n_ctx, n_batch) and an assign phase that re-maps nodes into reserved regions without touching the device. Round pool sizes to a size-class ladder (e.g. 16/64/256 KiB steps) so distinct shapes share. Add real memory accounting (bytes_reserved, bytes_live, bytes_device_free) and make the prefill path consult it before allocating, with CPU execution as an explicit fallback decided at build time (consistent with the "no silent fallback" rule — the fallback would be visible in the backend assignment).

2.4 L4 — KV cache and sequence state 🔴

Today. Each layer owns two persistent contiguous regions, K and V, created on first use by ensure_kv(layer, backend, size) (alloc.rs:385-393), sized n_kv_embd × n_ctx f32 (or f16 when the type flag is set). They live in the allocator inside GraphCache and survive graph rebuilds (graph/cache.rs:69). Positions are data, injected per step through the positions input node, so the topology never depends on n_past (ops.rs:94-99, COMPUTE-GRAPH-DESIGN.md §1.4).

That design is a genuine strength — it is the precondition for graph reuse — but it stops at the single-sequence append-only case. What is missing:

Gap. 🔴 This is the single largest structural gap, because it blocks four separate user-visible capabilities at once: multi-slot serving throughput, cross-request prompt caching, long-conversation context handling without re-prefill, and any state-space/hybrid model family.

Two concrete defects live here as well:

• ensure_kv ignores the requested size after the first call (alloc.rs:385-393: the early if let Some(&pair) = self.kv.get(&layer) returns without comparing size). CParams.n_ctx is part of the reuse identity, so a session that changes n_ctx on the same GraphCache silently keeps the old region. The CPU backend then errors on out-of-range positions (cpu_backend.rs:170-172); the GPU backends do not check at all (below). Today no caller changes n_ctx on a live cache, so it is latent rather than live — but it is a landmine directly under the "grow the context" feature that a server wants.
• GPU KvcacheStore had no bounds validation — closed. A3 moved the check to the backend-agnostic GraphAllocator::fill_input_i32 (check_positions_bound), the single point where positions/cells become graph data, so CPU, CUDA and Metal all refuse an out-of-range row before execution; the Metal gap-table G1 ticket #38 then added the arm-level guard for fill paths that bypass fill_input_i32 (MetalBackend::check_kv_store_rows, covering KvcacheStore, FusedQKV and FusedQkvNorm). A caller violating the documented contract now gets an Err naming the cell and the arena instead of an out-of-bounds device write.

What C1/C2 landed (2026-09-16). src/graph/kvcache.rs now owns a per-layer arena with an owner per cell, and GraphAllocator::kv_rm(start, len, &rope) removes a row range and re-bases the rows after it — a physical operation, so cell == pos survives and no backend needed a new kernel. The conversation's overflow path uses it instead of dropping turns and re-prefilling them. Its one approximation is inherent and recorded in ARCHITECTURE-EXECUTION-PLAN.md §5 (C2 record): the retained rows hold the values they were written with, so rows that attended to the dropped turns keep that influence — no shift that avoids re-prefilling can avoid it (llama.cpp's context shift behaves the same way). What is exact is the mechanism, and that is what the tests pin bitwise: a tail removal leaves the retained head byte-identical (so continuing from it matches a fresh prefill exactly), a middle removal copies V verbatim and re-ropes only K, and the whole operation keeps the identity cell mapping C1's scheduler gate requires.

Recommendation. Redesign the KV layer as a sequence-addressable cell store before adding batched attention, because batching without per-sequence KV addressing cannot be correct. Concretely: a KvCache owning a per-layer arena of n_ctx cells, each carrying the set of sequence ids that own it; the store node resolves (layer, seq_id) → cell index on the host and passes an index array to the kernel; the attention kernel receives an explicit per-query allowed-cell mask instead of deriving the bound from positions. E1 landed that last part on CPU: attn_span carries each query's [lo, hi) cell range (resolved from ownership + the query's position), and a representation as a range is complete because a sequence's cells are contiguous — a per-cell mask is what a hole-creating layout would need (C3/D1). Build one such cache with an optional window parameter and an optional recurrent state, rather than a family of per-variant caches. That single abstraction simultaneously delivers prefix reuse (cells already owned by the matching prefix), eviction (drop cell ownership rather than re-prefill), and defragmentation (a cell-copy op that a strided-view IR makes expressible).

2.5 L5 — Batching and serving 🔴

Today. Single sequence by default on CPU — batching exists and its default follows the device (E6: on for CUDA, off for CPU/Metal; MINFER_BATCH=0/1 forces it), slower on CPU (0.49x) and 1.9x faster on the GPU, where the E2 acceptance is met. GraphParams carries no sequence count at all: E2 deleted n_seqs, which A7 had kept as reserved for item 3, once item 3 landed and showed the count is data rather than topology (CParams.explicit_span carries the only topology decision it can force — see ARCHITECTURE-EXECUTION-PLAN.md §8). The other dead identity field, CParams.n_batch, was deleted in Phase A7; chunked prefill (item 10) will reintroduce it with its real semantics. Attention derives its causal bound from the per-token positions input (cpu_backend.rs:411-416, cuda_backend.rs:1100-1102), so two sequences in one batch would attend to each other.

The prefill is one graph covering the entire prompt (main.rs:862-890): n_ctx = max(--n-ctx, prompt_len), one forward with nt = prompt_len. The server rejects any prompt longer than the slot context (chat.rs:118-124); new_slots (slot.rs:33-36) divides n_ctx_total equally per slot. worker_loop drains the queue serially, one slot at a time (chat.rs:452-518), and every request starts from a fresh GraphCache (chat.rs:77). Default --n-slots is 1 (main.rs:211), so the default server is strictly serial with a cold KV per request.

Gap. 🔴 This is the largest capability difference between what minfer does today and what a serving workload needs, and it costs throughput on exactly the workload the server exists for. With --n-slots N minfer gets N independent serial sessions, not N-way batching: aggregate decode throughput stays at batch-1 tokens/s × (fraction of time the GPU is busy).

Recommendation. Work in dependency order: (1) KV cells → (2) IR seq_id and mask inputs → (3) attention kernels with explicit masks → (4) batch composition in the scheduler/worker → (5) n_batch chunking. Do not start with (4); it cannot be made correct first.

Two smaller but immediate items sit in this layer:

• Persistent server context. chat.rs:491 discards the GraphCache per request. The comment justifies it as avoiding cross-request KV contamination (a real bug fixed in doc 97) — but the correct fix is sequence-aware KV invalidation (§2.4), not discarding the whole cache. As written, every request pays KV + pool re-allocation (≈235 MB for 7B at n_ctx=4096, f16) and invalidates the CUDA Graph capture warm-up (pool_gen bump, cuda_backend.rs:1342). Fixed in B2: the slot keeps its cache and a record of the tokens its rows hold; a request reuses the KV only when its prompt starts with exactly that sequence. Measured end-to-end (0.5B Q4_0, 219-token second turn): prefill 219 → 16 tokens, time-to-first-token ≈2.8 s → 0.25 s (≈11×), with the cold-slot turn unchanged.
• No panic isolation. Fixed in A4. The guard that existed (guarded_forward, chat.rs:585-597) only covered forward_graph_cached; the speculative path calls both models' forwards directly (spec.rs) and the tokenizer, sampler and stop-string paths were bare. A panic in any of them unwound worker_loop, dropping the bounded job channel (capacity 64, server/mod.rs:66) — later requests got 503 — and the queued jobs' StreamEvent senders, ending their SSE (Server-Sent Events) stream after an empty [DONE] rather than an error (server/mod.rs:198-216). The whole per-job body now runs under run_job_isolated, which turns a panic into a logged 500 for that request and keeps the worker draining the queue.

2.6 L6 — Backend abstraction and device reach

Today. Backend is a Copy/Hash/Eq/Ord handle over a fixed id space (graph/registry.rs, F4) and a trait (graph/backend.rs:21-98). The ids are a KV-session file-format contract, and the name-keyed registry carries each backend's priority, capability matrix and pool hooks; consumers read it instead of matching. Before F4, adding a backend meant teaching nine #[cfg]-laden match sites about it: GraphAllocator::supports (alloc.rs:384-386), alloc_in_pool/alloc_fresh_in/free_in_pool (:819, :955, :959), sync_backend (alloc.rs:2400), copy_across (alloc.rs:2449), and the scheduler's execute match (scheduler.rs:274-292); the functions are registry-driven now.

GPU participation is decided by an all-or-nothing model-level gate: every weight must be registered on the GPU or the model runs on CPU (ARCHITECTURE.md §5.2; CUDA-BACKEND-DESIGN.md:251-255). The engine enumerates devices and honours --gpu N but uses exactly one (src/cuda/methods/init.rs:63-81); there is no tensor_split, no peer copies, and no remote-device backend.

Gap. 🟠 Two axes.

Pluggability — F4 (2026-09-24) landed the registry half: the backend set is data (a name-keyed table of entries registered at startup) rather than nine match sites, and a name that is unknown, not compiled into this build, or not usable on this machine is a loud startup refusal (BACKEND-REGISTRY-DESIGN.md). What is still closed at compile time is the set itself: there is no dlopen/plugin path, so a Vulkan backend (the only route to Windows/Linux AMD and Intel GPUs) or a remote/RPC backend is added by registering an entry — but it must still be compiled in. That is what remains of this axis.

Placement policy — a model that does not fit in VRAM could not run at all. The docs already record this as a real limitation on 8 GB devices (DEVICE-ADAPTATION-PLAN.md §9.2: "7B Q8_0 will not fit (7.2 GB weights alone)"). E5 (2026-09-23) added the missing granularity and the fit: CNode.layer (stamped by the model builders), an OffloadPlan (graph/offload.rs) that the loader's registration filter, the builder's device-only fused gates and the assignment pass all read, --gpu-layers N / MINFER_GPU_LAYERS=N, the startup report, and a verified mixed CPU+CUDA run on the 0.5B — then S2's auto: per-block weight bytes from the GGUF index, the pure fit_blocks prefix search against a weight budget (MINFER_GPU_MEM, else three quarters of the device's free bytes) with a quarter held back for the KV arenas and the activation pool. The registry/BackendId half is done (F4), and the free-bytes query for Metal landed as #53's DeviceMemory answer (recommendedMaxWorkingSetSize), so auto and headroom_bytes() work on macOS too.

Recommendation. (a) Introduce a BackendRegistry with register(Box<dyn Backend>), supports(op, dtype) -> Option<BackendId> and per-backend allocation/sync as trait methods; replace the enum with an opaque BackendId. Done in F4 (2026-09-24) — the shape that landed is a name-keyed Registry of BackendEntry values (priority + caps + the pool hooks, registered at startup) and a Backend handle over a fixed id space, with the assignment order made an explicit pinned number (docs/BACKEND-REGISTRY-DESIGN.md). It keeps the capability matrices in the trait's own modules (the trait methods forward to them), so the registry and the trait cannot disagree. (b) Add a layer-offload budget policy that consumes the memory accounting from §2.3.

2.7 L7 — Models, quantization, tokenizer, sampling

Architectures. Two (qwen2, qwen3 dense), dispatched on general.architecture (models/mod.rs:105-121). Every other family is absent: MoE, MLA, sliding-window/hybrid, state-space (Mamba/RWKV), multimodal, and every non-Qwen dense family (Llama, Mistral, Phi, GLM, InternLM). The per-family decision, its prerequisites and its cost live in MODEL-SUPPORT-ROADMAP.md; the caveat that matters here is §2.1 — its Tier 1 estimate holds only while a family needs no new IR structure. 🟠

Quantization coverage. Eight quant types; no Q2_K/Q3_K/Q8_K, no I-quants, no MXFP4/NVFP4 (SUPPORT-MATRIX.md §Not Yet Supported; CUDA_OPTIMIZATION.md §1.4 marks IQ/Q2/Q3 "not planned"). BF16 landed as a weight type: the writer and the CPU path in #142, the device halves (CUDA and Metal) in #208, 2026-10-06. The notable one still for reach is Q2_K/Q3_K (running large models on small machines). 🟠 but legitimately deprioritized.

Quantizer tooling — closed by F6 (#49, 2026-09-24), and its f16-weight gap closed by #141 (2026-09-25). A GGUF v3 writer (src/gguf_write.rs), byte-verified weight encoders (src/quantize.rs), a HuggingFace Qwen2 converter that satisfies the strict loader (src/convert.rs), and the convert / quantize / split subcommands (src/tooling.rs, docs/GGUF-TOOLING.md). A converted f16 model runs on the CPU and on CUDA — the CPU graph path gained f16 weight dispatch in F6 and #141 vectorized that dot (AVX2 F16C / NEON FCVTL) and moved the multi-token row loop onto the shared CPU pool (3.2 → 207 tok/s prefill on the 0.5B), while CUDA registers the raw 2 B/element weights and converts in-register. Both supported architectures now: #141's f16 registration and graph type gate landed in qwen2's loader/graph only, so #167 completed the qwen3 pair and gave the two loaders one shared registration rule (models::weight_reg, which also carries the q4_K W_dsc plane gate). Metal refuses f16 weights until #164, and a rewrite/split round-trip is bit-exact, and the HF conversion's f16 output is byte-identical to llama.cpp's converter on the same checkpoint, per tensor. What remains is scope F6 did not claim: the K-quant/I-quant encoders are still refused by name (#140) and bf16 output is refused (#142). 🟢 for conversion/quantization of the supported set.

CPU SIMD. AVX2 (Advanced Vector Extensions 2) dot kernels exist for Q4_0/Q8_0 (src/quants/dot_q4_0.rs:1-9, src/quants/dot_q8_0.rs:28), and since #56 (2026-10-08) the K-quants too: Q4_K/Q5_K/Q6_K × Q8_K have AVX2+FMA kernels (src/quants/avx2.rs) and AVX-512/VNNI variants (src/quants/avx512.rs), dispatched AVX-512 → AVX2 → scalar at runtime (MINFER_NO_AVX512=1 drops to AVX2, MINFER_NO_AVX2=1 to scalar — the x86 counterparts of MINFER_NO_NEON) and gated bitwise against the scalar reference by quants::avx2_correctness; Q4_1/Q5_0/Q5_1 remain scalar on x86 and NEON-only on aarch64. The support matrix is updated (SUPPORT-MATRIX.md, AVX2 column) and CPU_OPTIMIZATIONS.md §P1 is closed. What remains is weight repacking — the standard fix (repack the quant blocks at load time into SIMD-friendly interleaved layouts) that would expose the kernel gain end-to-end, since the current matmul re-reads each weight row per token; there is also still no aarch64 i8mm path and no AMX. 🟡 for x86 (the dots landed; the end-to-end win awaits repacking).

Tokenizer. Byte-level BPE only, loaded from GGUF metadata (tokenizer.rs). F7 (#50, 2026-09-24) made tokenizer.ggml.pre authoritative with two hand-written rule sets (qwen2, qwen35), replaced the silent unwrap_or(0) with a checked byte fallback, and made every other pre-tokenizer value (including a missing one), a non-gpt2 model, an empty merge table and an incomplete byte vocabulary a loud load refusal. What remains: no SentencePiece/unigram or WordPiece path, and no ignore_merges/multi-regex pre-tokenizers (llama3, default, deepseek-*, …) — those refuse by name, so the coverage gap is explicit rather than a wrong split. Also no NFC normalization (llama.cpp does not apply one for BPE either) and special-token matching is still a per-model hand-extension (the DeepSeek-R1 fix noted in AGENTS.md), which does not scale. 🟡 for the Qwen/Llama-3-shaped families this engine supports; 🟠 for anything SentencePiece-based. Follow-up: #132.

Chat templates. Landed in F7 (#50, 2026-09-24). minijinja 2.21.0's set_unknown_method_callback is the extension point, so Qwen3's chat_template renders (think-block extraction, tool-call formatting, enable_thinking) with no dependency change; the reference renderings are committed and byte-for-byte gated, and any template the engine cannot render is a loud refusal naming the construct and line — validated at load, so the CLI exits and the server refuses to start. What remains: a --chat-template <FILE> override (a user-supplied template for a GGUF that has none or a wrong one), strftime_now, and the Python str methods outside the implemented set (each refused loudly). Design + accepted/refused sets: docs/CHAT-TEMPLATE-AND-TOKENIZER-DESIGN.md. 🟢 for the supported models.

Sampling. One SamplerConfig pipeline in sampler.rs (#48, landed 2026-09-24): logit bias → penalties → DRY → greedy shortcut → top-k → typical → top-p → min-p → XTC → temperature or mirostat v1/v2, with every new knob defaulting to a no-op so the pre-F3 chain is bit-identical. Landed from the old gap list: min-p, typical, XTC, DRY, mirostat v1/v2, logit bias (and frequency/presence penalties from the OpenAI plan). Grammar / JSON-schema constrained decoding landed in F2 (2026-09-24): src/grammar.rs compiles a GBNF grammar (or a JSON Schema, via a generated GBNF) into a pushdown automaton and masks the logits inside that one pipeline, between DRY and the greedy shortcut; CLI --grammar/--json-schema, server response_format (json_object / json_schema) plus a grammar field. Design and the exact accepted subset: docs/GRAMMAR-DESIGN.md. What it deliberately does not do (and refuses loudly rather than guessing): pattern/minLength/maxLength, multipleOf, real-valued numeric bounds, object property permutations, and a device-side mask — the mask is host work before any backend runs. Still missing: top-n-sigma, adaptive-p, infill. 🟡

No LoRA (Low-Rank Adaptation) / adapter support. There is no adapter path at all. minfer's weights_version field in GraphParams (params.rs:88-98) exists precisely to break reuse on a weight swap, but nothing produces such a swap. 🟡

2.8 L8 — Ops coverage, safety, testing, observability

Op coverage. The Op vocabulary is deliberately small; six variants are parity-only stubs that no architecture emits (Scale, Softmax, View, Reshape, Permute, AttnMode::Mha — COMPUTE-GRAPH-DESIGN.md §1.3), and View/Reshape/Permute execute as copies. CUDA additionally refuses transpose_b matmul (cuda_backend.rs:1454-1459) and FusedQkvNorm (cuda_backend.rs:1905), and Metal refuses QkvBiasRopeStore (metal_backend.rs:700) — so the two GPU backends do not implement the same op set, and a model's decode path differs by platform. Neither is documented in SUPPORT-MATRIX.md. Fixed in A8: SUPPORT-MATRIX.md now carries an "Operator Coverage by Backend" table with the four asymmetric rows and their consequences.

Guard asymmetry. Both halves are closed on a Mac (2026-10-06). The decode-fusion debug_assert!s (#39, gap-table G2): FusedFFN / FusedQKV / FusedQkvNorm each return Err naming the node and the observed nt before the weight lookup, so a release build no longer passes an invalid shape through. The weightless-RMSNorm substitution (#40, gap-table G3): Op::RmsNorm / Op::QkNorm now call MetalBackend::norm_weight, which returns Err naming the node and the missing tensor for both None meanings (weight_name absent, or unregistered on the device) instead of running the weightless kernel where CUDA errors (cuda_backend.rs norm_weight). Neither violates the project's rule that "kernel-invariant violations return Err, never a silent fallback" (AGENTS.md §GPU Safety) any more.

Testing. Five integration files, four of them #![cfg(target_os = "macos")] Metal kernel isolation tests; CPU/CUDA correctness rests on inline unit tests plus real-model tests that skip when a GGUF is not cached (tests/conversation_cli.rs:1-12). CI is cargo build --release on macOS only — no test run, no CUDA build, no Linux build. Fixed in A2: CI now runs cargo test --release on Linux/CPU, compiles the CUDA backend in NVIDIA's devel image, and keeps the macOS build. 🟠 The remaining highest-value addition is a systematic op × dtype × backend correctness matrix (every op checked on every backend against a CPU reference); it would have caught several items in §4 automatically — that is ticket A1.

Observability. F8 landed 2026-09-24 (#51): GET /metrics (Prometheus text) exports live KV/arena occupancy, queue depth and the running/in-flight counts; per-op timing is available under MINFER_OP_TIMING (off by default, measured at the scheduler's per-node dispatch); and SIGINT/SIGTERM drains gracefully, bounded by MINFER_DRAIN_MS. Still missing: structured logging and levels (the server prints free-form startup lines), and worker-restart supervision (a panicking job is isolated per request by run_job_isolated, but a dead worker thread is not restarted). The trace/viz stack remains a development instrument, not an operations one. 🟡 for a research engine, 🟠 the moment it is deployed.

3. Recommendation backlog

Effort: S ≤ 2 d · M ≤ 1 w · L ≤ 2 w · XL > 2 w. Backlog items tracked on GitHub link their issue (the open list); an item whose work predates the tracker has no issue, and the plan is its record.

P0 — structural, unblocks a class of use cases

P1 — material capability or performance

P2 — coverage

P3 — hygiene and operations

4. Concrete defects found (verified in code)

Ordered by severity. Items 1–6 are behavioural; 7–12 are hygiene; 13–14 are behavioural defects found while executing the plan, already fixed.

1. GPU KV store bounds check — closed. A3 moved the check to the backend-agnostic GraphAllocator::fill_input_i32 (check_positions_bound), so CPU, CUDA and Metal all refuse an out-of-range row before execution, and #38 added the arm-level Metal guard (MetalBackend::check_kv_store_rows) for the fill paths that bypass it — §2.4 states the same closure in the gap list.
2. ensure_kv ignores a changed size (alloc.rs:995-1002). The KV region is frozen at first allocation while CParams.n_ctx remains part of the reuse identity, so a size change is neither honoured nor detected.
3. Metal weakens kernel guards. Both halves are closed: the decode-fusion debug_assert!s (#39) return Err naming the node and the observed nt, and the silent weightless RMSNorm when a weight is missing (#40, MetalBackend::norm_weight) returns Err naming the node and the missing tensor instead. Neither contradicts docs/GPU_SAFETY.md and AGENTS.md's no-silent-fallback rule any more.
4. Server worker had no panic isolation outside the forward call (src/server/chat.rs::guarded_forward_batch): a panic anywhere else unwound the worker, permanently degrading the server (503 for new jobs, empty 200/SSE for queued ones) with no log. Fixed in A4 — the whole per-job body is now guarded.
5. Backend op-set asymmetry drives silent path changes. FusedQkvNorm is Metal-only (MetalBackend::supports_op) but absent from CUDA's supports_op, so Qwen3 decode takes the fused path on Metal and the unfused path on CUDA. QkvBiasRopeStore is the mirror case (MetalBackend::supports_op returns false for it). Neither is documented in SUPPORT-MATRIX.md. Documented in A8: SUPPORT-MATRIX.md now has an "Operator Coverage by Backend" table; the asymmetry is visible rather than silent.
6. CUDA RoPE is non-interleaved only (cuda_backend.rs:1522), so any model needing the interleaved style splits every layer between CUDA and CPU, producing two host round trips per layer (§2.2); docs/ARCHITECTURE.md:455-459 now states that limitation rather than advertising both styles.
7. Stale unreachable!("CUDA pool not implemented") in the non-CUDA arms — the string and any unreachable! are gone from alloc.rs. Fixed in #57 (cdf41b2).
8. Dead fields in the reuse identity: CParams.n_batch and GraphParams.n_seqs are compared by params_match (cache.rs:95-101) but no builder reads them; every construction site hard-codes 1 / n_tokens. Fixed in A7, closed in E2: n_batch is deleted; n_seqs was kept for item 3 and then deleted by it (the sequence count is data — the topology decision lives in CParams.explicit_span), with sequence_count_is_data_not_topology pinning that a sequence-count change no longer rebuilds an otherwise identical graph.
9. Single-entry cross-backend staging (GraphAllocator::cross_staging), mitigated by the consumer-side filter at scheduler::cross_input_ready. Fixed in A5 — keyed by (graph uid, node, dst_backend); two foreign consumers can now be served.
10. read_host returns None on CUDA (cuda_backend.rs:2268), so the trait's host-read contract is backend-dependent; the allocator compensates with copy_to_host (alloc.rs:2137-2143).
11. CUDA pool never releases device memory (cuda_backend.rs:2148-2156), documented as accepted debt but reasoned about for fixed-shape CLI runs; a varying-n_tokens workload accumulates one buffer set per distinct shape.
12. Tests: no Linux/CUDA/CPU in CI, four of five integration files macOS-only (.github/workflows/ci.yml, tests/*.rs). Fixed in A2 for the CI half (Linux/CPU tests, CUDA compile, plus --no-run test compilation); the four macOS-only kernel-isolation files remain macOS-only by nature.
13. Op::Softmax returned unnormalised values on the CPU backend (cpu_backend.rs): it called vec_soft_max_f32, which writes exp(x - max) and returns the sum, and discarded the return. No architecture emits a standalone Softmax node, so nothing exercised it. Found by the A1 op matrix and fixed: the arm now scales by 1/sum, with an op-matrix case pinning it.
14. The CPU attention was not nt-invariant. Each token's scores were padded to the batch-wide nkv, and the softmax, the normalisation and the weighted sum all ran over nkv instead of the token's own causal window vl = pos+1. The padded entries are -inf → 0, so the arithmetic reads as equivalent — but the reduction length, and with it the rounding, depended on how many tokens shared the batch. Measured on 0.5B: a one-token decode step differed from a single-shot prefill by max|Δlogits| = 0.41, and the same token's K rows differed from layer 3 on (nt=6 vs nt=13). Found while testing prefix reuse for B2 and fixed: restricting the window to vl makes incremental prefill, the decode loop and a single-shot prefill bitwise identical, and removes the padding pass. It also matters beyond B2 — the speculative-decoding identity gates assume exactly this property.

5. Suggested execution order

The dependency structure is more informative than the priority table alone:

                       ┌──────────────────────────────────────────┐
                       │ 1. KV cell store (seq-addressable)       │
                       └───────────────┬──────────────────────────┘
                                       │
        ┌──────────────────────────────┼───────────────────────────────┐
        ▼                              ▼                               ▼
 2. IR seq_id + mask         4. persistent server context    21. quantized KV
        │                              │
        ▼                              │
 3. continuous batching  ◄──────────────┘   (4 alone already removes
        │                                    per-request realloc/re-warm)
        │
        ├──► 10. chunked prefill
        └──► 9. layer offload ◄── 8. allocator reserve + accounting

  7. IR views/multi-output ──► 17. MoE, 11. CPU SIMD (independent)
  12. backend registry      ──► 14. async copy ──► multi-device execution

Items 5, 6, 13, 26–28 are independent, small, and can land at any time — they are the ones worth doing first simply because they are cheap and they remove hazards that the larger work would otherwise have to work around.

Appendix — evidence index

minfer Architecture Execution Plan

Status: Phase A complete (9/9, 2026-09-16); Phase B complete (3/3, 2026-09-16); Phase C complete (8/8) — C1, C2, C3, C4 (quantized Q8_0 cache, CPU), C5 (session save/restore), C6 (logical positions), C7 (+C7b) and C8 (cross-sequence cell sharing) are done. C6 merged 2026-09-20 as 001b8cc; C7 landed 2026-09-20 (the partition is elastic, and growth moves runs in both directions, so a busy neighbour above the slot no longer blocks it); C8 split into C8a (shared prefill, duplicated rows: no IR change) and C8b (paged sharing: a block map and a gather in every attention kernel — S1a/S1b/S2/S3/S4/S5 landed, closed on all three backends: Metal's share path landed after the G5 port, in #362); C4 and C5 landed 2026-09-22 (C4's fused dots and the CUDA kernels are #87, Metal's packed read #310; the CLI/server surfaces C5 enables are #89). Phase D complete (3/3) (D1 done: views, multi-output via split_parts, D2, D3); Phase E complete (7/7) (E1, E1b, E2, E3, E4, E5, E6 all done); Phase F in progress (7/8) (F2, F3, F4, F5, F6, F7, F8 done; F1 waits for an x86 host); Phase G complete (7/7) — G1–G7 all landed on a Mac (the KV port G5 and the measurement bookend G7 last; #54 recorded the round's final baseline). Next: the open Linux-side tickets on dgxspark (aarch64, GB10 sm_121), beginning with #428's three fixture producers, and F1's weight-repacking increment, which still wants an x86 host. The order was deliberate: the Metal KV port (G5) came after the CUDA arena stopped changing shape (C7, C7b, C8), so those semantics were written into Metal once — it landed on a Mac 2026-10-06. Per-ticket evidence is in each phase's record and in the §14 open-risks table. Companion to: docs/ARCHITECTURE-ROADMAP.md (what is missing, why, and how it is ranked). This document is the how: phase-by-phase tickets with deliverables, acceptance criteria and dependencies. Baseline: HEAD = f32daa7 (2026-09-16); Phase A landed on architecture-phase-a (PR #1). This status was refreshed against master = 53b3ace (2026-10-06, the macOS/Metal round's final master); it is refreshed when a phase or round closes, not on every PR. Derived facts: the phase counters above, the next: sentence and the two commit ids are not hand-maintained prose — scripts/status.toml is the source of truth, and scripts/check_status.py --check (CI job check-docs) fails when this block disagrees with it, naming the file, the line and both values. Edit the source, not the counter. The suite counts are the sibling ledger scripts/test-baselines.toml, read by scripts/check_baselines.py: its --check-live compares the one CI-verifiable box against the test-linux-cpu log, while the CUDA / real-model / sanitizer rows are labelled recorded measurements because CI has no GPU. The per-ticket ✔ marks in the §11 diagram are not derivable from (done, total) and are deliberately out of scope for the checker.

Issue links. Work tracked on GitHub carries its issue link in its table row (prose sections carry it in the heading), and the record written when that work lands keeps the link. Tickets that predate the tracker — Phase A and B, C1–C3, C6/C7/C7b, C8a, D1–D3, E1/E1b/E2/E6 — have no issue, and their record here is the only one; the tickets below that do have one are exactly the open list in the tracker. Doc debt and test health are filed the same way: #62 (docs/USAGE.md staleness). The docs build was #63 and is closed: CI job check-docs builds the book with the same toolchain Docs deploys with and runs scripts/check_docs_links.py, so a renamed file now fails on the PR that renames it instead of rotting silently (it found four dead links the day it landed). Test health was #82 — the three #[ignore]d real-model tests that failed on master — and it is closed: two were writing into a directory nothing created, and the third asserted a model behaviour (the greedy 0.5B stopping on EOG within 16 tokens) instead of the engine's rule that ties need_insert_eot to the stream.

0. Decisions already taken

1. Standing rules

Every ticket is rejected if it breaks one of these. They restate the project's own invariants rather than inventing new ones.

1. Params-only reuse. Anything that changes graph topology must join GraphParams / CParams and be compared in GraphCache::params_match (src/graph/cache.rs:57-64); n_past never enters the identity. Decision: ADR-0002.
2. No silent fallback. A kernel-invariant violation returns Err from execute_node with the actual values; assignment is decided at build time. Decision: ADR-0009.
3. Identity gate. Any change to a kernel or an execution path is A/B'd against the existing path. Bitwise-identical is the default bar; where that is impossible, the ticket must name the tolerance class and its cause (the project's existing example: nt ≤ 8 bitwise, nt = 9 tolerance-class). Decision: ADR-0010.
4. GPU safety. Bounded waits, status checks, device limits queried at runtime, no hardcoded device constants. Decision: ADR-0008.
5. Index docs move with the code. README.md, AGENTS.md, docs/SUPPORT-MATRIX.md and docs/ARCHITECTURE-ROADMAP.md are updated in the same commit as the change they describe.
6. Deferred-Metal marking — retired. A ticket whose cross-backend design changed Metal had to add a Phase G line in the same commit. Metal is in scope since ADR-0005 (2026-09-20), so the rule has no subject.

2. Verification matrix (what dgxspark can prove)

dgxspark — the box every recorded measurement in this plan was taken on — is a DGX Spark (GB10), aarch64 Linux, CUDA 13.0 (/usr/local/cuda-13.0), with cached Qwen2.5-0.5B/7B/14B and Qwen3-0.6B GGUFs. Measured records name it absolutely, never as "this box" (gate contract rule 5).

A0 verdict (2026-09-16) — superseded (2026-09-18). It read "CUDA is compile-only in this environment", and for two days every CUDA ticket's acceptance was "compiles + reviewer-inspected". The reading was wrong: those probes ran under an agent file sandbox whose Landlock rules denied open() on /dev/nvidia* even though the nodes exist and are world-writable, so cuInit failed with err 304 for a reason unrelated to the driver. The device has been available since 2026-09-18, and CUDA work is now measured on it — E1b's windowed attention, E2's 1.9x, and the f16 windowed-prefill fix recorded in §14 row 0 all carry device evidence.

The design consequence worth keeping is about coverage, not capability: CI has no GPU (§14 row 1), so a device-gated assertion is a local, manual run. Where a guard can live in the backend-agnostic layer (the allocator) instead of cuda_backend.rs, put it there so CI still proves it — that reasoning is independent of whether a device happens to be present (applied in A3, and again by C3 below).

3. Phase A — instrument, then hazard removal

Why first: the roadmap's §5 calls items 5, 6, 13, 26–28 "cheap and independent". This plan pulls item 23 (op matrix) and item 24 (CI) to the front of that batch — they are the instrument that keeps Phases B–E honest, and one of them (the op/dtype/backend matrix) would have caught several of the roadmap §4 defects automatically.

A0 — CUDA access spike — SUPERSEDED (2026-09-18): the device is available

• Original verdict (2026-09-16): cargo build --release --features cuda succeeds (1m23s, targets sm_75…sm_121, PTX compute_121), but at runtime cudaGetDeviceCount returns err 304 and the engine logs CUDA: no CUDA devices found (cudaGetDeviceCount err 304, count 0) then CUDA: not available, using CPU fallback. The CPU path was unaffected (Qwen3-0.6B Q8_0: 120 tok/s prefill, 64.7 tok/s decode).
• Correction (2026-09-18). The agent's execution sandbox is enough to produce every symptom A0 recorded, on a device that works. Under the harness's default file sandbox (Landlock) every open("/dev/nvidia*") returns EACCES even though the nodes are crw-rw-rw-, so cuInit fails with 304 and NVML prints Failed to initialize NVML: Unknown Error — the exact signature A0 recorded. With the sandbox widened (2026-09-18, after a reboot that also cleared a driver-upgrade state the maintainer had flagged): NVIDIA GB10, sm_121, 121.6 GiB, driver 580.178.04, CUDA 13.0, cuInit → CUDA_SUCCESS. Every A0-era probe was therefore run inside a sandbox that cannot reach a GPU even when the GPU is healthy, so those probes could not establish anything about the driver's state: "no device" was unsupported rather than merely pessimistic. (The maintainer recalls a driver upgrade without a reboot at the time; that account and the sandbox are both consistent with the record, and only the sandbox is reproducible today — which is why the correction is to re-run the verification, not to assume it always would have passed.) What is certain now: the CUDA half of every ticket between A0 and this session was compile-verified only, and in this session it is device-verified (see the E1b record's device section).
• Consequence for the plan: tickets that were closed "compile-verified only" because of A0 are re-opened as verification, not as code: E1b's kernels, the A1 matrix's CUDA column, and the CUDA-side test suite all get their first execution on hardware below.

A1 — Op × dtype × backend matrix · item 23 · M — DONE

• Files: new src/graph/op_matrix.rs, registered as #[cfg(test)] mod op_matrix;. It needs the crate's internals (the allocator, the backends), so it lives in the module tree rather than under tests/ — the crate is a binary target, so a file under tests/ is a separate crate that cannot name the binary's items at all (there is no lib target to link against). The test must therefore be an in-crate module.
• Three tests:
  1. matrix_cases_match_their_reference — 17 cases (Add, Mul, Scale, Silu, SwiGLU, Softmax, RmsNorm, QkNorm, MatMul, GetRows, View/Reshape/Permute, RoPE, Attn, KvcacheStore/Load) run on every backend that claims the op, each compared against an analytic reference written in the test — never against another backend. Unavailable backends report SKIP (reason), never PASS: on dgxspark that is 17 CPU cells + 34 skips.
  2. support_table_matches_support_matrix_doc — supports_op for 23 op rows against the table published in SUPPORT-MATRIX.md, so the A8 doc and the code cannot drift. The CPU column is checked here; the Metal/CUDA columns check themselves wherever they are compiled in.
  3. every_op_has_a_matrix_decision — every Op variant is either covered or excused in EXCUSED. op_label matches with no wildcard arm, so adding an Op variant is a compile error until the matrix is updated.
• Found a real defect, fixed here: the CPU Op::Softmax arm called vec_soft_max_f32 (which writes exp(x - max) and returns the sum) and discarded the return, so the op produced unnormalised output. Nothing caught it because no architecture emits a standalone Softmax node. The arm now scales by 1/sum; verified by neutering that line (the Softmax cell goes red) and restoring it.
• Acceptance met: the matrix runs under cargo test; the A8 asymmetry table is now machine-checked; one real defect found and fixed; cargo test --release 162 passed / 0 failed.
• Left for a GPU-verifiable run: the Metal and CUDA columns, and the four GPU-only fused ops (FusedQKV/FusedFFN/FusedQkvNorm/QkvBiasRopeStore) which are excused on a CPU-only box.
• CUDA column executed on hardware (2026-09-18). With the sandbox corrected (A0) the CUDA column finally ran, and every cell it claimed passed: Add, Mul, Silu, SwiGLU, RmsNorm, QkNorm, MatMul, GetRows, View, Reshape, Permute, RoPE, Attn, KvcacheStore, KvcacheLoad (Scale/Softmax stay SKIP (Cuda does not claim …), which the support table asserts). Running it exposed four harness defects, all fixed here — the kind that only a device can show:
  1. the CUDA column silently depended on test order: CudaState::get() is None until something calls init(), which main does at startup and the model tests do through the loader, so a filtered run reported SKIP (no CUDA device) while a full-suite run used the device. The harness now initialises the state itself (in backend_claims and run_on).
  2. the synthetic weights were registered only on CPU (alloc.register_weight delegates to the CPU pool; in the product the loader registers each weight into CudaState), so every weight op failed with "weight not registered on CUDA". The harness now registers them into CudaState too.
  3. Attn (hd 2) and QkNorm (hd 2) used a head dim the CUDA kernels reject (must be a nonzero multiple of 4) — cells that could only ever pass on CPU. Both fixtures now use hd 4.
  4. the KvcacheStore/Load case wrote two of its four region rows and expected the rest to be zero: true for a fresh CPU pool, false for device memory. It now writes every cell, so the expectation is defined on every backend.
• Acceptance met for the CUDA column too (2026-09-18): on GB10 (sm_121) the matrix is 17 CPU cells + 17 CUDA cells with no FAIL in either column.
• Recorded limitation (2026-09-18): support_table_matches_support_matrix_doc compares the code against a mirror table hard-coded in Rust and only names SUPPORT-MATRIX.md in its failure message — it does not parse the markdown. A stale row label therefore survives a green suite (it did: the Attn row still said multi_seq after the field was renamed to explicit_span). Parsing the published table would make the drift impossible rather than merely visible; it is a hardening item for A1/A8, not a defect.

A2 — CI · item 24 · S — DONE (verified in CI)

• Files: .github/workflows/ci.yml.

• Deliverable: three jobs — test-linux-cpu (cargo test --release, the runtime net), build-linux-cuda (cargo build --release --features cuda in nvidia/cuda:12.8.0-devel-ubuntu22.04; no GPU on a hosted runner, so compile-only by necessity), and the unchanged macOS/Metal build.

• Acceptance: the workflow parses and declares all three jobs; a broken commit now fails test-linux-cpu (unit tests run there, including the allocator/scheduler tests added in A3/A5).

• Verified outcome (PR #1, run 35080755969):

  job	result	time
  test-linux-cpu	✅	1m09s
  build-macos	✅	1m26s
  build-linux-cuda	✅ (incl. cargo test --features cuda --no-run)	4m36s

• The x86_64 question is answered: the Linux/CPU suite passes. running 163 tests → 160 passed; 0 failed; 3 ignored, plus 3 passed; 6 ignored for the integration file. That is 2 fewer than the aarch64 dev box, and the difference is exactly the quants::neon_correctness module, gated #[cfg(all(test, target_arch = "aarch64"))] — no test is silently missing.

• First CI run failed, and that was the point. Run 35080525201 came back build-linux-cuda ❌ at the dtolnay/rust-toolchain step: curl: command not found — the CUDA devel images are minimal, so rustup could not bootstrap and the CUDA build never started. Fixed by installing curl ca-certificates build-essential before the toolchain step (the last also supplies nvcc's host compiler). Only the fix commit's push turned all three green.

• Container base moved (2026-09-30): the image tag above is now nvidia/cuda:12.8.2-devel-ubuntu24.04 — deliberately the same image the x64 release job builds in (.github/workflows/release-build.yml, linux-x64-cuda). The 22.04 tag was glibc 2.35 / gcc 11 against the artifact's 2.39 / gcc 13, so the compile check ran on a base no artifact was ever built on. The rest of this record stands as written.

A3 — KV bounds guard + ensure_kv size check · item 5 · S — DONE

• Files: src/graph/alloc.rs (only — see the deviation note).
• Deliverable, part 1: ensure_kv now records the element count and the backend each region was allocated for, and returns Err when a later graph asks for a different size or assigns the layer elsewhere. Before this, the early return silently handed back a wrongly sized region; the CPU backend then failed with a position error and the GPU backends wrote out of bounds.
• Deliverable, part 2 (deviation from the ticket as written): the pos < n_ctx guard was not added to cuda_backend.rs. Because A0 found CUDA unverifiable here, the guard went into GraphAllocator::fill_input_i32 — the single point where positions become graph data. It is structural (an I32 input consumed by a KV-writing or attention op is bounded by the graph's n_ctx, taken from the kv_load node), so it covers all three backends including Metal without touching metal.rs, costs one O(nt) host scan, and is fully testable on CPU. token_ids is deliberately exempt (vocabularies exceed n_ctx).
• Acceptance: kv_region_size_change_is_a_loud_error, position_beyond_n_ctx_is_rejected, token_ids_are_not_bounded_by_n_ctx (all fail before this change); cargo test --release 155 passed / 0 failed.
• Defers to G: nothing — the Metal path is covered by the same allocator guard. Phase G keeps only the style asymmetry (debug_assert! vs Err). Superseded in part by gap-table G1 (#38, 2026-10-06): the allocator guard covers the fill_input_i32 path, and the Metal arm guard (MetalBackend::check_kv_store_rows) was added for the fill paths that bypass it.

A4 — Worker panic isolation · item 6 · S — DONE

• Correction to the ticket as written: the worker did have a catch_unwind — guarded_forward (chat.rs:438-450) contains a panic inside forward_graph_cached. It is just too narrow: the speculative path calls both models' forwards directly (spec.rs:385/417/447/471/509), and the tokenizer, sampler, stop-string and streaming paths were unguarded too. One panic there unwound worker_loop, dropping job_rx (every later request is rejected) and the already-queued jobs' event senders (an empty 200 instead of an error).
• Files: src/server/chat.rs — run_job_isolated + panic_message, called around the whole per-job body in worker_loop; the inner guarded_forward stays for the better message on the common case.
• Deliverable: any panic in a job becomes a 500 on that request's stream, is logged, and the worker keeps draining the queue with the slot released.
• Acceptance: isolated_job_turns_a_panic_into_an_error_event, isolated_job_forwards_a_normal_error, isolated_job_passes_success_through_silently; cargo test --release 158 passed / 0 failed.
• Found while writing the test: panic_message(&payload) on a Box<dyn Any + Send> downcasts against the Box (which is itself Any) and always reports a non-string payload; &*payload is required. Verified against a throwaway rustc probe.
• Not fixed (needs a device): a panic while the CUDA capture window holds the process-wide stream mutex poisons it, so every later request would panic too — now contained per job, but the server would still fail every request. Recorded for the CUDA-verifiable phase.

A5 — Staging map re-key · item 27 · S — DONE

• Files: src/graph/alloc.rs (cross, copy_across, cross_buffer, plus a new write_pool helper that removes the duplicated four-arm write) and src/graph/scheduler.rs (the consumer-side filter).
• Deliverable: cross is keyed by (NodeId, Backend). The scheduler's .filter(|cb| cb.backend == split.backend) is gone — the map itself can no longer offer a consumer the other backend's copy, and one node feeding two foreign backends now gets one staging buffer each instead of a single entry that only one of them could use.
• Acceptance: staging_is_keyed_by_destination_backend (stages one node for CPU, asserts the CUDA consumer gets None, then stages CUDA and asserts the CPU entry survives); the existing scheduler tests still pass; cargo test --release 159 passed / 0 failed.
• Honest limitation: a behavioural test needs a second usable backend, which dgxspark does not have (A0: CUDA unavailable; Metal not compiled). The test asserts the new keying contract directly through a #[cfg(test)] hook rather than through a real split boundary. Phase G should re-test it on a Mac.

A6 — CPU per-op allocations · item 28 · S — DONE: measured, refuted, reverted

• Files (measured, not changed): src/graph/cpu_backend.rs — the per-node Vec<&[f32]> of resolved inputs, and the K/V source clone in the KV store arm.

• What was tried: replace the per-node input Vec with a fixed [&[f32]; 4] array (every op the builders emit has ≤ 4 inputs) plus a heap spill for anything larger — i.e. zero allocation per node execution instead of one.

• Measurement (Qwen2.5-0.5B Q4_0, CPU, 20 threads, bench -r 3, three interleaved before/after pairs, minfer-a6-before vs minfer-a6-after):

  pair	before pp512	after pp512	before tg128	after tg128
  1	76.08 ± 0.03	74.79 ± 0.03	40.30 ± 0.09	39.60 ± 0.28
  2	76.05 ± 0.01	75.33 ± 0.02	40.25 ± 0.16	39.07 ± 0.40
  3	76.14 ± 0.03	75.39 ± 0.05	40.00 ± 0.31	39.68 ± 0.29

  The change is consistently slower (−1.2 % prefill, −1.8 % decode; every pair separated, and the intra-run error is ±0.05 or less). Reverting restores the baseline (76.10 / 76.06 pp512, 40.79 / 40.21 tg128), which confirms the cause.

• Why it does not matter anyway: the decode loop is weight-streaming bound. At ~250 nodes/token the per-node Vec was one small allocation each, on the order of 0.1 % of the step — below what the harness can resolve. The K/V clone was left alone for the same reason: it is ~2 KB per layer and the buffers are needed for the borrow structure (the store arm needs disjoint access to four pool slots).

• Outcome: no code change; the negative result is recorded so nobody re-opens it. Item 28 is closed as "not worth doing" rather than done.

A7 — Dead identity fields · item 26 · S — DONE, closed in E2

• Decision taken: option (a) — delete CParams.n_batch, keep GraphParams.n_seqs marked reserved for item 3 (rationale in §8).
• Files: src/graph/params.rs, src/graph/cache.rs:57-64, src/graph/json.rs:40, and the n_batch construction sites in src/models/*/graph.rs.
• Deliverable: n_batch gone from CParams, from params_match and from the graph JSON export; n_seqs carries a comment naming the item that will read it.
• Acceptance: no occurrence of n_batch remains in src/graph/; every field compared by params_match has at least one reader; cargo test green.
• E2 closure (the second half of the ticket). Keeping n_seqs reserved was the right call only until item 3 landed. It did — and the field turned out not merely unread but redundant: the only topology decision a multi-sequence batch can force is the attention instantiation, and CParams.explicit_span, derived from the KV reservations (the authority on where each window starts), already carries it. E2 therefore deleted GraphParams.n_seqs, satisfying the acceptance by the "or deleted" branch. Measured justification (sequence_count_is_data_not_topology, models/qwen2/graph.rs): with the field present, a 2-sequence batch and a 1-sequence batch with the same n_tokens/n_out/gtype/explicit_span rebuilt the graph (uid 3 → 4); without it the same graph is reused and its logits are bitwise-identical to a fresh single-sequence forward. The test is permanent, so the claim cannot silently regress.

A8 — Guard symmetry · item 13 · S — DONE (docs route)

• Files: docs/SUPPORT-MATRIX.md — a new "Operator Coverage by Backend" section.
• Route taken: the ticket offered "CUDA gains FusedQkvNorm or the matrix records that it does not have it". A0 made CUDA unverifiable here, so writing a new CUDA kernel blind is strictly worse than documenting the asymmetry: the table now lists every op against CPU/Metal/CUDA, generated from the three supports_op implementations, with the four asymmetric rows called out and their consequences (Qwen3 decode is fused on Metal and unfused on CUDA; QkvBiasRopeStore is the mirror case; interleaved RoPE is CPU-only).
• Acceptance: the roadmap §4 defect 5 asymmetry is now visible in the support matrix; A1's matrix must agree with this table.
• Defers to G: the Metal half (debug_assert! → Err; the weightless RMSNorm fallback) and any decision to port FusedQkvNorm to CUDA. (The Metal half landed: debug_assert! → Err in #39 and the weightless RMSNorm fallback in #40 — both gap-table G2/G3. The FusedQkvNorm-to-CUDA question is #52.)

4. Phase B — persistent server context

Why second: it is independent of the KV redesign, and it is the first user-visible win — a chat client that resends the conversation stops re-prefilling it every turn.

B1 — Reproduce contamination — DONE: the documented mechanism does not reproduce

• What the docs claimed. chat.rs reset the slot cache on every request because "re-prefilling a DIFFERENT prompt over the same regions leaves stale rows inside the new attention window" — the message of 39eceaa (doc 97), which fixed a cross-request KV bug and added the reset.
• What was tested. reused_cache_across_prompts_matches_a_fresh_cache (src/models/qwen2/graph.rs) runs real prompts through one GraphCache and compares bitwise against a virgin cache, in three orderings:
  1. long prompt → short prompt (the stale-row case);
  2. short → long (the append case B2 wants);
  3. prefill → 3 decode steps → shorter prompt (the server's actual sequence, including the generated rows the comment is about). All three are bitwise equal.
• Why. A prefill writes rows 0..nt and attention reads [0, max(pos)+1) = [0, nt) — i.e. only rows this request just wrote. Stale rows survive above the window, where nothing reads them. The original bug is consistent with the state recorded in the OpenAI plan's revision notes: the cache was process-global at the time, so two slots shared one set of regions. That is already fixed by per-slot caches; the reset was belt-and-braces whose stated mechanism does not hold on the current tree.
• Residual uncertainty (why the reset stays until B2). This box can only run the CPU path. The fused GPU stores (FusedQKV, FusedQkvNorm, QkvBiasRopeStore) write K/V inside their own kernels and are unverified here (A0: CUDA compile-only; Metal not built). The chat.rs comment now records the finding and this caveat instead of the unverified mechanism.
• Consequence for B2. Reuse is safe by construction if it is gated on an exact token-prefix match: every row read is then a row whose contents were verified to be the same tokens. That is what B2 implements.

B2 — Retention + prefix reuse — DONE

• Files: src/server/slot.rs (Slot.cached_tokens), src/server/chat.rs (common_prefix_len, prefill_span, the suffix prefill, the per-token record, worker_loop no longer resets the cache).
• Deliverable: the slot keeps its GraphCache and a record of the token sequence its KV rows hold; generate_seq reuses the cache only when the new prompt starts with exactly that sequence (else it prefills from position 0), and always feeds at least the last token because its logits seed the sampler.
• Why it is safe by construction: the reuse gate is the verification — every row attention reads is a row common_prefix_len just proved to hold the same token. It is also bitwise: the CPU attention is now nt-invariant (item 14), so feeding a suffix at its own positions reproduces a single-shot prefill exactly.
• Deliberate limits: the speculative path neither reuses nor records (a verify round writes rows past the committed tokens); a panicking job clears the record; MINFER_NO_PREFIX_REUSE=1 restores the pre-B2 behaviour for A/B.
• Tests: common_prefix_len_finds_the_exact_match, prefill_span_always_feeds_the_last_token (pure), and prefix_reuse_matches_a_full_prefill (real model, bitwise).
• Not covered here: the fused GPU stores are unverified on dgxspark, so the gate stays conservative; Phase G re-tests on Metal.

B3 — Measurement — DONE (interleaved A/B, same binary)

• Setup: serve --n-ctx 1024 --n-slots 1 on Qwen2.5-0.5B Q4_0, one long system prompt (203 tokens) plus a second turn that appends an assistant reply and a new question (219 tokens), max_tokens=1 so the wall time is essentially time-to-first-token. Four alternating runs; after vs before differ only by MINFER_NO_PREFIX_REUSE=1.

  mode	turn	prompt tok	fed	reused	wall
  after	1	203	203	0	2.66 s
  after	2	219	16	203	0.24 s
  before	1	203	203	0	2.60 s
  before	2	219	219	0	2.71 s
  after	1	203	203	0	2.66 s
  after	2	219	16	203	0.26 s
  before	1	203	203	0	2.88 s (turn 2)

• Result: turn 2's prefill drops from 219 to 16 tokens and its time-to-first-token from ~2.7–2.9 s to ~0.24–0.26 s — ≈ 11× — while turn 1 (cold slot) is unchanged, as it must be. The per-request line [server] prefill fed N/M prompt tokens (R reused) is what the numbers come from, and it stays in the server so the behaviour is observable in production.

5. Phase C — KV cell store (item 1)

Five sub-steps, each keeping the tree green. C3 depends on Phase D.

C1 — Cell store, one sequence — DONE

• Landed: src/graph/kvcache.rs (KvCache, KvLayer, SEQ_MAIN/FREE, cells_for, is_identity, and the ownership/write bookkeeping C2 will use); the allocator's kv map is now the store, so ensure_kv validates through it (it gained the n_ctx argument) and kv_pair/copy_kv_to_cpu read it.
• The gate has teeth: BackendScheduler::execute refuses any graph whose mapping is no longer the identity, because no backend consumes the resolved cell array yet — so a half-ported C2 fails loudly instead of writing the wrong row. a_non_identity_kv_mapping_is_refused proves it fires.
• The resolver has a live consumer today: A3's position guard (fill_input_i32) now resolves through cells_for instead of re-deriving the bound from a node's shape, so when C2 changes the mapping the check keeps meaning the same thing without being touched.
• Verified bitwise, as the ticket requires: cargo test --release 172 passed / 0 failed — including the real-model bitwise tests (reused_cache_across_prompts_matches_a_fresh_cache, prefix_reuse_matches_a_full_prefill, graph_logits_match_forward_real_model) — and a pre-C1 vs post-C1 binary A/B on Qwen2.5-0.5B Q4_0 greedy (-n 24) produced byte-identical generated text; only the timing lines differ.
• Deliverable (as specified): per-layer cell arenas with an owner set; the KV store node resolves position → cell on the host and passes an index array to the kernel; attention still derives its bound from positions, so behaviour is unchanged.
• Delivered in this pass: the arenas, the owner set and the host-side resolver. The index array reaching the kernel is deliberately not done — it is only needed once the mapping stops being the identity, and the scheduler gate refuses to execute in that state, so no backend can silently index the wrong row in the meantime. Handing the array across the Backend trait was planned as C2's first step; C2 landed without it because it removes rows physically and therefore never leaves the identity mapping (see the C2 record below) — so the hand-off stays available, not urgent.
• Acceptance (as specified): bitwise-identical greedy output vs the pre-change binary on the full smoke set (0.6B Q8_0, 7B Q4_K_M, 14B Q4_K_M) at several context lengths; the graph topology is unchanged (no new topology-affecting params).
• Acceptance met for: 0.5B Q4_0 byte-identical end to end, plus the in-tree bitwise model tests. Deferred at the time: the 0.6B/7B/14B smoke rows — minutes-long CPU jobs and C1 does not touch a kernel, so they were left to the C2 landing rather than claimed here. Discharged at C2 (2026-09-16): the pre-C2 (C1) binary and the C2 binary generate byte-identical text greedily (-n 16, -t 8, single-shot) on Qwen3-0.6B Q8_0, Qwen2.5-7B Q4_K_M and Qwen2.5-14B Q4_K_M; only the timing lines differ.
• Defers to G: nothing — the Metal backend keeps the old regions until G.

C1 design (written before the code, 2026-09-16). The shape below is what C2 needs, so C1 builds it rather than a type with no consumer:

1. KvCache (new src/graph/kvcache.rs) owns, per layer, the two persistent regions plus owner: Vec<SeqId> — one entry per cell, FREE for a cell no sequence holds. One implicit sequence for now, so every cell of a written row is owned by it.
2. KvCache::cells_for(positions: &[usize]) -> Result<Vec<u32>, String> — the host-side resolver. Today it is the identity (cell == pos) and returns Err for a position outside the arena; C2 changes exactly this function to consult owner and the layer's window start.
3. KvCache::is_identity() -> bool — true until C2 introduces a hole or a window. The scheduler consults it: while true the existing positions input reaches the backend exactly as today (so no kernel changes), and once false the resolved cell array must be passed instead. A backend that has not been ported returns Err instead of indexing the wrong row (standing rule 2). On dgxspark that means CPU first: CUDA was compile-verified only at the time (A0 — superseded 2026-09-18, the device is available) and Metal stays untouched (G), so C2 must either keep the mapping identity for them or refuse to run there.
4. The index array travels as a graph input (kv_cells, I32), filled by the allocator from positions — positions stay data, so the topology and the params-only reuse identity are untouched.
5. Tests: resolver edge cases (out-of-range → Err, identity while there are no holes), is_identity flipping exactly when C2 introduces one, and the existing model tests (B1, B2's prefix reuse, graph_logits_match_forward) staying bitwise — that is the refactor's acceptance gate.

The first thing to change is therefore the resolution, not the kernels: the Backend trait gains the resolved cell slice alongside kv_pair, and each backend uses it for row indexing while still using positions for the causal bound and RoPE.

C2 — Removal and shift — DONE

• Deliverable: seq_rm (drop a range) and seq_add (shift positions), surfaced as graph-level operations; conversation.rs overflow switches from "drop the oldest turns + full re-render" to a context shift.
• Acceptance: the conversation overflow test passes; a turn that overflows no longer re-prefills the whole conversation; the shift path is a named tolerance class with the reason recorded (positions change ⇒ RoPE inputs change ⇒ not bitwise).
• Defers to G: n/a.

C2 record (2026-09-16)

What landed. GraphAllocator::kv_rm(start, len, &KvRope) removes KV rows [start, start + len) from every layer and re-bases the rows after them by -len, re-roping their K; kv_shift(drop, rope) is the start == 0 case (KvCache::after_rm / after_shift do the bookkeeping). rope_shift_kv now takes a signed delta so a rotation can be undone (the tests use that).

The removal is physical, which is the whole point: cell == pos survives, so is_identity() stays true, C1's scheduler gate stays satisfied, and no backend gains a kernel. The work is a host-side memmove plus one re-rope pass per layer, both through the existing copy_kv_to_cpu / write_host pair — which is also why CUDA/Metal need no change here (copy_kv_to_cpu has no Metal arm, so a Metal session falls back loudly rather than shifting wrongly; the port is G5).

Engine::kv_rm(start, len) is the conversation-facing hook (Err by default, so a mock or a non-shiftable backend is handled explicitly). Conversation plans the drop on a copy of the message list, resolves the dropped turn's token span from the chat template, and verifies both boundaries against the KV stream before touching anything; an unverifiable boundary or an engine that refuses falls back to the exact drop-and-re-render path, with the reason logged. MINFER_NO_CONTEXT_SHIFT=1 forces that path.

The retained rows keep the context they were computed with. This is the ticket's named tolerance class, and its cause is not rounding. Re-roping is exact to the rope tolerance (measured max|Δ| = 2.38e-7, relative 1.14e-7, by rope_shift_matches_roping_at_the_new_position), but a shifted row's value is whatever it was when it was written — it attended to the tokens that were later dropped. Measured on Qwen2.5-0.5B Q4_0 with the a/b/c probe (drop a, keep b, continue with c):

Last-token logits after continuing with c: max|Δ| = 16.75 vs a fresh b+c prefill, and the argmax changes. This is inherent: recomputing the retained rows is exactly the re-prefill the shift exists to avoid. llama.cpp's context shift (llama_kv_cache_seq_rm + seq_add + llama_kv_cache_update) has the same property. What C2 guarantees instead is exactness of the mechanism, and that is what the tests assert bitwise:

• removing a tail range leaves the retained head byte-identical, so continuing from it is bitwise equal to a fresh prefill of that head (kv_rm_is_exact_and_the_window_shift_is_a_named_tolerance_class);
• removing a middle range copies V byte-for-byte, leaves [0, start) untouched, and moves K by exactly the re-rope (undone to < 1e-4 per layer);
• removing everything written empties the arena, len == 0 is a no-op, and removing past the end is an Err, not a silent truncation.

Model-path A/B (discharging C1's deferral). The pre-C2 binary (HEAD at C1) and this build produce byte-identical generated text greedily (-n 16, -t 8, single-shot, CPU) on Qwen3-0.6B Q8_0, Qwen2.5-7B Q4_K_M and Qwen2.5-14B Q4_K_M — only the timing lines differ — which is the evidence that adding the removal path did not perturb the model path.

Conversation measurement. context_shift_real_model_measurement (0.5B Q4_0, n_ctx = 192, 12 turns with a system prompt) overflows on turns 8–11 and shifts on every one:

turn 8:  dropped 22 KV rows at 17 (2 messages), prefill 14 tokens instead of 185
turn 9:  dropped 23 KV rows at 17 (2 messages), prefill 14 tokens instead of 180
turn 10: dropped 23 KV rows at 17 (2 messages), prefill 14 tokens instead of 175
turn 11: dropped 23 KV rows at 17 (2 messages), prefill 14 tokens instead of 170

— a ~13× cut in prefill tokens per overflowing turn, start = 17 (the system prompt is retained and not even re-roped), no full re-render, and the answers stay correct on the shifted window (Stockholm,/Athens,/Warsaw,/ Lisbon, for Sweden/Greece/Poland/Portugal). prefill_tokens is the observable that makes "no longer re-prefills the whole conversation" checkable, so the ticket's acceptance is a test, not a claim.

Two pre-existing bugs the C2 tests found and fixed (both outside the ticket but blocking it, and both user-visible in --cnv):

1. The first user_turn never prefilled pre-existing messages — it wrote only the new message's delta, so the --system prompt never reached the KV while still being rendered into every later delta. The first turn now prefills the whole canonical render (first_user_turn_prefills_the_system_prompt).
2. A pending EOT was written before the overflow check, so a full context with a reply that had not reached EOG called forward at position == n_ctx and tripped the KV-region guard. The EOT is now deferred into the overflow path (it belongs at the end of the stream, so it commutes with the removal), and ContextFull keeps it pending for the next attempt.

Not done here. The logical mapping (a hole or a window with cell != pos, which is what would let a backend keep the old rows in place) is deliberately still refused by C1's scheduler gate — nothing consumes the resolved cell array yet, and the physical removal makes it unnecessary for the sliding-window case. C3's defragmentation, C4's quantized KV and C5's state save/restore are untouched. The server path is untouched too: a full slot is still reported as finish_reason: "length" and relies on B2's prefix-matched reuse, so a server-side shift policy (and the n_keep-style rule that protects a system prompt) belongs with E2's batching work.

C3 — Defragmentation

• Deliverable: a cell-copy operation that compacts the arena; triggered when fragmentation exceeds a threshold.
• Deps: D1 (a copy needs either a view or an explicit copy op).
• Acceptance (as resolved, 2026-09-19): node-count and arena-utilisation counters before/after; the moved bytes are exact — V verbatim and K exactly rope_shift_kv(old, delta), both asserted — and a mid-session compaction leaves the continuation's greedy token intact (asserted; a wrong re-rope flips it). Bit-identical logits were not claimable before C6 (a cell move shifted every RoPE angle, so the tail moved by the amplified rounding of that shift — §14 row 9 carries the four probes that attributed it). C6 landed the construction that removes it: positions are sequence-relative and the allocator resolves cells, so a compaction now changes no angle and the continuation is bitwise (gated by a_compaction_between_steps_keeps_the_continuation and the inverted offset tests).
• Follow-up: C6 does exactly this — sequence-relative positions plus an allocator-resolved cells input — which makes a cell move change nothing (a token's RoPE angle is its index within its sequence) and drops C3's K re-rope. When C6 lands, this ticket's acceptance tightens from the named amplified-rounding class back to bit-identical logits.

C3 design (written before the code, 2026-09-19)

The problem the arena has. KvCache::reserve_seq is first-fit over [0, n_ctx): it needs a contiguous free run of cap cells, while release_seq frees whatever run a sequence held. Two sequences that finish out of admission order therefore leave a hole in the middle, and the next admission can fail with no free run of N cells while the arena has far more than N free cells. That is fragmentation, and it is the one failure mode of E2's per-sequence reservations that no amount of capacity fixes.

Scope. C3 compacts downward: live runs are packed from cell 0 in ascending old-start order, each keeping its cap, so free space coalesces at the top of the arena. Compaction never changes a run's cap, never reallocates a KV region (the copy is within the same K/V buffer) and never changes a query's logical position — but it does change the cell a sequence's rows live in, so the caller that reserved the run must be told (the contract below).

Where the copy executes. The K/V regions are backend buffers (host memory on CPU, device memory on CUDA), so the copy is a backend operation: Backend::copy_cells(&mut self, dst: BufRef, src: BufRef, rows, elems_per_cell) — a new trait method, implemented by CPU and CUDA in this ticket and refused loudly by Metal (Phase G), which keeps standing rule 2: a backend that cannot move cells fails the request instead of skipping the compaction.

The contract is dst_row <= src_row (downward only) and the implementation must be correct for overlapping ranges, which the two obvious implementations are not:

• CPU: copy_within (memmove semantics) is exactly right.
• CUDA: device-to-device cudaMemcpyAsync is documented undefined for overlapping ranges, so the kernel is one block that walks rows in ascending order with a __syncthreads() between rows. Ascending is the safe direction when dst <= src (the row a write could clobber is already copied), and a row's elems_per_cell (<= n_kv_embd) parallelises across the block's threads. No staging buffer and no second pass.

The planner is pure. KvCache::compaction_plan(need) is a host-side function over the run table with no backend calls: it returns the moves (seq, from, to, rows) that compact the live runs, or — with Some(n) — the shortest prefix of that plan which opens a free run of n cells (None compacts fully). rows is the sequence's written length (last owned cell + 1 - start), not cap: unwritten cells hold no data worth copying, and moving them would shuffle another sequence's stale bytes. Because the planner is pure, its policy is unit-tested without a device — that is the half CI can prove.

The contract with the caller. Only whoever reserved a run caches its start: E2's server keeps one per slot (SlotState.start) and passes start + current_pos as the store position, so a move must be reported. GraphAllocator::kv_defrag(need) returns the applied moves plus before/after stats, and BatchEngine applies the new start to the slots it owns; everything else follows, because the graph's window (attn_span) is derived from seq_slot on every forward. A caller that ignores the report is not silently wrong in a hard-to-find way: its next store would write to the old cells and the span check would reject the query.

Trigger. GraphAllocator::kv_reserve_seq_with_defrag retries once through a compaction when first-fit fails, unless MINFER_NO_KV_DEFRAG (presence-checked) disables it — the A/B gate standing rule 3 requires: the same workload must produce identical output with and without the compaction, and the counters must show what it bought. The plain kv_reserve_seq stays pure, so a caller that keeps no run bookkeeping (the single-sequence path, the graph tests) is unaffected; the defrag-capable variant returns the moved runs precisely because a caller that does keep one (E2's server) has to follow them.

Counters (the ticket's other deliverable). KvArenaStats carries n_ctx, reserved/owned/free cells, the number of free runs (the "node count" C3 compacts), the largest free run, the live run count, and cumulative defrags/cells_moved. The acceptance test records them before and after; the same struct is what F8 (observability) will export.

Increments.

1. Planner + counters + KvCache bookkeeping (apply_moves) with unit tests — pure host-side logic, so CI proves it.
2. Backend::copy_cells (CPU + CUDA), GraphAllocator::kv_defrag, the reservation retry and the MINFER_NO_KV_DEFRAG gate.
3. Integration: fragment the arena on purpose (--n-slots N, prompts of different lengths) so an admission needs the compaction, and assert the server's replies stay bit-identical to the serial path, plus a device run for the CUDA copy.

Deferred (recorded, not silently skipped). Compaction does not grow a sequence's run, does not evict, does not move a sequence up to make room for a larger one, and does not run inside a forward (it is a between-forwards, host-driven operation, so CUDA Graph capture is unaffected). Compacting the arena per layer separately is also out of scope: one plan moves every layer together, which is what keeps owner[] consistent across layers.

C3 record (2026-09-19) — increments 1 and 2

Landed: the planner, the counters and the bookkeeping (increment 1), and the copy path (increment 2) — Backend::copy_cells on CPU (copy_within, i.e. memmove) and on CUDA (a new kv_move_rows kernel: one block, ascending rows, a barrier between them, no staging buffer), Metal refusing loudly (G5), GraphAllocator::kv_defrag + kv_reserve_seq_with_defrag + the MINFER_NO_KV_DEFRAG gate, and the server applying the returned moves to its slot starts.

Evidence:

• CPU end to end: kv_defrag_moves_the_bytes_and_opens_the_run fragments a real 16-cell arena (K/V regions allocated by a kvcache_store graph), asserts the 8-cell reservation is refused, compacts, then checks the bytes at the new cells, the counters (free_runs 2 -> 1, defrags/cells_moved), and that the refused reservation now fits — and repeats the fragmentation so the retry helper has to hand the moves back. CPU suite: 207 passed / 0 failed / 5 ignored.
• Device: cuda_copy_cells_moves_overlapping_rows_down moves rows [1, 4) to [0, 3) — two of the three rows are read and overwritten — compares the whole buffer, and checks that the downward-only at the time (C7b later added the upward direction — see §5) contract is refused before a launch. CUDA suite: 255 passed / 0 failed / 5 ignored.
• Coverage split, stated rather than implied: the allocator's CPU arm is proven end to end and the CUDA primitive is device-proven, but "the allocator drives the CUDA copy" is glue that only increment 3's integration run exercises.

Increment 3 (2026-09-19) added the part that makes a compaction correct rather than merely well-bookkept, and the model-level gate:

• A compaction must re-rope the K rows it moves. positions are cells today, so RoPE angles are absolute: a K row copied from cell from to cell to still carries the angle of from, and the next decode step attends to it at the wrong relative angle. kv_defrag(need, rope) now re-ropes the moved K rows through the model's own rope_shift_kv — the C2 path, one host pass per moved run, with delta = from - to (C2's convention: rope_shift_kv(d) means new_pos = old_pos - d). V carries no rotation and moves verbatim.
• The gate caught its own sign bug. a_compaction_between_steps_keeps_the_continuation (real 0.5B: prefill + one step at a non-zero start, release the holder below it, compact, then one more step) first failed on its argmax assertion because the delta was written as to - from; the allocator's unit test had passed anyway, because it computed its expectation with the same wrong sign. That is the whole argument for a behavioural gate next to a byte-level one.
• The offset-alone effect, measured while gating. The control run differs from the compacted one only in the subject's cell offset, and that alone moves the logits by 2.6% relative (0.43 absolute) on the 0.5B before any compaction; the compaction's own contribution is the remaining ~0.2pp. A uniform RoPE shift is attention-neutral and V is unrotated, so an unidentified op is offset-sensitive. This is now §14 row 9 — it is why the C3 acceptance is a surviving greedy token plus byte-exact K/V rather than bit-identical logits, and it is the floor the compaction test prints.

Coverage: the planner, the counters and the bookkeeping are unit-tested; the compaction's bytes, ownership and counters are asserted end to end on CPU (kv_defrag_moves_the_bytes_and_opens_the_run, with the K re-rope computed independently through rope_shift_kv); the CUDA copy primitive is device-proven (cuda_copy_cells_moves_overlapping_rows_down); and the continuation gate runs the real model. Still open: driving the compaction from the server's dynamic runs (today every slot is reserved once at startup and packed, so the trigger cannot fire there), and the logical-positions change that would remove the re-rope and the offset sensitivity altogether.

C4 — Quantized KV cache, Q8_0 first · #42 — DONE (CPU: S1 + S2a; CUDA: S2b) 2026-09-24

Why. f32 (and the GPU's f16-in-an-f32-region) is all the cache could be, so context length was bounded by KV memory with no way to trade precision for it. Q8_0 is the first packed layout: one cell is ceil(n_kv_embd/32 * 34) bytes instead of 4 * n_kv_embd, which is a real footprint reduction — unlike f16, whose regions stay f32-shaped and buy bandwidth, not memory.

What S1 landed (src/graph/kvformat.rs is the new authority):

• KvFormat { F32, F16, Q8_0 } with a strict MINFER_CACHE_TYPE parser and a device policy (resolve, pure so CI covers the matrix): an unknown value is refused on every device — CUDA used to read anything that was not f16 as f32, which is exactly the silent fallback the ticket forbids — and q8_0 is refused loudly on CUDA and Metal, whose attention kernels address f32/f16 rows (both since lifted: CUDA in C4 S2b below, Metal in #310). f16 on CPU stays the documented f32 (a process whose env is set for a GPU run must not fail a CPU one); the load path (models::load_model_ns) turns any refusal into a failed load with the reason printed.
• Packed rows stay addressable by the C3/C8b machinery. A cell is rounded up to a whole number of f32 words, so elems / n_ctx is still one cell's width and copy_cells (the copy-on-write shift, the compaction) keeps moving rows verbatim with no change at all. KvcacheMeta::row_elems carries the cell width, ensure_kv sizes the region by it, and the node's shapes stay logical ([n_kv_embd, n_ctx]), which is what the store's K/V input means.
• CPU store quantizes, CPU attention reads the packed blocks — S1 dequantized the window into a scratch; S2's fused read is below. The CUDA/Metal kernels are still issue #87.

Acceptance, as measured (CPU, cargo test --release, 2026-09-22):

• Footprint: the real-model gate prints the persistent regions, f32 against q8_0 — 0.5B (n_kv_embd 128): 6 291 456 B → 1 671 168 B (3.76× smaller); Qwen3-0.6B Q8_0 (n_kv_embd 1024, head dim 128): 58 720 256 B → 15 597 568 B (3.76×).
• Tolerance class (never bitwise — the store rounds every K/V cell): at the reference's argmax the logits move by ≤ 1.0 (measured 0.60 on the 0.5B, 0.55 on Qwen3), and over the whole 152k-way vector by ≤ 3.0 (measured 2.50, ≤ 8 % of the 37.8 spread). The greedy continuation agreement is reported, not asserted: the CLI diverges at the 5th token on a chat-templated 0.5B prompt, which is the format's honest cost, not a defect.
• The format itself is pinned one level down: a_packed_kv_region_answers_like_the_f32_one_and_is_smaller asserts a stored cell is bitwise the Q8_0 quantizate of the row it was given (and that the f32 region holds the row verbatim), plus the 3× footprint and an attention comparison over three causal queries. The gate was mutation-checked: disabling the packed read path turns that comparison into max |Δ| = 3.2e38.
• Refusals: MINFER_CACHE_TYPE=q8_0 on dgxspark (CUDA available) ends the load with minfer: MINFER_CACHE_TYPE=q8_0 is not supported on cuda yet: …, and a typo (banana) with … is not a KV cache type (f32, f16, q8_0); refusing rather than silently running with f32. ensure_kv backstops both: a packed width on a non-CPU backend, and a width Q8_0 cannot express, are Err where the region is sized. (Superseded: CUDA stops refusing it in C4 S2b below, Metal in #310.)

Coverage, stated rather than implied. The packed layout, the parser/policy matrix, the store-exactness check, the refusal and the region accounting are unit-tested and run in CI. The real-model gate is #[ignore]d (at this record's time it flipped a process-wide KV policy; since #99 the format is per engine and it no longer mutates shared state — see the #99 record below) and was run on both cached models. Not verified here: any packed region on a GPU (by design, S1 refuses it), and the fused dots' performance (S2).

C4 S2 — the fused read, the quantize-aware shift, and what is handed off (2026-09-23).

#87 named three items. The first two are CPU-only and land here; the GPU kernels are the follow-up at the end of this record.

• The fused read (attn_heads_q8, src/graph/cpu_backend.rs). S1 dequantized the union of the batch's windows into a reusable f32 scratch — two allocations plus a write-then-read pass per attention node, per layer, per forward — and ran the unchanged f32 kernel over it. S2 reads the packed blocks directly: the K score is dot_q8_0_q8_0 between the stored K blocks and the Q8_0-quantized query row (llama.cpp's form for a quantized cache), and V accumulates out of the cell block by block (kvformat::accumulate_q8_0_row). No scratch, no second pass. MINFER_NO_FUSED_Q8_KV (presence-checked) restores S1's path — the A/B standing rule 3 asks for — and the S1 path stays as the fallback for a hd_kv narrower than one Q8_0 block.
• The quantize-aware shift. kv_rm/kv_shift used to refuse a packed region because they re-rope K in f32 and a packed cell is not f32 rows. S2 keeps the property S1 already relied on — a cell is a whole number of words, so the survivors move verbatim, which is why copy_cells and the CoW/compaction paths needed no change at all — and then maps each survivor through kvformat::map_q8_0_cells: dequantize → re-rope → requantize. Stated honestly, the shifted K is quantize(rope(dequantize(quantize(row)))): C2's re-rope class composed with C4's packing class.
• The store allocates once per node, not once per row (kvformat::pack_q8_0_cell_into): a 2048-token prefill packs ~98k cells across 24 layers, and the per-cell vec! was the packed store's one avoidable cost.

Acceptance, as measured (CPU aarch64 on the GB10 box, 20 threads, Qwen2.5-0.5B Q4_0, n_ctx 4096, minfer bench -r 3, greedy):

• Decode is the measured effect: the fused read is 1.16× (ctx 512) and 1.31× (ctx 2048) faster than S1's dequantizing read, which is what the deferred item was for. Against f32 the packed cache went from 0.89× / 0.79× (a real regression, which is why "does not regress by more than a named factor" was the acceptance) to 1.03× (both contexts). A second run at ctx 2048 read 174.89/50.20 f32, 162.10/50.22 fused, 156.62/37.66 S1 — the same decode ratios; the f32 baseline itself moves 47.0–50.2 between runs, so only within-run ratios are quoted.
• Prefill is not measurably affected: the three configs overlap inside their own stddev (3–6 tok/s).
• Tolerance class: the real-model gate's whole-vector bound moved 3.0 → 4.0 because the fused read adds a new term — the query's own Q8_0 quantization. Measured on the same gate: 3.029 fused vs 2.505 S1 over the vector, and at the reference's argmax 0.646 vs 0.604; the decoding-relevant bound stays 1.0. MINFER_NO_FUSED_Q8_KV=1 reproduces the S1 column, so the delta is attributable to the term rather than to run-to-run noise.
• The shift: after a 4-row shift of a 20-token prefill, the first decode step — the one whose input is the shifted context — is 2.466 from the f32 shift (0.064 at the argmax) on the fused path and 3.165 (1.074) on S1's; the greedy continuation agrees 4/8 and 1/8 respectively, reported, because a flipped token makes the two runs different sequences from that step on. (An earlier version of this gate compared the last step instead and read 13.9 for a difference that is 0.06 at the decision point — the measurement, not the code, was wrong.)
• Gates, each mutation-checked, one level below the model: the_fused_q8_read_matches_the_dequantizing_reference (two KV heads whose K rows differ strongly, so a wrong head block cannot pass; using head 0's blocks for every head fails with max |Δ| 1.238 of a 1.313 spread — the shipped path measures 6.0e-5) and a_packed_physical_shift_moves_v_verbatim_and_requantizes_k (V verbatim compared as bits — a packed word is an f16 scale plus int8 quants, so as f32 it is frequently NaN and == on it is never true; K exactly the Q8_0 quantizate of the re-roped row; the tail zeroed; skipping the re-rope fails at row 0). The shift gate's first fixture filled only positions and left the builder's own cells input at zero, so all three rows went to cell 0 and the gate passed on zeros: that was found by mutating the re-rope and watching the mutant survive. The fixture now fills cells and asserts every row is non-zero.
• Suites: CPU 282 passed / 0 failed / 15 ignored (was 280/0/15), green both with and without MINFER_NO_FUSED_Q8_KV.

Handed off, still on #87: the CUDA and Metal kernels — MINFER_CACHE_TYPE=q8_0 is still refused on both, and the refusals stay loud. CUDA is the one that matters for throughput, and it is a kernel project rather than a read-path change: kv_ld4<KV> and stride_kv address elements, so a packed cell needs byte-based addressing plus a block-dequantizing load in every attention kernel (three window modes each), a Q8_0 store, and the fused decode QKV epilogue's own store. Metal stays at G5 by the round's own decision. (Superseded: CUDA landed in C4 S2b below, Metal in #310.)

C4 S2b — the CUDA Q8_0 kernels: landed 2026-09-24 (#87). The subsection below is the plan of record as it was written before the kernels; the design decisions, the measured results and the honest scope are recorded after it, so the plan and the outcome can be read against each other.

The plan of record (2026-09-24, written before the first line of kernel code).

The CPU half of #87 landed (S2a). The device half is a kernel project, not a read-path change, so it lands as its own increment against the same issue. This is the map it starts from: where the CUDA backend touches a KV region, the design that covers those sites, and the dispatch cuts that keep the first increment bounded. It came from walking the three files below at a5f7961 + S2a, so the line numbers are the ones to read, not to trust.

What is there today. No CUDA code knows KvFormat. The layout is one process-wide bool (cuda.rs's KV_F16), and every kernel addresses a cell as row * nkt elements of a float*/__half* — there is no byte addressing anywhere, so a 34-byte-per-32-element cell cannot be expressed by any current kernel signature. set_kv_cache_type maps anything that is not exactly f16 to false, so flipping the load gate before the kernels exist would silently mean f32.

The design. A layout tag plus byte-addressed row accessors in cuda_kernels.cu:

• KV_LAYOUT_F32 / KV_LAYOUT_F16 / KV_LAYOUT_Q8_0, kv_row(base, cell, row_bytes), and kv4<LAYOUT>(row, elem) -> float4 as the one load idiom: f32 is the old float4 load, f16 the old two-__half2 pair (both bit-identical to today's instantiations), and Q8_0 reads block elem/32's f16 scale plus four quants at 2 + elem%32. A 4-element group never straddles a block, because a KV head's base is hd-aligned and hd % 32 == 0 — the property ensure_kv's packed-width check already enforces;
• kernels take const void* k/v + size_t row_bytes instead of a typed pointer plus stride_kv = nk * hd, templated on int LAYOUT rather than typename KV;
• store_kv_q8_0: one thread per (row, 32-element block), with the CPU store's own quantizer (amax/127, f16 scale, round_ties_even), so both backends store the same bytes.

The access sites to convert (src/cuda_kernels.cu, plus the host rows below):

The dispatch cuts that keep the first increment bounded — build-time choices (no silent fallback), each stated when the format is enabled:

• decode (nt == 1) takes the split-K path through the converted 1-warp body; its hybrid 4-warp dispatch (hd == 128, nkv >= 1921) is f16-typed and is skipped for Q8_0 (rpw_gate = 0);
• 1 < nt ≤ 16 routes to gqa_attn_f32 instead of the batched split kernel, which exists for spec-verify's bitwise identity contract with sequential decode. A speculative session must therefore refuse a packed cache loudly (a draft keeps its own KV, and the two reduction schedules would no longer agree) — consistent with S2a, which already refuses a draft in the session container;
• prefill (nt > 16) routes to gqa_attn_f32: the FA path (fa_prefill_f16kv, f16-typed shared-memory staging) is not offered for Q8_0 in this increment, so the prefill is correct but off its tuned path;
• the fused decode epilogue (Op::FusedQKV / Op::QkvBiasRopeStore) is not built for a packed cache — the model builders' layer_gpu gate gains && !packed, and a Q8_0 decode takes the unfused bias/rope/store chain through the converted store kernel.

Acceptance (the issue's, made concrete): cuda_map_window_matches_the_span_over_the_same_rows (cuda_backend.rs:3678) extended to Q8_0 rows — it writes its own region contents and calls exec_ids directly, so it needs a packed cell encoder and a packed-word region size; a Q8_0 store→attention round trip (cuda_kv_f16_roundtrip_attn, :5150, is the pattern); copy_cells under the packed stride (cuda_f16_kv_cell_move_strides_by_row_bytes, :3946); the real-model gate a_packed_kv_cache_answers_like_the_f32_one run with the cuda feature on GB10 (since #123 that gate is CPU-forced on a CUDA build until these kernels land, so this acceptance includes re-pointing it at the device; CI has no GPU, and its MINFER_C4_MODEL arm covers a second model); and an A/B against f16 on the same model/context with the numbers recorded. The honest expectation, stated before the work: the win on the device is memory (3.76× less than f32, 1.88× less than f16); whether decode bandwidth also improves depends on the block-dequant instruction count, so the bar against f16 is "no worse than a named factor", not "faster".

S2b design decisions, frozen before the first line of kernel code (2026-09-24). The plan above is the map; these are the choices it leaves open, decided up front so the implementation cannot drift into a silent fallback:

1. Layout tag and accessors. KV_LAYOUT_F32 = 0 / KV_LAYOUT_F16 = 1 / KV_LAYOUT_Q8_0 = 2 are #defines in src/cuda_kernels.cu, and 0/1/2 are a host contract: the Rust side stores the same codes (the registry's KvFormat discriminants) and every launcher takes the tag as an int. Two accessors are the only place a KV address is formed: kv_row(const void* base, int64_t cell, size_t row_bytes) (byte address of a cell) and kv4<LAYOUT>(const char* row, int elem) -> float4 (the single load idiom). F32 is the old float4 load and F16 the old two-__half2 pair, both bit-identical to today's kv_ld4<float> / kv_ld4<__half>; Q8_0 reads block elem/32's f16 scale plus four quants at 2 + elem%32. A 4-element group never straddles a block because a KV head's base is hd-aligned and hd % 32 == 0 — the property ensure_kv's packed-width check already enforces.
2. No typed pointer survives. Every attention kernel that can see a packed row takes const void* k, const void* v plus size_t row_bytes, and templates on int LAYOUT instead of typename KV. The row_bytes == nk * hd * 4 value the f32 path passes is the same arithmetic as the old stride_kv * sizeof(KV), so the f32/f16 instruction streams are unchanged.
3. Dispatch cuts, each stated when the format is enabled (build-time; a packed row never reaches a kernel that would address it as f32):
   • decode nt == 1 → the converted split-K 1-warp body (gqa_attn_split_partial), with rpw_gate = 0: the hybrid 4-warp dispatch is f16-typed (attn_split_h4w_body takes const __half*) and is not converted in this increment;
   • 1 < nt <= 16 → gqa_attn_f32. The batched split kernel (gqa_attn_split_partial_bt) exists for spec-verify's bitwise identity contract with sequential decode; rather than claim that contract without measuring it, a speculative session refuses a packed cache loudly (the draft already keeps its own KV, and C5 already refuses a draft session);
   • prefill nt > 16 → gqa_attn_f32. The FA path (fa_prefill_f16kv, f16-typed shared-memory staging and tensor-core QK^T) is not offered for Q8_0, so the prefill is correct but off its tuned path — measured and reported, not hidden;
   • the fused decode epilogue (Op::FusedQKV / Op::QkvBiasRopeStore) is not built for a packed cache: the model builders' layer_gpu gate gains && !packed, so a Q8_0 decode runs the unfused bias → rope → store chain through the converted store_kv_q8_0.
4. The store is the CPU's quantizer, byte for byte. store_kv_q8_0 maps one thread to one (row, 32-element block), computes amax, d = amax/127 as an f16 (nearest-even), and round_ties_even for each quant — the same three steps quants::quantize_row_q8_0_into uses — and writes d then the 32 quants into the packed cell at 34-byte stride. Both backends therefore store the same bytes for the same f32 row, which is what makes the CPU/device Q8_0 comparison a layout check rather than a tolerance question.
5. The gate flips only after the kernels exist, and through the registry. READS_PACKED_KV stays false until every dispatch cut above is in place; then it becomes true for CUDA and KvFormat::supports(Cuda) follows automatically (registry::reads_packed_kv is the one authority). CudaBackend gains an int layout field fed by the resolved KvFormat, so MINFER_CACHE_TYPE=q8_0 can never again mean "f32 with a packed region".
6. Acceptance bar, named before measuring. The device win is memory: a Q8_0 cell is 34/128 bytes per element versus 4 (f32) and 2 (f16), i.e. 3.76x less than f32 and 1.88x less than f16. Whether speed also improves depends on the block-dequant instruction count in the load path, so the bar against f16 is stated as a factor, not as "faster": Q8_0 decode tokens/s must be no worse than 1/1.30 of the f16 decode rate on the same model and context (the same 1.30x already recorded for the CPU's fused Q8_0 read at ctx 2048), and the prefill is reported as its own number because it takes the untuned gqa_attn_f32 route.

Why it was not in the S2a increment. ~10 kernel sites, 3 launchers, ~15 host sites and 41 test sites, each needing an nvcc iteration and — for the gates — a serial device run. It was recorded here rather than half-wired. Metal's half stays at G5 — since landed as #310.

What actually landed (2026-09-24). The real counts are close to the estimate: 8 kernel sites, 3 launchers (launch_gqa_attn_split_q8_0, the layout-tagged launch_gqa_attn_f32, launch_store_kv_q8_0) plus two re-templated existing ones, ~25 host sites (cuda.rs's FFI + CudaState methods, cuda_backend.rs's store/attention/fused-epilogue dispatch, copy_cells, the registry entry), 4 new gates and 2 strengthened ones. The two files the plan expected not to change did not: kv_move_rows is a plain word copy (verified — the packed cell's word count is what the host already passes, and the Q8_0 move gate pins it), and the CPU path is untouched.

• Layout and accessors, exactly as designed. KV_LAYOUT_F32/F16/Q8_0 are 0/1/2 on both sides of the FFI; kv_row + kv4<LAYOUT> are the one address/load idiom. The f32 and f16 kernels compile to their pre-C4 instructions (same loads, same cells, same order).
• The gate flipped through the registry. BackendCaps::reads_packed_kv became true for CUDA and KvFormat::supports(Cuda) followed automatically; CudaBackend carries an int kv_layout fed by the resolved KvFormat, so the pre-S2b bool that mapped anything but f16 to f32 can no longer turn a packed region into f32 rows.
• The store is the CPU's quantizer, byte for byte. cuda_q8_0_store_matches_the_cpu_quantizer asserts the device's packed words equal kvformat::pack_q8_0_cell's for two block widths, and that unwritten cells stay zero. Mutation: amax / 126 instead of / 127 fails it at nkt=64 row 0.
• All three acceptance gates, mutation-checked. cuda_map_window_matches_the_span_over_the_same_rows now sweeps f32/f16/q8_0; cuda_kv_q8_0_roundtrip_attn (store → attention) and cuda_q8_0_kv_cell_move_strides_by_row_bytes (copy_cells under the packed stride) are the two new device gates. compute-sanitizer --tool memcheck over all three reports no memory error (the 2 pre-existing CUDA API errors it did report were root-caused in S2c below, #145). Mutations: dropping the Q8_0 block base in kv4 left the map-window gate green (every comparison there is between modes over the same bytes, so a value-level fault shifts both sides together — the F6 lesson) and made the round-trip and the real-model device arm fail (max |Δlogit| 2.48 → 26.34 of a 37.8 spread); halving the packed stride in copy_cells fails the move gate; flipping READS_PACKED_KV back to false makes the device arm refuse the region loudly. The map-window gate was strengthened: it now also compares one single-row Q8_0 window against the dequantized V cell (max |Δ| 1.71 under the block-base mutation), so it cannot pass on consistency alone.

The A/B against f16, measured on dgxspark (aarch64, GB10 sm_121) — the bar was named before the work as "Q8_0 decode no worse than 1/1.30 of f16" and it is not met against the fused f16 baseline:

Superseded 2026-10-07 — read the numbers as measured, not as current. This table is the pre-#144 S2b measurement (2026-09-24; its baseline was re-measured in the #144 round at master 85c712e, which reproduced it within a few percent), and three later changes moved it: #144 items 1+3 (041de15, 2026-09-26), #186 (cc19b4f, 2026-09-27, the dp4a packed K dot) and #202 (798fd32, 2026-09-27, the wide u16 block load). The current A/B is the 2026-10-07 comment on #310, measured at 740e0ff on dgxspark (aarch64, GB10 sm_121): 1.23x decode / 1.24x prefill at hd 64, and 1.01x / 1.04x at hd 128 — the 15.3x is gone, and q8_0 is still never faster than f16 on CUDA.

The f16 column is pinned, not auto. auto_device_format needs n_layers × n_kv_embd ≥ 8192 (AUTO_F16_MIN_KV_ELEMS, src/graph/kvformat.rs), so today's resolver answers f32 for the 0.5B (24 × 128 = 3072) and f16 for Qwen3-0.6B (28 × 1024 = 28672); both arms in this table were run with MINFER_CACHE_TYPE=f16 explicitly, and the header's older "(default policy)" wording does not describe the 0.5B arm on the current resolver.

Read honestly [as measured pre-#144; the clauses that no longer hold are marked inline]: decode on the 0.5B misses the named 1.30x because ~1.18x of the gap is the stated cut (no fused QKV epilogue for a packed cache — the unfused chain adds launches on a model whose decode is launch-bound) and the remaining 1.25x is the packed load itself (kv4<Q8_0> costs four int8 converts and four multiplies where f16 costs two __half2 converts; there is no dp4a packed dot in this increment — #186 landed it on 2026-09-27, cc19b4f]). On Qwen3-0.6B, where the fused epilogue is not on the f16 default path in the same way, decode is 1.13x — inside the bar (the superseding run at 740e0ff measures 1.01x). The prefill is a different story: at hd 128 the f16 path is fa_prefill_f16kv (8604 tok/s) and the packed path is the general kernel (563 tok/s), i.e. the packed cache is correct but off its tuned route by 15x — #144 item 3 then landed the packed FA prefill, so the superseding run measures 1.04x here]. The win is memory, as stated up front: measured on both models, the f32/f16 region is 6 291 456 B and the Q8_0 region 1 671 168 B — 3.76x smaller than f32 and, against f16's actual 2 B/element payload (3 145 728 B), 1.88x smaller.

Acceptance results.

• Real model, device arm: a_packed_kv_cache_answers_like_the_f32_one runs both arms — the CPU one (kept, coverage on every build) and, on a CUDA build with a device, one that asserts device() == Cuda. 0.5B: f32 6 291 456 B vs q8_0 1 671 168 B (3.76x), CPU max |Δlogit| 3.03 / argmax 0.65, CUDA 2.48 / 0.60 of a 37.8 spread; the physical-shift arm: CPU 2.47 / 0.064, CUDA 3.34 / 0.42 — inside the bounds the CPU gate had already fixed.
• Suites: CPU cargo test --release 432 / 0 / 28 unit + 10 / 0 / 6 integration (unchanged); CPU serial ignored 28 / 0; CUDA serial unit 490 / 0 / 31 (was 486/0/31: three new device gates + one new cuda.rs layout test; the packed gate now also runs on the device); CUDA serial ignored 0.5B 31 / 0; Qwen3-0.6B 30 / 1, the one failure the then-pre-existing #130 f16 session-container gap (closed in C5 S3: 31 / 0).
• The Q8_0 session container works. a_session_resumed_from_disk_continues_bitwise with MINFER_CACHE_TYPE=q8_0 on the device prints live KV format: q8_0, saves 24 layers / 256 cells / 5 written / 1 696 464 B and continues bitwise (max |Δlogit| = 0). The container's FLAG_PACKED bit is what encodes it; f16 had no flag then — the gap #130 closed in C5 S3 (FLAG_F16), below.

Handed off after S2b: the packed fused decode epilogue, a dp4a packed K dot and the FA prefill on packed cells — taken up as #144 and recorded in the next subsection (items 1 and 3 landed; item 2 then landed as #186 below, DONE 2026-09-27). Metal's half stays at G5 — since landed as #310.

C4 — #144: the packed fused decode epilogue and the packed FA prefill · #144 — DONE (items 1 + 3) 2026-09-26

Why. C4 S2b's packed cache is correct and 3.76x smaller than f32, but three paths were off their tuned route: a Q8_0 decode ran the unfused bias/rope/store chain (the builders' layer_gpu gate carried && !packed), kv4<Q8_0> paid four int8 converts where f16 paid two __half2, and an hd-128 packed prefill could not enter fa_prefill_f16kv's f16 shared-memory staging, so it fell to the general layout-tagged kernel — 15.3x off on Qwen3-0.6B pp2048. This ticket took the first and third (a dp4a K dot is a numerics change with its own accuracy statement; it is filed separately).

What landed.

1. The packed fused decode epilogue. attn_bias_rope_store_q8_0 keeps the f32/f16 epilogue's q section (bias + rope in place, one thread per pair) and re-maps K and V to one thread per (head, 32-element block): the thread computes all 32 K roped values itself (from the unroped row plus the pair partner d <-> d + hd/2) or the 32 V bias-added values, and hands them to q8_0_quantize_block — the same device quantizer store_kv_q8_0 now calls, factored out so the two cannot drift. Neither K nor V is written back (both buffers are dead in the fused classes) — which is also why the packed epilogue has no rope race. The builders dropped && !b.kv_is_packed() from the Qwen2 family's fuse_qkv gate. FusedQkvMeta / QkvBiasRopeStoreMeta / FusedQkvNormMeta gained a row_elems field (the packed cell width, KvFormat::row_elems(nkt)) and the allocator sizes the region from it: without that, the first packed fused graph died with the node declares 34 words per cell but the q8_0 layout packs one cell of 512 elements into 136.
2. The packed FA prefill. fa_prefill_f16kv became fa_prefill_kv<CAUSAL, MAP, LAYOUT>: the staging is the only layout-dependent part (kv8_q8_0 dequantizes each packed 32-element block into the same f16 tile, one 16-byte smem store per 8 elements); the tensor-core QK^T, the fragment-resident softmax and P·V are untouched. gqa_attn_kv_prefill serves f16 and Q8_0 from one call and keeps the general layout-tagged kernel as the documented fallback; the launcher's failed-smem arm and MINFER_NO_FA_PREFILL=1 both reach it. The chokepoint bumps testfail::note_checked("cuda_fa_prefill_q8_0"), so the gate can prove the packed prefill took the FA route rather than silently falling back.

The A/B protocol and the numbers. All numbers: GB10 sm_121, cargo build --release --features cuda, minfer bench -p 2048 -n 128 --n-ctx 4096 -o json, MINFER_CACHE_TYPE pinned per arm, 5 interleaved rounds, medians, 2026-09-26. The bars were named before measuring (gate contract rules 3 and 5); the full protocol lives in CUDA_OPTIMIZATION.md and cuda_optimization_steps/107.

The re-measured baseline (master 85c712e, its own binary) reproduced the S2b table within a few percent: 0.5B q4_0 tg128 237.37 f16 / 161.49 q8_0, MINFER_NO_FUSE_QKV=1 f16 200.80 (a 1.182x cut — the ticket's 1.18x); Qwen3-0.6B Q8_0 pp2048 8323.24 f16 / 564.47 q8_0.

Item 1's miss is a result, not a failure. The ticket's "~1.18x" was the fusion cut measured on the f16-weight arm; on the q4_0 packed arm the same binary's fused-vs-unfused A/B is 1.052x. The saving is the ~6 launches/layer the epilogue removes (~13 µs/layer at 24 layers ≈ the 0.32 ms/token the two arms differ by), so the launch-overhead component is what this cut actually buys; the residual 1.39x f16-to-packed gap is the packed load (kv4<Q8_0>'s four converts) and the 1-warp split-K decode body's rpw_gate = 0, which the dp4a item targets.

Verification. The two new device gates and the five named Q8_0 gates are green: cuda_q8_0_fused_epilogue_matches_the_cpu_quantizer (new: byte-exact K/V against the CPU quantizer at pos = 0, a value q check at pos = 7, plus V byte-exact), cuda_q8_0_fa_prefill_attention_parity (new: 100-token hd-128 prefill against the CPU attention over the same packed bytes, 5e-3 class, max err 4.8e-4, plus the note_checked observation arm), cuda_q8_0_store_matches_the_cpu_quantizer, cuda_kv_q8_0_roundtrip_attn, cuda_q8_0_kv_cell_move_strides_by_row_bytes, the Q8_0 arm of cuda_map_window_matches_the_span_over_the_same_rows, and cuda_fa_prefill_attention_parity (the f16 route, unchanged, max err 2.8e-4). Mutations: reverting the fused kernel's block offset (d = blk*32 + i → d = i) turns the byte arm red (it found exactly that bug during development); shifting kv8_q8_0's block base (elem >> 5 → elem >> 4) turns the FA parity red (max err 3.33); making the packed prefill skip the FA launch leaves the parity arm green (max err 5.8e-5) but turns the observation arm red — which is the whole reason it exists. MINFER_TEST_ISSUE162=1 drives all 124 audited launch sites (six new: three fa_prefill_kv layouts × modes and the packed epilogue) and is green.

Deliberately not taken. The dp4a packed K dot (item 2): a per-head-quantized query against the packed K accumulates in int, which is a numerics change needing its own accuracy statement and a re-measured real-model tolerance — filed as a follow-up rather than assumed. Qwen3's packed fused QKV (Op::FusedQkvNorm) also waits: that op has no CUDA kernel at all (it is Metal + CPU today), so the !packed gate there is not what keeps Qwen3's decode off the device; its 1.13x packed gap is unrelated to this ticket's cuts.

C4 — #186: the dp4a packed Q8_0 K dot · #186 — DONE 2026-09-27

Why. [#144]'s residual was 1.393x on Qwen2.5-0.5B q4_0 tg128 (170.11 q8_0 vs 236.97 f16), and [#144] names two components: the 1-warp split-K body's rpw_gate = 0 plus the kv4<KV_LAYOUT_Q8_0> load. Only the load is addressable by a dp4a K dot. The trap is that a whole-kernel q8_0-vs-f16 gap would mix the two: at hd 128 the f16 arm takes the 4-warp hybrid body while the packed arm cannot. hd 64 (the 0.5B) is the shape where both take the identical gqa_attn_split_partial<LAYOUT,true,false> on the same grid, so the delta there is the load alone.

The bar, named before measuring. The load-attributable share of the Q8_0 decode step ≥ 10%. Reasoning: only one of the residual's two components is addressable; a free load recovers at most share / 1.393; dp4a removes only the K side's converts and multiplies, adds a query quantize, and still pays the memory traffic; [#144]'s landed item 1 was 5.3% on this arm, the measured floor for "worth landing". Measured 20.3% (nsys, real decode, node tracing: the incumbent packed partial kernel is 67424 ns median per layer/token vs f16's 18528, = 1.182 ms of a 5.818 ms/token step). Bar cleared.

What landed. attn_split_1w_body<LAYOUT,CAUSAL,MAP,Q8DP4A> (the Q8DP4A default false, so the f32/f16 and verify instantiations are untouched) quantizes each lane's four query values against its 32-element block's amax (the eight lanes sharing a block reduce it with three shfl_xor), reads the K quants as int8 through kv4_q8_0_packed, accumulates __dp4a(qk, ki, 0) and scales by d_q * d_k once per block. V still dequantizes; only K changes. The launcher launch_gqa_attn_split_q8_0 takes an int dp4a and picks the instantiation; cuda::q8_kv_dp4a_enabled() resolves MINFER_NO_DP4A_Q8_KV=1 once per process (the same-binary A/B control), and the Rust call site bumps testfail::note_checked("cuda_q8_kv_dp4a") only when it launched the dp4a arm.

The numbers (GB10 sm_121, cargo build --release --features cuda, minfer bench -p 2048 -n 128 --n-ctx 4096 -o json, 5 interleaved matched rounds, medians, 2026-09-27):

The 0.5B f16 arm in the same round is 240.03, so the packed/f16 residual went 1.396x → 1.242x; on Qwen3-0.6B the f16 arm is 136.73, i.e. the packed decode now matches f16. The isolated probe (same 1-warp geometry, load only) cuts the delta from +96% to +40% at hd 64, and ncu (collected as root, the module parameter unchanged) names the residual: dp4a removes 494984 instructions (−15.3%) and leaves the L1 load-sector count unchanged at 792904 — 1.71x f16's, while its L2 read sectors are 0.57x f16's — so the remaining 1.242x is the 34-byte block layout's L1 request count, not DRAM traffic and no longer the arithmetic.

Tolerance, re-measured not assumed. a_packed_kv_cache_answers_like_the_f32_one's CUDA arm reads max |Δlogit| 2.479504 of a 37.82 spread, at the argmax 0.552662, greedy 9/9 (incumbent arm: 2.479504 / 0.596050 / 9-of-9) — inside the ≤4.0 / ≤1.0 class. The max is unchanged because it is the nt = 512 prefill step's; the decode steps' deltas moved and nsys on the gate itself names gqa_attn_split_partial<(int)2,(bool)1,(bool)0,(bool)1> under the default and ...,(bool)0> under the control, so the class is not "both arms ran the old path".

Verification and refusal-to-over-claim. The dp4a arm was added to cuda_kv_q8_0_roundtrip_attn (two cells through an explicit span, an exactly Q8_0-representable query, plus the observation-counter arm); the first version used one cell and was vacuous — a mutated block base still passed, which is how the two-cell form was found. Mutation: elem >> 5 → elem >> 4 → red at max |Δ| = 0.35126442. Prefill/verify is deliberately not taken: the nt > 1 general kernel also reads kv4<Q8_0> and measures 1.32x off its f16 arm on the 0.5B prefill attention, but that is a different kernel (per token and head query, no shared block scale) and not the ticket's residual — named rather than assumed. Full record: cuda_optimization_steps/108.

C4 — #202: the packed Q8_0 KV cell's L1 request count · #202 — DONE 2026-09-27 (counter-only; throughput bar missed)

Why. [#186]'s ncu attribution named the residual: the packed decode arm issued 1.71x f16's L1 load sectors (~792904 vs 463008) with 0.57x its L2 sectors, because a 34-byte Q8_0 block's 4-element group at blk + 2 + (elem & 31) is 4-byte aligned only on odd blocks, so kv4<Q8_0> spends four s8 loads plus a scattered scale load per group. The ticket listed three candidates and said to prefer the one with no layout change if it clears the bar (a split plane or a 36-byte block would move the CPU store/read path, copy_cells, map_q8_0_cells, the FA-prefill staging and the C5 session version).

What landed. The no-layout candidate: 34k + 2 + 4m is always 2-byte aligned (34k is even for every k, and 4m is), so two unsigned short loads replace the four s8 loads for the same bytes. q8_0_load4<WIDE> / kv4<LAYOUT,WIDE> / kv4_q8_0_packed<WIDE> / a Q8WIDE template parameter on attn_split_1w_body, and a third arm in launch_gqa_attn_split_q8_0, selected once per process by cuda::q8_kv_wide_enabled() (MINFER_NO_Q8_KV_WIDE=1 is the same-binary control). Three new launch sites, driven by the #162 test and added to the audit fixture. No layout, CPU, copy-stride, session-format or map_q8_0_cells change.

The bar, named before measuring. (a) packed/f16 L1 load-sector ratio ≤ 1.30x; (b) same-binary tg128 ≥ +2% over the #186 dp4a baseline 193.30, pp2048 within 2%. Measured: (a) met with margin — 792904 → 455840, ratio 1.712x → 0.9845x, instructions −2.17%; (b) NOT met — 193.13 → 193.65 (+0.27%, 5 interleaved medians), the kernel itself only 1.4–1.7% faster (nsys 40320 → 39744 ns median).

Honest result: a partial refutation. The L1 request count moves exactly as the mechanism predicted and the decode step does not follow, so the 1.23x packed/f16 residual is not L1-request-bound — not L2 traffic (0.564x), not instruction count (1.10x), not load sectors (0.9845x). Landed counter-only, byte-identical (the nt > 1 general kernel keeps the byte form; only the dp4a decode arm opts in), with the latency hypothesis (per-block scale load + cvt + V dequant on a 1-warp block's critical path) recorded for the next attribution rather than claimed. Mutation: swapping the two u16 halves in q8_0_load4_wide turns cuda_kv_q8_0_roundtrip_attn red at max |Δ| = 2.5326836. Full record: cuda_optimization_steps/109.

Test infrastructure — #207: the CUDA count-row drift · #207 — DONE 2026-09-27

Why. The CUDA rows are live_check = false recorded measurements (CI has no GPU), but they count the same test binary plus the device-gated tests: a CPU-only ticket that adds a feature-independent test moves both. [#140]'s K-quant encoder tests (7 passed / 1 ignored) and [#142]'s bf16 writer + CPU-path tests (7 passed / 2 ignored) did exactly that and left the row reading 548 / 0 / 39 while a device run on the branch reads 562 / 0 / 42 and both CUDA real-model rows read 42 / 0; this was the second consecutive occurrence.

What landed. Mechanism (c), chosen over a pending field and a rule line: docs/status.toml's recorded rows may carry projection_key / projection_box / projection_base_passed, and scripts/check_status.py --check prints (never fails on) passed + (cpu_now − base) when the CPU twin has moved. The relation is inexact — the device-gated tests move independently — which is exactly why it is a hint and not a comparison; the point is that a stale recorded row cannot look current without a projection beside it. The unit/sanitizer rows project from cpu-unit (aarch64) and the two real-model rows from cpu-real-model (aarch64), each with that row's passed at this measurement as the base, so all four hints are silent today and fire the moment a CPU ticket adds tests. --selftest gained a moved/not-moved pair (12/12 cases).

Why. minfer bench on a CUDA build printed CUDA kernel launch error: 1 between its two loops — pre-existing on master, not an S2b regression — and compute-sanitizer --tool memcheck over the unit suite reported 36 CUDA API errors. CudaState::sync polls cudaGetLastError, which reports whatever an earlier call on the thread latched, so an API error from several operations ago was printed as a failure of the kernel that had just run.

Root cause: two origins, both unchecked return values.

1. gemm_prefill_smem_init (1 of the 36; #145). The eager dynamic-smem opt-in looped over every gemm_f16_nt_kernel_t<tm, ks, af32> and asked cudaFuncSetAttribute(.., cudaFuncAttributeMaxDynamicSharedMemorySize, N) for a stale N: its formula assumed a 512-thread TM=256 launch and always added the AF32 mirror, so gemm_f16_nt_kernel_t<256,64,true> requested 131072 B against GB10/sm_121's cudaDevAttrMaxSharedMemoryPerBlockOptin of 101376 B. The rejected call's return value was never read, so the error latched. The same formula existed a second time, inline in launch_gemm_f16, and that copy dropped the AF32 mirror — the launcher declared 16384 B less than the kernel's own Bs/Cs offsets need at the default KS=32, an out-of-declaration shared-memory access that only worked because the block's smem happened to be carved where nothing else wrote.
2. cudaGraphDestroy on a cudaGraphExec_t (26 of the 36; #128). CudaState::graph_destroy is the only destroy path for an exec handle and called the graph destructor. It returned cudaErrorInvalidValue, leaked the exec, and the latch surfaced at the next sync — the actual source of the bench line. The remaining 9 errors were cudaGetLastError observations of one of the two.

What landed.

• One formula, read twice. gemm_dynamic_smem_bytes(tm, ks, af32) is the single source for the kernel's byte layout (As 2*TN*KS halves + Am 2*TN*KS floats [AF32] + Bs 2*TM*KS halves + Cs NW*256 floats, TN = 64, NW = blockDim.x/32 = 8 at every launch site); launch_gemm_f16 and gemm_prefill_smem_init both call it.
• The eager opt-in is checked and deliberate. Every cudaFuncSetAttribute return value is read; a failure is named once, at init, with the function, attribute, requested bytes, device limit and cudaGetErrorName, and cleared there. A request above the queried device limit is skipped without calling it, with the reason printed — on GB10 that is gemm_f16_nt_kernel_t<256,64,true> (122880 B > 101376 B), which cannot launch on this device at all. CudaState::try_new reports the counts.
• sync attributes honestly. latched_api_error_message names the observer (cudaGetLastError), the cudaGetErrorName symbol and the code, and states it is not attributed to a kernel; the error is counted (latched_api_error_count) and cleared — still visible, never dropped. debug_sync's label follows.
• graph_destroy calls cudaGraphExecDestroy; the cudaGraph_t / cudaGraphExec_t distinction is stated at the extern and the call site, and a failure is named there and cleared.

#218 forward note (2026-09-29): the eager sweep this record describes is gone. #188 deleted gemm_prefill_smem_init's CudaState::try_new call site with no mention in any commit message or document, and #218 removed the orphaned function (and the checked/skipped introspection) instead of leaving it allow(dead_code)-annotated. The per-launch opt-in this record already describes — gemm_smem_optin → the shared minfer_smem_optin, reached by launch_gemm_f16 — is now the whole mechanism; the invariant it must uphold (the attribute is in force before a capture window opens, never set inside one) is upheld by the 3-run capture warmup (capture_warmup), cudaStreamCaptureModeThreadLocal and the per-instantiation cache, and gated by cuda_prefill_smem_optin_is_done_by_production, the control arm cuda_prefill_smem_optin_refusal_fails_the_prefill, the coverage arm cuda_prefill_smem_lazy_optin_admits_every_launchable_instantiation, and cuda_prefill_smem_optin_is_never_set_inside_a_capture_window — see the #218 record in §test-infrastructure. gemm_dynamic_smem_bytes remains the one formula; the "read twice" wording above is historical (the launcher reads it once).

#223 forward note (2026-09-29): the eager half is back, but not as the sweep this record describes. CudaState::try_new now calls gemm_prefill_smem_prewarm_one(tm, ks, af32) for every launchable combination — a thin dispatcher onto the same production gemm_smem_optin cache the launcher reads — so there is one opt-in mechanism and one cudaFuncSetAttribute site, not a second sweep with its own copy of the formula (the pre-#145 bug class). The placement is the argument: at try_new no CudaBackend (and so no capture window) can exist, so the attribute is set outside any window by construction. The checked/skipped counters and the startup banner are still gone; a successful pre-warm is silent, and each failure or deliberate over-limit skip is named per instantiation. The lazy path stays as defence in depth, and the #218 gates keep their claims by running their children against the documented control MINFER_NO_GEMM_PREWARM=1. On the real binary the loop costs ~2.2 ms (the fatbin's one-time module load), which prewarm_prefill() already paid — net new startup cost ≈ 0, hot path within ±1% — see the #223 record in §test-infrastructure.

Acceptance results (GB10, sm_121, CUDA 13.0, driver 580.178.04).

The five new gates, all mutation-checked.

• the_latched_error_message_never_blames_a_kernel (pure) — the message names cudaGetLastError and cudaErrorInvalidValue and does not contain "kernel launch". Mutation: the old message fails it.
• the_gemm_smem_formula_matches_the_kernel_layout (pure) — pins gemm_dynamic_smem_bytes against the kernel's byte layout for all 12 (tm, ks, af32) combinations. This is the value-level arm the F6 lesson asks for: a mode-vs-mode comparison is blind to a shrunk formula (launcher and init shrink together, so a consistency check stays green). Mutation: dropping the AF32 term fails it.
• cuda_prefill_smem_optin_covers_every_launchable_instantiation (device) — every >48 KiB request the device admits reads back opted in through cudaFuncGetAttributes().maxDynamicSharedSizeBytes; every over-limit one is skipped, never called. Mutation: continue-ing one admitted combination fails it.
• cuda_graph_exec_destroy_leaves_no_latched_error (device) — capture → instantiate → destroy returns true and leaves the latch at 0. Mutation: cudaGraphDestroy fails it. The gate asserts the destroy call's own result (graph_destroy now returns bool), not just the latch: the failure is named and cleared inside graph_destroy, so a latch-only assertion passed with the wrong destructor — the first cut of this gate passed for the wrong reason and was fixed here.
• cuda_sync_surfaces_a_latched_error_as_latched (device; env-gated behind MINFER_TEST_LATCH_ERROR=1 because it deliberately latches a real API error, which a compute-sanitizer run must not see) — sync reports the injected error once and clears it. Mutation: dropping the report fails it.

The >48 KB capture path stays covered by the existing prefill gates — prefill capture is ON by default (R3-B) and cuda_prefill_capture_defaults_on, cuda_prefill_capture_bit_parity_pp16_pp300, cuda_multisplit_capture_bit_parity and the real-prefill cuda_graph_generation_replay_parity_real_model are all green in the 495 / 0 / 31 run.

The deliberate-failure check. The pre-fix over-limit request was forced back with the sanitizer-clean skip bypassed, so cudaFuncSetAttribute really failed: the init printed cudaFuncSetAttribute(gemm_f16_nt_kernel_t<256,64,true>, cudaFuncAttributeMaxDynamicSharedMemorySize, 131072 B) failed: cudaErrorInvalidValue (1); device opt-in limit 101376 B, and cuda_prefill_smem_optin_covers_every_launchable_instantiation failed (failures = 1). Reverted; the reverted src/cuda_kernels.cu is byte-identical to the committed one (sha256sum, in the closing comment on #145).

Honest scope. The device limit — and therefore which instantiation is skipped — is measured on GB10/sm_121 only; the decision is a runtime query and the skip is printed, so another device skips a different subset. MINFER_GEMM_TM=256 + MINFER_GEMM_K64=1 + the f32-A path is the only combination that lands on the skipped instantiation, and it had no working opt-in before either. The remaining unchecked CUDA calls in the same file are enumerated in #147 rather than silently fixed here.

C4 — S2d: the remaining unchecked CUDA attribute / launch / destroy returns · #147 — DONE 2026-09-25

Why. S2c (#145) fixed the two latched-error origins compute-sanitizer could still see and left the CUDA unit suite at 0 API errors — but the audit it performed listed a whole family of call sites that still discard a return value which gates a later launch or allocation, the class that let #122, #128 and #145 hide in the first place. None was known to fail on sm_121, so this is latent hardening, not a live bug: the value is that the next driver, device or tile-config change cannot turn one of them into another phantom "kernel launch error".

The re-derived site list (verified against the tree, not the ticket). The S2c table named six rows; re-deriving them against the post-#145 source gives eight code sites, one of which the table understated (the wide-NT launcher already read its own return but reported nothing) and two of which were already fixed by #145 (the eager gemm_prefill_smem_init opt-in and graph_destroy's cudaGraphExecDestroy). What remained, and how each is closed:

The mechanics (one place, no per-launcher copy). cuda_kernels.cu gained a small block of shared helpers — minfer_smem_optin (skip an over-limit request without calling, otherwise call once, name the function/attribute/bytes/device-limit/cudaGetErrorName, clear the latch, return false → refuse), minfer_launch_prelude (report a latch that predates the launch, so the post-launch read is a launch check and not attribution by position), minfer_launch_smem, and minfer_launch_ok (name the launch's own error and clear it) — plus minfer_site_fail_* introspection for the gates. minfer_smem_optin queries cudaDevAttrMaxSharedMemoryPerBlockOptin once and skips a request above it, which is what keeps compute-sanitizer clean (calling it would only observe cudaErrorInvalidValue).

Two signature changes, both followed through. launch_mmq_raw_nt and launch_mmq_nt went void → int, and launch_gemm_f16 / launch_gemm_f32a went void → int (1 = launched and accepted, 0 = refused). The Rust callers read the result and return Err — a failed launch is an error, never a silent pass over an unwritten output. The fallback launchers keep their documented 0 = clean fallback contract, so which kernel runs is unchanged. Nothing numeric or dispatch-related moved.

A measured correction to a documented belief. docs/CUDA-BACKEND-DESIGN.md said cudaFuncSetAttribute is illegal under cudaStreamCaptureModeGlobal and that a lazy opt-in inside a window fails. A probe on CUDA 13.0 / driver 580.178.04 / sm_121 (/tmp/fix147_attr_capture_probe.cu) shows the call now returns cudaSuccess inside an open Global capture window (both for the set value and a new one), so the eager init is defence-in-depth rather than the only working form; the launcher caches one answer per instantiation, so it is not re-asked on the hot path either way. Recorded in the design doc.

Acceptance results (GB10, sm_121, CUDA 13.0, driver 580.178.04).

The five new gates, all mutation-checked.

• the_graph_destroy_failure_message_names_the_matching_destructor (pure) — the Rust formatter names cudaGraphDestroy, the cudaGraph_t from cudaStreamEndCapture handle, the cudaGetErrorName symbol and the ticket, and must not name cudaGraphExecDestroy or "kernel launch" (a gate that only asserts "a message appeared" cannot see a wrong message). Mutation: naming the exec destructor fails it.
• the_injection_matcher_matches_only_the_named_site (pure) — all, exact comma-separated tokens, whitespace-tolerant, and no substring rule in either direction (small must not arm attr:mmq_nt; attr:mmq_nt_extra must not either). Mutation: the substring form fails it.
• cuda_issue147_attribute_sites_name_the_call_and_refuse_the_launch (device; env-gated behind MINFER_TEST_ISSUE147=1 because it makes a real CUDA call fail) — for each of the ten dynamic-smem tokens the injected over-limit cudaFuncSetAttribute is reported once, at that site, with the API, the attribute, the device limit, a request above it, the exact instantiation and cudaErrorInvalidValue; the launcher returns 0; take_last_error() is 0. Then a positive control (knob off) launches and leaves no latch.
• cuda_issue147_launch_sites_name_the_call_and_refuse_the_launch (device; env-gated) — the same for the ten launch tokens (an over-limit dynamic-smem <<<>>>, which the launch call itself rejects with cudaErrorInvalidValue — probed before use, /tmp/fix147_launch_fail_probe2.cu — so the kernel never runs), plus a positive control.
• cuda_issue147_graph_destroy_failure_is_named_and_the_exec_survives (device; env-gated) — the injection re-creates #145's bug at this site (cudaGraphDestroy(exec)), the site clears the latch, and the valid exec is still returned and still destroyable (a failed graph destroy leaks only the graph handle, so refusing the exec would be wrong).

Mutation evidence. Every hardening was reverted one at a time and the corresponding gate re-run; each reverted version failed, the file was restored, and the restored files were byte-identical (sha256sum, in the closing comment on #147). Round 1 (per-site attribute guards): A1 attr:mmq_raw_nt_kd4, A2 _kd8, A3 the mmq_nt macro, A4 attr:mmq_raw_nb, A5 attr:mmq_raw_nb_bt, A6 attr:mmq_raw_nb_bt_q6k, A7 attr:mmq_raw_wide_kd4, A8 _kd8 — all eight trivially failed the attribute gate at that site. Round 2: A9 the gemm opt-in guard, B the shared minfer_launch_ok check, C the gemm message naming a wrong instantiation (the "prints a wrong message" check), D minfer_smem_optin reporting but admitting the launch (the shape a message-only assertion would miss — the gate's ret == 0 catches it), E the Rust destroy read discarded, F the Rust formatter naming the wrong destructor, G the injection matcher weakened to a substring.

Honest scope. The over-limit skips and the device limit are measured on GB10/sm_121 only; the decision is a runtime query and every skip is printed, so another device skips a different subset. All deliberate-failure evidence is device-gated and driven by test-only env knobs, which invert a site's own answer rather than exercising a genuinely failing driver call in production — the injection makes the call fail for real (over-limit attribute / over-limit dynamic smem / wrong destructor), so the "latch cleared" half is exercised for real, but the production paths remain latent by construction. The MMQ launch check is shared by the mmq_nt and raw-NT launchers, so the per-launch check was mutated once (B) rather than per launcher; every launch token's own report was observed by the gate. Other void launchers in the file — 65 of them (launch_dequant_f16, launch_convert_f16, launch_gemm_qb_nt, every store/rope/attention/MVQ/MVQ-multi wrapper) — still do not read their own <<<>>> error and were left alone: a failure there is currently reported by the next hardened launch's prelude or by sync as a latched API error, and converting all of them is filed as #162 rather than smuggled in here. And the pre-#147 fallback semantics are preserved deliberately: launch_mmq_raw_nb_nt / _nb_bt_nt / _q6k_nt / _wide_nt still return 0 on any opt-in refusal, so the dispatch falls to the next kernel — the failure is named, not fatal, because a fallback is the designed behaviour and changing it would change which kernel runs.

Every <<<>>> reads its own launch error · #162 — DONE 2026-09-26

Why. The #147 audit enumerated the wider case: 104 <<<>>> sites in src/cuda_kernels.cu (the 65 wrappers in the ticket plus multi-site families and two wrappers the #147 list had not counted) enqueued a kernel and never read the launch's error. A launch that failed for real — an illegal grid/block shape, an out-of-resources configuration, a stale context — latched the error, and it surfaced at CudaState::sync or at the next hardened site's pre-launch check as a latched API error (#145's honest label) with no indication of which launch produced it. grep -c '<<<' on the tree is 122; two of those are the <<<>>> in prose comments, so the audited surface is 120 sites in 76 launch_* owners (the wrappers plus the static launch_gqa_attn_split_batched_kv helper).

What landed.

• Every site reads its own error, through the shared helper. Each <<<>>> is preceded by minfer_launch_prelude(site, kernel) (which reports any pre-existing latch as not this launch's) and followed by a minfer_launch_ok / minfer_launch_ok_opt read that clears the latch it named. The site token starts with launch: and the report names the kernel instantiation and cudaGetErrorName — e.g. kernel launch gemm_f16_nt_kernel_t<128,32,true> failed: cudaErrorInvalidValue (1) — the launch is refused (#162/launch:gemm_f16_a32).
• The decision per launcher is severity in the helper, not a signature change. minfer_launch_ok is required: it records a sticky failure that CudaBackend::execute_node — one Rust-side check (CudaState::take_launch_failure), not 104 signature changes and 104 Rust Err arms — drains on both arms and turns into an Err naming the site and the node, so the op never proceeds with a stale output (and the drain keeps a stale record from blaming the next node). minfer_launch_ok_opt names and clears without the sticky for a path with a documented fallback. Split of the 120 sites: 107 required → Err (67 launchers — the matmul, elementwise, rope/store-KV, embedding, dequant/convert, MMVQ, attention and split-attention families, plus the MMQ terminal launchers launch_mmq_raw_nt / launch_mmq_nt and the #147 "no later gate" ones) and 13 documented fallbacks (6 launchers — launch_fa_prefill_f16kv's three window modes fall back to the legacy attention kernel; launch_mmq_raw_nb_nt / _nb_bt_nt (kernels + k-split reduce) / _q6k_nt (kernels + k-split reduce) / _wide_nt are #147's clean fast-path fallbacks; launch_kv_move_rows returns non-zero to its Result caller). The choice per family is a comment at each family in cuda_kernels.cu.
• The injection lever is shared geometry, so a site's coverage is data. minfer_launch_block(site, dim3|unsigned) replaces the block argument at every ordinary site; when MINFER_TEST_CALL_FAIL names the site the block becomes 4096 threads (over the 1024/block limit) and the launch itself returns cudaErrorInvalidValue for real, the kernel never runs — probed on GB10/sm_121. The dynamic-smem launchers keep #147's minfer_launch_smem lever. Adding a site's token is therefore a one-line string in the driver, not a bespoke mechanism.
• A scripted audit, wired into CI. scripts/check_cuda_launch_returns.py parses the source (comments and string literals blanked, so a commented-out occurrence is not a site) and requires, for every <<<: an enclosing minfer_launch_prelude("<site>", …) before it, an minfer_launch_ok/_opt("<site>", …) after its statement naming the same token, a token that starts with launch:, and a real-failure lever (minfer_launch_block or minfer_launch_smem) in the launch geometry. It prints its result and exits 1 on any offending line. It runs in the check-docs job (python3 present; the CUDA container's is not guaranteed) together with its own --selftest and a --check-fixture against tests/fixtures/cuda_launch_sites.tsv, the committed site list (line, owner, token, kernel fragment) the device gate compares its driven set against.
• Two k-split reduce sites are separate tokens (launch:mmq_raw_nb_bt_ksplit, launch:mmq_raw_nb_bt_q6k_ksplit). They share the launcher with the kernel site, whose _opt read returns 0 before the reduce; arming the kernel's token would leave the reduce site unreachable, so the reduce has its own prelude/read/token and the driver arms it alone.

Acceptance results (GB10/sm_121, CUDA 13.0, driver 580.178.04; serial device runs).

The sanitizer row's command is compute-sanitizer --tool memcheck --target-processes all target/release/deps/minfer-<hash> --test-threads=1 (wrapping the test binary, not cargo — the cargo/test build tree is not instrumented and --target-processes all around the whole bash scripts/cuda_test.sh chain stalls the harness; the binary form is the one that completes, in 350.17 s).

Mutation checks (each reverted; src/cuda_kernels.cu restored to sha256 da2e00fb79442fcd03bd3618301b4014cd835f832bbafa4153b4af4283dcdbfb byte-identically). Full transcripts in the closing comment on #162; the list: (a) deleting the read at one single-site family (launch_add_f32) fails the audit and the device coverage test (the armed site reports nothing); (b) the same at one switch case (launch:embed_rows__q4_k) and (c) at one branch of a templated family (launch:gqa_attn_split_f16kv__hybrid_causal); (d) making minfer_launch_read report but admit the launch (return true) is caught by the severity test's assert_eq!(rc, -1) on the fa-prefill fallback — a message-only assertion misses it; (e) the message naming a wrong instantiation is caught by the fixture's fragment check; (f) the sticky removed from minfer_launch_ok is caught by cuda_issue162_required_sites_set_the_sticky_opt_sites_do_not and by the node-level test; (g) execute_node's unconditional drain removed is caught by cuda_issue162_the_err_arm_also_drains_the_sticky (an f16 matmul whose Rust wrapper returns Err leaves the sticky pending, so a drain only on the Ok arm would blame the next node); (h) minfer_launch_block made a no-op is caught because no armed site fails and every armed set observes the empty set; (i) an _opt site that also sets the sticky is caught by the severity test's take_launch_failure().is_none(); (j) the audit's lever check disabled is caught by the selftest's "read but no lever" case; (k) the fixture's kernel fragment changed is caught by --check-fixture. An equivalent mutant is recorded too: deleting the trailing cudaGetLastError() in minfer_launch_read changes nothing (the read's own cudaGetLastError already resets the latch), and the gate correctly stays green — the trailing call is belt-and-braces, not the clear.

Honest scope. Nothing was failing on sm_121 before or after: the sanitizer was already 0, so the production paths remain latent and the evidence is that every site can be shown to refuse and name a real failing launch. The injection is a test-only knob (MINFER_TEST_CALL_FAIL + MINFER_TEST_ISSUE162=1, unset in every default, bench and sanitizer run) that drives the site's own geometry illegal; it does not exercise a genuine driver fault. The source audit is static: it proves the read and the lever are written, not that they run — a <<< inside a string literal or a macro the parser cannot resolve is reported, never silently accepted, but the parser resolves a site variable only through a plain = assignment in the enclosing function (the launch_site ternary of launch_gemm_f16 resolves to its first arm, which is why the driver arms gemm_f16_a32). The device gate's coverage assertion compares sets of site tokens; where two source sites share one token (the two mmq_raw_nb_bt kernel instantiations) the gate proves the token is reached, not that both branches were — the audit, not the gate, is what guarantees each source site has its own read. The compute-sanitizer and real-model rows are recorded measurements on dgxspark (CI has no GPU); the x86_64 CPU row is unaffected because no pure-Rust test was added.

C5 — Session save and restore · #43 — DONE 2026-09-22

Why. A session's KV rows, ownership and run table lived only in memory, so every restart re-prefilled the whole context.

What landed:

• src/graph/kvsession.rs is the container. An 8-byte magic, a version, flags, the backend tag and the shape (n_layer, n_ctx, n_embd, row_elems), then one K/V blob per layer as raw little-endian pool words, then the bookkeeping, then an FNV-1a checksum. It streams both ways: a save never holds a second copy of a 6 MB–1 GB arena, and a load materializes one layer at a time.
• KvCache::session_state / restore_session is the bookkeeping half — the owner table, n_used, the run table (SeqSlot: start, cap, shared prefix, written extent), each sequence's span list, identity and the C3/C8b counters — kept separate so it round-trips in unit tests without a device. restore_session validates before it applies: arena capacity, the owner table's length, every reservation and span inside the arena, and every live sequence carrying a span list.
• GraphAllocator::kv_save / kv_load move the bytes through the backends' read_host/write_host (CUDA included) and enable the pool a file names if it does not exist yet, because a restore happens before the first graph — which is also how the CUDA run of the gate found the gap. The header records the KV element type (f32/f16/q8_0) in its flags word (0/FLAG_F16/FLAG_PACKED — C5 S3, #130), so a file written under one width cannot be resumed under another.
• A failed load is a no-op. kv_load runs kvsession::verify — a full pass over the header, every layer's declared length, the bookkeeping, the checksum and end-of-file — before ensure_kv creates a single region. Every refusal path is asserted to leave no arena behind (kv_n_used(0).is_none()).
• Truncation is caught by construction, not only by the checksum: every read is exact against a length the header fixes, so a short file fails wherever it stops with "truncated — the file ends inside …". A version bump, a bad magic, trailing bytes, an unknown flag, a header whose packed flag and cell width disagree, and a file whose backend / n_ctx / n_embd / element type does not describe this run are each refused with the reason.

Acceptance, as measured (2026-09-22):

• Container, in CI: 9 unit tests — a round trip including the run table, a packed session, truncation, a version mismatch, a flipped byte (checksum), trailing bytes, a foreign file, a header whose flag and width disagree, and a writer short a layer.
• Allocator, in CI: the region bytes and the run table survive a save and a load into a fresh allocator, the restored table resolves the same cells, and four refusal paths (another n_ctx, another row width, another backend, another element type) plus a truncated file each leave the allocator untouched.
• Real model (0.5B q4_0, #[ignore]d, run on the CPU and on CUDA): prefill, save (24 layers / 256 cells / 5 written / 6 316 748 bytes), drop the cache, restore into a fresh one, then 8 greedy steps against the session that never left memory — max |Δlogit| = 0, i.e. bitwise, which is the strongest form of "the same continuation".
• Handed off: resuming the CLI's --session (which still re-prefilled its history JSON) from this container, and an E2 slot-table snapshot for the server, are #89.

C5 S2 — the CLI resumes, the host state rides along (2026-09-24).

The container described the KV rows; the rows belong to a host state, and a restore that brought one back without the other would be a different session. S2 closes that gap for the CLI and records what the server still needs.

• The container gained a host-state section (version 2). An opaque, length-prefixed blob after the bookkeeping and inside the checksum (KvSessionWriter::set_host, KvSessionReader::finish → KvSessionBody), so the two halves travel together or not at all; a version-1 file has no such section and is refused loudly, which is exactly what the caller's fallback path is for. kv_save_with_host/kv_load_with_host carry it through the allocator; kv_save/kv_load stay as the no-host wrappers every existing gate uses.
• Conversation::snapshot / restore_snapshot is the host half: the message list, stream_tokens, current_pos, turn_pos, prev_tokens and need_insert_eot, as a versioned JSON blob (SNAPSHOT_VERSION), with validation before it is applied — a snapshot whose current_pos disagrees with the token mirror, or that does not fit this run's n_ctx, is refused rather than half-applied.
• --session FILE (under --cnv) writes FILE.kv on exit and resumes it on start. Everything that could make a resume wrong falls back to re-rendering the JSON with the reason printed: another n_ctx/model/MINFER_CACHE_TYPE, a history the user edited between runs (the snapshot and the JSON must agree), an unknown snapshot version, a missing companion, an engine that cannot hand its KV to the host (a speculative session — the draft keeps its own KV — and the mocks), and any file the container itself refuses.

Acceptance, as measured (2026-09-24, Qwen2.5-0.5B Q4_0, CPU, --n-ctx 1024, 1530-character history, greedy, -n 8):

• Continues alike: turn 2 in a new process resumed from the companion answers The answer to 2+2 is — byte-identical to the same turn in the process that never left memory, and to a run with the companion moved away (the re-seed path). So the resume is behaviour-preserving, not merely fast.
• And prefills nothing at startup: the resumed run prints resumed 2 message(s) and 379 KV row(s) from …seed.json.kv (25 268 624 bytes) — 0 tokens prefilled. Wall clock to the identical continuation: 0.47 s resumed vs 2.36 s re-seeded (5.0×), the difference being the ~370-token history the JSON path re-renders.
• Mismatches are loud: --n-ctx 512 against the 1024-cell companion prints KV session: the file describes a 1024-cell arena, this run has 512 (--n-ctx); re-seeding the history instead and continues correctly.
• Gates, each mutation-checked: a_resumed_snapshot_prefills_nothing_and_continues_alike (the restored run issues the same engine calls, at the same positions, as the in-memory run — forgetting prev_tokens in restore_snapshot fails it), the_host_state_round_trips_and_is_covered_by_the_checksum (a flipped byte inside the host blob is refused; dropping the host from the writer fails it), a_version_1_file_is_refused_so_the_caller_can_re_seed, a_snapshot_that_contradicts_the_host_mirror_is_refused, and an_engine_without_a_kv_refuses_the_session_calls. Suites: CPU 287 passed / 0 failed / 15 ignored (was 282/0/15).
• Not verified here: a CUDA or Metal session companion (the container is backend-tagged and the CPU path is what dgxspark measured; the CUDA half of C5's own gate was run at S1).

C5 S2b — the server's slot snapshot (2026-09-24).

server::batch keeps a per-slot table (seq, start, cap, cached_tokens) that admission rebuilt from scratch, so a restart dropped every slot's context even though the rows were recoverable in the same container.

The design decision, stated because the issue's wording admits two readings. A Run holds the response channel to a client that a restart has already disconnected, plus its RNG and the row's logits — none of it survives a process boundary, and "resuming" a stream nobody is reading would be a fiction. So the snapshot carries the context, not the in-flight request: each slot's reservation and the token sequence its rows hold. A request that arrives after the restart with the same prefix is admitted onto the restored rows and prefills only its own delta — B2's cross-request prefix reuse, with the rows coming from disk. That is the property the acceptance measures ("resume the slot, and get the same continuation").

• SlotRow / SlotsSnapshot (server/batch.rs) are the table, as versioned JSON in the container's opaque host section — the same mechanism S2a added, so the two halves cannot drift apart and a version this build does not know is a refusal.
• BatchEngine::save_slots(path) writes the table and the arena through kv_save_with_host; load_slots(path, model) loads it back, validates it against this run (another --n-slots/--n-ctx, another model, another KV element type, a table whose sequence ids or written extents disagree with the arena it rode in with), and installs the slots.
• When it is written: after every completed request (finish). That is when a slot's context is stable and it is the last moment the rows are known-good — so a server killed without warning still resumes the conversations that had finished, which shutdown-time saving would not give. The cost is one arena write per completed request (--slots-file prints the size at startup; ~12 MiB for the 0.5B/512-row fixture), which is why the flag is opt-in.
• --slots-file <PATH> (CLI) → server::run → the worker, which loads at startup and rewrites on completion, printing what it resumed or why it started empty. The serial (non-batched) path has no shared arena to snapshot and says so loudly instead of writing nothing quietly; the batched engine is the one with cross-request prefix reuse to restore.

Acceptance, as measured (2026-09-24, Qwen2.5-0.5B Q4_0, CPU, 512 rows over 2 slots):

• Resumes without re-prefilling: the cold run feeds 5/5 prompt tokens; the restored engine feeds 1/5 (4 reused) for the same prompt, and its next turn — a continuation carrying the whole conversation — feeds 6/11 (5 reused). A re-render would have fed all 11.
• Same continuation: the restored run's generated text and token count equal the run that never stopped (a_slot_snapshot_resumes_the_context_without_re_prefilling, #[ignore]d, on the cached model).
• Refusals, naming both numbers: a 2-slot snapshot into a 1-slot engine (--n-slots), and a 512-row snapshot into a 1024-row server (--n-ctx).
• Gate, mutation-checked: dropping the token mirror from load_slots makes the resumed prompt feed 5/5 again and fails the gate. Suites: CPU 288 passed / 0 failed / 16 ignored (was 287/0/15).
• Sharing: kvformat.rs gains expect_for(model, n_ctx) / backend_of(device), and the CLI's --session engine now uses them too, so "what this file must match" has one definition.

C5 S3 — the container encodes the f16 element type · #130 — DONE 2026-09-25.

The header already carried the KV element type in its flags word, but only Q8_0 had a bit (FLAG_PACKED). An f16 region therefore wrote flags == 0 and read back as f32, so kv_load's header.format != live check refused the file its own writer had just produced — on any CUDA box whose model crosses the f16 auto-policy threshold (Qwen3-0.6B: 28 × 1024 = 28672 ≥ 8192), for --slots-file and --session alike. The real-model gate showed it as "the file was written with the f32 KV element type, this run uses f16".

The design decision: a flag bit, not a format field. FLAG_F16 = 1 << 1 is added, and flags_of / format_of_flags are each other's exact inverse. A format field would have had to either renumber FLAG_PACKED — changing the meaning of every Q8_0 file already on disk — or move the type elsewhere in the header, shifting the byte layout a version-2 reader is already parsing; both are the silent misread the container exists to prevent. A new bit keeps the flag word's layout (and every existing file's meaning), and the compatibility story follows for free: a pre-#130 build reading an f16 file sees an unknown bit and refuses it loudly at the unknown-flags check instead of decoding the region as f32, while a pre-#130 f32 file (flags == 0) still loads as f32. The two element-type bits are mutually exclusive — packed | f16 describes two incompatible cell layouts, so it is refused by name, never resolved by preferring one bit. No version bump: VERSION stays 2, because a bump is for a layout change and this is an additive flag whose older-reader behaviour is a loud refusal — exactly as FLAG_PACKED was when it landed.

What landed.

• src/graph/kvsession.rs: FLAG_F16, KNOWN_FLAGS, the flags_of / format_of_flags pair, the writer encoding the format into the flag word, and the reader's unknown-bit + mutual-exclusion refusals before it decodes.
• Six new gates, five in kvsession.rs (no device — the CI-covered half) and one in alloc.rs (kv_save → kv_load under an f16 policy): the f16 round trip, the flag-word sweep over all three formats (encode and decode), the both-bits refusal, the unknown-bit refusal, the legacy flags == 0 → f32 load, and the allocator-level f16 round trip.

Measured (GB10 sm_121, CUDA 13.0, serial).

Mutation checks (each reverted byte-identically, sha256sum): breaking the encode (f16 → 0) fails the flag sweep, the container f16 round trip and the allocator f16 round trip (435/3); breaking the decode (dropping the FLAG_F16 branch) fails the same three; removing the mutual-exclusion check fails only the both-bits gate — and it fails on the message, because the width check would otherwise refuse the file for a different reason (that is the "passes for the wrong reason" hazard this gate's assertion closes); widening KNOWN_FLAGS to !0 fails only the unknown-bit gate; and encoding f32 as FLAG_F16 fails the legacy gate plus the two f32 gates. rustfmt --edition 2021 --check clean on both files (stable rustfmt 1.9.0 — the pinned 1.97.1 toolchain has no rustfmt component here; CI runs no fmt job).

C6 — Logical positions (positions ≠ cells)

Why. §14 row 9 closed the cell-offset sensitivity by measurement: a sequence's logits' tail changes when its run moves because positions is both the RoPE angle and the KV cell index, so a cell move shifts every rotation. The intervention experiment proved that RoPE is the only entry (offset_divergence_is_caused_by_the_rope_rounding_alone: injecting run A's 48 rope outputs into run B makes the logits bitwise identical), and the sweep showed the effect saturates at a ~1e-6 distributed perturbation (a_distributed_rope_perturbation_saturates_the_logits_tail). Two consequences drive this ticket: C3's compaction needs a K re-rope (and can never be bit-identical), and any future cross-slot sharing inherits the same coupling.

Semantics. positions becomes what a caller naturally has — the token's index within its sequence — and the allocator resolves, per token, the KV row it must be written to:

The single-sequence case is unchanged by construction: its run starts at cell 0, so cells[t] == positions[t], and the classic path stays bitwise.

Invariants / gates.

1. every existing single-sequence test stays bitwise (regression);
2. batched forward == the same sequences run one at a time, same layout, bitwise (the existing gate, migrated to relative positions);
3. a compaction between steps leaves the logits bitwise identical — the literal acceptance C3 could not claim before, and the reason this ticket exists;
4. a backend that cannot resolve cells refuses the node (Err, no silent fallback); Metal is out of reach by construction because it already refuses multi-sequence attention (supports_attn_span() == false), so it never sees a non-zero run start and keeps positions == cells (recorded as G5, no Metal code change in this ticket);
5. decode timing on GB10 does not regress (the fused QKV family is the only hot path this touches, ported in S3 with its own A/B).

Change surface.

• IR: GraphBuilder::kvcache_store wires its row input to a new cells node (created like attn_span); the fused QKV family (FusedQKV, QkvBiasRopeStore, FusedQkvNorm) is gated off while explicit_span is set (S1) and takes cells in S3.
• Allocator: fill_batch_inputs / fill_attn_inputs / fill_seq_ids stop treating positions as cells — own_range, kv_note_used and attn_span resolve through the run table — and fill cells. (History: E1's test-only fill_attn_inputs and the kv_note_used → own_prefix C1 remnant were deleted in #228; fill_batch_inputs is the one production fill entry point.)
• KV store: KvCache::cells_for (today the identity) and attn_span (today requires position ∈ run and owner[position] == seq) become the resolver; check_positions_bound / check_attn_span check the cells form.
• Models: forward_batch's explicit_span decision and positions stay as they are semantically; the store wiring is inside kvcache_store.
• Server: positions = slot.start + current_pos becomes current_pos; SlotState.start drops to reporting/diagnostics.
• Backends: the KvcacheStore arms need no change — they already write at the rows their third input names; only that input's producer changes. This is what makes S1 landable with the suite green on all three backends.
• C3: kv_defrag(need, rope) loses the rope parameter, and the KvRope plumbing added for the re-rope goes away.

Staged plan (each step: local CPU suite, plus the CUDA suite when it touches CUDA, then a docs progress update, then a commit on feat/logical-positions).

The docs build gate is the CI job check-docs (.github/workflows/ci.yml): mdbook build with the deployed toolchain plus scripts/check_docs_links.py. It did not exist when S0/S4 ran — the docs were checked by hand then — and it landed with #63; the rows above name it so the record and the gate agree.

C6 progress (2026-09-19)

S0 done (71b86da). S1 + S2 landed — the semantics switch and the re-rope removal went in as one commit, as the merged-step row above requires. Suites: CPU 213 passed / 0 failed / 5 ignored, CUDA 261 passed / 0 failed / 5 ignored.

What S1 changed (production):

• GraphBuilder::cells_input; kvcache_store creates and consumes it, so callers no longer pass a row buffer (26 call sites migrated);
• KvCache::attn_span resolves cell = run.start + position; GraphAllocator::kv_cells_for_seq fills cells, with the classic single-sequence identity fallback when no run is reserved;
• fill_batch_inputs / fill_attn_inputs / fill_seq_ids use relative positions, and fill_attn_inputs only resolves cells when the graph has that input (a rope-only fixture must not need a KV arena). (History: fill_attn_inputs was deleted in #228, which moved its corpus onto fill_batch_inputs; the rope-only branch survives as the #[cfg(test)]-scoped GraphAllocator::fill_attn_inputs_without_cells.)
• qwen2/qwen3 gate the fused QKV family off while explicit_span is set (S3 ports it for CUDA), and server/batch.rs passes relative positions;
• kv_defrag no longer re-ropes and the KvRope plumbing is gone (S2).

What the tests became, i.e. the new gates: the minimal hand-built attention graph asserts bitwise invariance to the cell placement (it used to be "within the rotation's rounding"); the four offset-sensitivity tests invert into "a cell placement changes nothing" (logits, per-layer V, per-node q/k/attention); and the perturbation probe was deleted, because its premise — that the offset effect needs explaining — is gone.

S3 done — the CUDA fused decode QKV family takes cells:

• attn_bias_rope_store_f32 gained a cells parameter: pos rotates q/k and row = cells[0] addresses the four KV store writes. The launcher, the FFI declaration and CudaState::attn_bias_rope_store carry it; the builder wires the shared cells input into fused_qkv (sources [x, pos, cells]) and qkv_bias_rope_store ([q, k, v, pos, cells]), so no model call site changed;
• the qwen2 gate becomes (cuda_on || !explicit_span): CUDA keeps the fused chain under an explicit span, Metal keeps the pre-C6 gate (it has no explicit-span attention at all, G5). Qwen3 needs nothing: its fused family is Metal-only, so no CUDA arm existed to port. The hand fixtures in cuda.rs::d38_probe_tests pass their positions buffer as cells — identity by construction, which is the case they assert.

Suites: CPU 213 passed / 0 failed / 5 ignored, CUDA 261 passed / 0 failed / 5 ignored. End-to-end gate on GB10 (Qwen2.5-0.5B Q4_0, --n-slots 2 --n-ctx 1024, one short + one long request submitted together): cap = n_ctx/2 = 512, so the long request ran in slot 1 (server log) and, once the short one answered in a token, alone — a single-token (nt == 1) step in a run whose start is not 0, which is exactly the path the S1 gate had closed. Fused vs MINFER_NO_FUSE_QKV=1 produced byte-identical completions (sha1 36c55ed81464, 319/319 expected words): the cells port and the S1/S2 KvcacheStore path agree on a non-zero run start.

The gate also caught a real server bug, now fixed: BatchEngine::submit_on still built positions as start + i — the pre-C6 "positions are cells" semantics — so a request placed in a non-zero-start slot fed position 512 into a 512-cell run and kv_cells_for_seq rejected it ("past sequence 2's reserved run"), rejecting the job. Both the batch prefill path and current_pos were already relative; only this entry point had not been migrated, and no earlier test exercised a second slot's run (the CPU suite batches nothing by default, so every server test ran the single-sequence path). It now passes feed_from..nt.

Regression coverage: server_batch_matches_serial_and_is_faster pins each of four requests to the slot it would occupy through submit_on, so it drives runs with non-zero starts — it is #[ignore]d only because CI has no cached model, and it must be run locally for any change to positions or the run table: cargo test --release -- --ignored server_batch_matches_serial. It passes both on CPU (byte-equality of batched vs serial, 1.43x) and on CUDA (structural assertions, 1.28x).

One measurement note worth keeping: an end-to-end A/B of the fusion on the batched decode path cannot isolate this gate, because fuse_qkv requires nt == 1 — a step carrying four streams is unfused in both configurations. The measurement that matters is the single-stream step in a non-zero-start run above; a 4-slot throughput A/B (master, 7B Q4_K_M, 4 concurrent requests) showed the fused and unfused chains within noise (92.6 vs 92.7 tok/s best case), which is expected for a bandwidth-bound model and is why the port's value is correctness and parity, not throughput.

The two pre-C6 experiments that established the attribution above — the RoPE injection (EXP1) and the distributed perturbation sweep (EXP2) — are archived with their code and measured numbers in experiments/logical-positions/.

Not in this ticket: sharing one cell range across sequences (owner[cell] becomes a set/refcount — the real cross-slot prefix reuse, which this ticket makes possible), C4 (quantized KV) and C5 (state save/restore) — both should follow this, because each would otherwise multiply the addressing surface.

• Explicitly not changed: C2's kv_rm/kv_shift re-rope stays: it changes positions, not cells, so the angles genuinely have to move. Only the compaction re-rope disappears.

C7 — Dynamic runs: one request may use the whole arena · follows C6 · M · #59

Why (user-visible). BatchEngine::new hands every slot a fixed cap = n_ctx_total / n_slots up front, and submit_on rejects anything whose prompt + generation exceeds that cap (prompt of N tokens exceeds slot context of M). With --n-slots 4 --n-ctx 8192, a single 6000-token request is refused while three slots sit idle and their rows are free: the arena is sized for the total, the partition is the limit. This is the most user-visible consequence of the fixed reservation and the reason C3's "server-side dynamic-run trigger" was recorded as a follow-up instead of being closed with C3.

What. Make the partition elastic instead of fixed:

• a slot may grow past its initial cap when the admitted request needs it, via kv_reserve_seq_with_defrag(seq, need) — C3 already plans the move, copies the rows through Backend::copy_cells and returns the moves;
• the shrinking slots are re-reserved around the growth, and every caller that cached a run start (E2's SlotState) applies the returned moves — the contract kv_defrag already documents. C6 is what makes this safe: a moved row keeps its sequence-relative position, so no other slot's logits change and nothing re-ropes;
• the repartition happens at one serialization point (admission), so no in-flight step can observe a half-moved arena;
• a genuinely full arena still rejects loudly with the existing message shape — one request never silently shrinks another's context.

Gates. (1) --n-slots 4 --n-ctx 8192 serves one request of ~8k tokens while the other slots are idle (the literal requirement); (2) after a grow, each idle slot's next request is bitwise identical to the same request on a fresh engine; (3) the 4-slot batched-vs-serial byte-equality test stays green (server_batch_matches_serial_and_is_faster, run locally with --ignored); (4) the CUDA batched-decode throughput stays within noise of the recorded numbers.

Not in this ticket: sharing (C8) — C7 only makes the partition elastic, two sequences still never look at the same cell.

C7 increment 1 (2026-09-20) — the partition is elastic

What landed. BatchEngine now sizes a slot from the request instead of from the startup partition: wanted_cells_from (pure, unit-tested) asks for prompt + max_tokens + 1 for a bounded request and prompt + slot cap + 1 for an unbounded one, clamped to the arena and never below the prompt. When that exceeds the slot's cap, ensure_slot_capacity reclaims the idle runs above it (they hold at most B2's prefix hint, which the slot re-prefills when it is next used), then re-reserves this slot's run at its existing start and re-owns its written prefix; the moves the KV store returns are applied to every slot's cached start (apply_run_moves, the contract C3 documented).

No new KV primitive was needed — kv_release_seq, kv_reserve_seq_with_defrag and kv_own_range compose into "grow in place" — and no backend changed, because the store already writes at the allocator-resolved cells (C6). A reservation that would land anywhere other than where the rows physically are is refused and the slot is left to re-prefill: the one outcome C6 exists to prevent (reading rows at an offset they do not have) cannot happen. The HTTP layer's request bound moved from a slot's share to the whole arena (AppState::n_ctx); the serial path (batching off) keeps its own slot-sized bound, because its graph region really is that size.

(Forward note, 2026-09-29: kv_own_range in the sentence above is GraphAllocator::kv_own_range, which #232 deleted as dead surface — it was a one-line forwarder to KvCache::own_range, the function production actually calls (through fill_batch_inputs → own_positions, and through kv_copy_prefix). The sentence records what this step landed with; it is not a live API.)

The boundary, and why it is no longer padded. tick forwards the committed token before advance can report length, so a generation that reaches its cap used to perform one more forward at the cell after its last token — a whole weight pass whose logits are discarded, plus a cell nothing reads. The reservation was therefore padded with one cell, and a reservation of exactly prompt + max_tokens rejected that batch (kv_cells_for_seq: position N is past sequence S's reserved run).

That padding is gone: advance now decides before committing whether another token can still be used (the request's token budget, and current_pos + 1 < cap), and ends the turn instead of committing one that could only be discarded. wanted_cells_from is exact — prompt + max_tokens, clamped to the arena — and the top-of-advance bound is the safety net it always was. The regression is a request with no token budget (max_tokens = -1) on a run sized to its prompt plus 48 cells: it must end with length and exactly 48 tokens, because a forward past the run is rejected by kv_cells_for_seq — before this change that request failed instead of finishing.

Gates. (1) wanted_cells_plans_from_the_request — the policy, no model needed; (2) a_long_request_may_use_the_whole_arena — a 302-token prompt on n_ctx = 366 with four slots (91 cells each) is served, and both admission paths (submit / prefill_group and submit_on, the one the server uses) produce a continuation byte-identical to a one-slot engine that needs no reclaim; (3) the four-slot batched-vs-serial byte-equality test still passes; (4) CPU and CUDA suites green; (5) on GB10, --n-slots 4 --n-ctx 8192 with a 2054-token prompt logs slot 0: capacity 2048 -> 2071 cells ... (released 3 idle slot(s) above) and answers the same bytes as --n-slots 1, which needs no reclaim at all. Gate (5) is the point of the ticket: C6 makes a moved row arithmetic-free, so where the partition puts a sequence cannot show up in its output.

C7b — both directions, so a busy neighbour is not a wall (2026-09-20). The increment-1 caveat is gone: Backend::copy_cells and CUDA's kv_move_rows kernel now move rows up as well as down (the kernel walks one row at a time with a barrier, descending when the run slides up), apply_moves accepts to > from and frees the vacated cells in either direction (guarded by the sequence's own stamp, so a plan can never clear another sequence's ownership), and KvCache::set_cap grows or shrinks a reservation and returns the plan — shrinking below a sequence's written rows is refused, and so is a growth the arena cannot hold with the other reservations. The allocator's kv_set_cap_with_defrag copies the rows and then renumbers, like kv_defrag.

The subtle half is order: wherever ranges overlap, a destination must never land on a row that has not been copied yet. kvcache::order_moves is the single rule — upward moves top-down, downward bottom-up, upward first — and both the data copy and the bookkeeping follow it (the owner table is a mirror of where the rows live, so they have to travel together). The unit test caught exactly this: applying the plan in its raw ascending order moved one sequence's stamps through cells another had already overwritten.

Gates. (1) growing_a_run_pushes_the_runs_above_it_up — a pure test of the plan, the upward move and the owner stamps; (2) a_resize_refuses_to_eat_rows_or_overcommit_the_arena; (3) the CPU and CUDA twins that used to pin "upward is refused" now pin the result of an overlapping upward move (a memmove's output) — the CUDA one on GB10; (4) on GB10 the CUDA twin pins the upward kernel path (the bytes a memmove would produce). Suites: CPU 216 passed, CUDA 264 passed (0 failed, 6 ignored each).

The cells bound is its own rule (2026-09-21) — #60. check_positions_bound classified any I32 input consumed by a KV-writing op as positions, and cells — which now feeds the same ops (C6) — was measured by that rule. It cannot false-reject (a cell always indexes the arena, which is exactly n_ctx cells), but the two bounds only coincide because a cell equals its position, which is the coincidence C8 removes. cells now has its own bound and its own message (input 'cells': cell N is past the M-cell arena), pinned by a test that also covers the no-arena case.

End-to-end (2026-09-20) — the engine scenario is verified. Four slots at --n-ctx 8192: a short request, then a long generation on another slot, then a 2148-token prompt admitted while that generation is still running. On GB10 the server logs slot 0: capacity 2048 -> 2157 cells for a request wanting 2157 (released 1 idle slot(s); 2 run(s) moved) — two runs moved, one of them the live generation above — and on CPU (forced with MINFER_BATCH=1, which makes the comparison deterministic) the same scenario moves one run. In both runs the live neighbour's continuation is byte-identical to the same request served alone: its rows travelled up under it and its answer did not change, which is exactly what C6 + C7b promise. The grown request itself is compared structurally (a shared 38-byte opening with its one-slot baseline), because it shares the batch with the live request and therefore takes the windowed (explicit_span) attention path while the baseline takes the causal one — a named tolerance class, not a bitwise property.

That gate was missed twice for an environmental reason worth recording: cargo test --release overwrites target/release/minfer with a CPU-only build, so the server logged batching: off (device cpu) and served the serial path while the script believed it was testing the engine. The script now asserts batching: on and the device line in the log before it measures anything — a gate that cannot check its own preconditions is not a gate.

C8 — Cross-sequence cell sharing (owner → set/refcount) · follows C7 · L · #41

Why. A prefix shared by several sequences (a system prompt on every slot; B2's prefix reuse, but across slots) is duplicated today: each slot stores its own copy and pays its own prefill, so N sharers cost N× the memory and N× the prefill for the same tokens. llama.cpp's unified cache shares one cell between sequences through a per-cell sequence set (seq[i] is a bitset, with seq_cp/seq_rm/seq_keep over it); this ticket is that capability, scoped to what minfer's server needs.

What. KvCache::owner[cell] becomes a set (bitset or refcount); attn_span and kv_cells_for_seq accept "owned by any of these sequences"; a new kv_seq_cp(src, dst, range) shares a prefix's cells with another sequence; kv_rm frees a cell only when its last owner drops it; and the store must not write into a shared prefix — the first divergent token copies the shared cell it would have overwritten into a private one (copy-on-write). The CoW rule is the delicate part: a store that wrote through a shared row would corrupt every sharer, so it is either implemented or refused, never ignored.

Gates. (1) two slots sharing a prefix produce continuations bitwise identical to the same two slots run with private copies (CPU first; CUDA may take a named tolerance class only if a kernel shape differs); (2) removing one sharing sequence leaves the other's logits bitwise; (3) a store that would write into a shared cell copies it or fails loudly; (4) KvArenaStats counts a shared cell once, not once per owner (memory-footprint measurement).

Depends on: C7 (elastic runs — shared prefix plus growth is what the server actually needs); C1's owner table and C6's resolver are the hooks.

C8 design (2026-09-21) — the read path, not the refcount, is the cost

Why. A prefix used by several sequences (a system prompt on every slot, a conversation resumed twice) is prefilled and stored once per sequence. B2 reuses a prefix within a slot; across slots there is no reuse at all. KvCache::owner[cell] holds a single SeqId, so a cell belongs to exactly one sequence.

The constraint that shapes everything. A sequence's rows are one contiguous run, and the read path depends on that: cells[t] = start + position, and attention reads a single [lo, hi) span from attn_span. The write path is already general — the allocator hands the backend a per-token cells vector, so a store may target any row (C6) — but a read cannot follow a sequence whose rows are not contiguous. Now let two sequences share a prefix and then diverge: only one of them can keep a contiguous layout (its tail sits right after the prefix); every other sharer's rows become [0, p) ∪ [private, ...) — two ranges. So true sharing is not "a refcount in the cell store"; it is a paged read path (a per-sequence block map and a gather in the attention kernels of each backend). That is the whole cost of this ticket, and it is why it was estimated L rather than M.

Two increments, because the cheap half is independent of the read path.

Invariants / gates.

1. C8a: a prefix copied from another slot produces continuations byte-identical to the same prefix re-prefilled, on CPU and on CUDA (the same named-tolerance caveat as any device comparison).
2. C8a: the copy must be measurably cheaper than the prefill it replaces (it is the point).
3. C8b: two sharers are byte-identical to the same two sequences with private copies.
4. C8b: kv_rm frees a block only when its last owner drops it; a store into a shared block copies it (copy-on-write) or fails loudly — never writes through.
5. C8b: KvArenaStats counts a shared block once, and a compaction moves it once (not once per owner); the memory saving is measured.

C8a increment 1 (2026-09-21) — the copy primitive. GraphAllocator::kv_copy_prefix(src, dst, rows) copies the first rows written K/V rows of one sequence's run into another's (per layer, through Backend::copy_cells, which handles either direction — C7b) and hands dst their ownership via own_range. It refuses a missing run, a source with fewer written rows than requested, and a destination that reserved fewer cells than that, and it is a no-op for zero rows or a same-run copy. Unit coverage is the region-free half of those refusals; the data path and the written/capacity bounds are covered by the S3 gate, because they need real KV regions and the honest test for a copy is "the answer does not change".

The engine is not wired to it yet. S2 does that at admission: placement already prefers the slot with the best-matching prefix, but only among idle slots and only for the slot's own rows — what is missing is that the donor may be another slot (possibly busy, and its rows are stable while it generates).

C8a increment 2 (2026-09-21) — admission uses it. The reuse source is now every slot, not just the admitting slot's own rows: when another slot holds the longer match, its written rows are copied into this slot's run (the donor may be busy — its rows are stable for the duration of the copy) and only the suffix is prefilled. The reuse is clamped to leave one token to feed, exactly as the slot's own reuse is, because a forward with nothing to run produces no logits. A failed copy falls back to the slot's own cache and prefills the rest — never a wrong-row read.

The test-only counter prefix_rows_copied makes the copy observable, so the gate asserts both that it happened and that it changed nothing: the same prompt served via a copy and on a private single-slot run answers byte-identically (CPU, where the comparison is exact; a device comparison takes the usual named-tolerance caveat). C8a cost gate (2026-09-21, GB10). With a 1450-token prompt on --n-slots 2 --n-ctx 8192 and the donor held busy (a second request generating), the same prompt served through the copy costs 0.013 s against 0.467 s with prefix reuse switched off (MINFER_NO_PREFIX_REUSE=1) — 36.6× — and both answers are identical. Two scenario traps had to be fixed before the number meant anything, and both are easy to repeat: (a) if the sharing slot is idle, the engine simply places the request there and reuses its own cache (B2), so no copy happens at all — the donor has to be busy; (b) if the donor's own request needs more than its share of the arena, C7's growth reclaims the idle slot the measurement needs, leaving one slot and self-defeating the scenario. Each copy is reported for an operator as [server] slot N: copied R prefix row(s) from slot M.

Order. C8a first: it delivers the user-visible half with the machinery that already exists (C3's row copy + C7's moves) and its gate is byte-equality. C8b is then a read-path project, and it can be scheduled on its own evidence rather than on this ticket's estimate.

Not in this ticket: C4 (quantized KV) and C5 (state save/restore) — each multiplies the addressing surface this ticket touches, which is why C6 preceded C7 for the same reason.

C8b design (2026-09-21) — a block map for the read path

What sharing needs that C8a does not. C8a duplicated the rows so the destination could keep one contiguous run. Sharing for real means a block exists once, several sequences refer to it, and a sequence's rows become a list of spans rather than one run (cells[t] = span_of(t).start + offset). The write path needs nothing new — since C6 it takes an arbitrary per-token cells vector. The read path is the whole change: attn_span hands each query a single [lo, hi) range, while a sharing sequence's window is a set of ranges.

Data. KvCache gains (a) a per-cell refcount at block granularity (64 cells — small integers, and the unit that compaction moves and the stats count), and (b) a per-sequence span list covering [0, written). A sequence that shares nothing has a one-entry span list, which is the property that keeps this incremental: its code path stays today's, bitwise.

IR — additive, never a replacement. A graph whose sequences share blocks gets a kv_map input (per query, the sequence's spans) alongside attn_span; attention uses the map when the input is present and the span otherwise. attn_span's single range is load-bearing — the CAUSAL fast path and Metal's refusal (G5) both rest on it — so it is not generalized in place.

Allocator. kv_cells_for_seq and attn_span resolve through the span list. A store that would land in a shared block takes a private row for that token (kv_private_row_for(seq, t)): the allocator decides before the forward, so no backend learns that sharing exists and a shared block is never written through. That is the copy-on-write rule, and it is why the store path did not have to change.

Backends. CPU attention gains the span-list gather first (it is the reference), then CUDA's gqa_attn_* inner loop, whose KV walk becomes (block, offset); Metal stays behind G5 and refuses the map, as it already refuses the explicit span.

Why S2 does not split the way S1 did (2026-09-21, found while sizing it). S1 split cleanly because its two consumers already existed — the write resolver and attn_span — so each half was a live path with a bitwise gate. S2's halves are coupled instead. Sharing means one cell has several owners, and the store's ownership representation is per cell: Layer::owner is a Vec<SeqId>, written_rows scans it contiguously for owner[cell] == seq, and attn_span's written check compares it against the query's sequence. A shared prefix cannot be expressed without changing what "written by this sequence" means (block refcounts plus the span list as the read-side authority), and until the read path gathers over several spans a sharing sequence would resolve a wrong window. So the store half and the read half have to land together: staging them separately would add machinery nothing reads, which is the shape A7 deleted. Concretely S2 = block refcounts (64 cells), kv_seq_cp, refcount-aware release/kv_rm, and a compaction that moves a shared block once and renumbers every sharer's span list — plus the additive kv_map input, the CPU gather, and admission wired to kv_seq_cp, accepted by gate 2 below.

Gates.

1. S1 changed nothing observable (S1a + S1b landed: both resolvers read the list, and the suites stayed bitwise); the refcounts arrive in S2 together with the readers that justify them.
2. Two sharers answer byte-identically to the same two sequences with private copies (CPU; CUDA takes the usual named tolerance).
3. kv_rm frees a block only when its last owner drops it, and a store into a shared block either takes a private row or fails loudly — never writes through.
4. KvArenaStats counts a shared block once; a compaction moves it once and renumbers every sharer's span list; the memory saving is measured.

Why S1 is a used generalization, not dead bookkeeping. The first draft of this plan put the refcounts in S1 and the span list in S2. That is the shape A7 deleted: the campaign's own note says an identity field earns its place only if some topology decision reads it — "a future feature will need it" is not enough, and n_seqs was removed for exactly that reason. Refcounts are only read by the sharing that S2 introduces, so they move to S2 and S1 becomes the span list plus the two resolvers that read it (kv_cells_for_seq and attn_span). Those are live paths for every request, which is what makes S1's gate meaningful: a single-entry span list must reproduce today's answers bit for bit, and any slip shows up in the existing suites rather than in a field nobody consults.

C8b S1a landed (2026-09-21) — the write path resolves through spans. KvCache carries a per-sequence span list (position base, first cell, length), and GraphAllocator::kv_cells_for_seq resolves every store row through KvCache::cell_of instead of slot.start + position. Every path that changes a run's start/cap or drops the run republishes the list (refresh_spans: reserve, release, resize, and each relocation apply_moves performs), and a position outside the list is a loud error, not a fallback to the contiguous form — which is what makes a missed maintenance point visible rather than silent. Today the list always holds exactly one entry, so nothing observable changed: the full suites keep their counts (CPU 218 / CUDA 267, 0 failed) and every byte-equality gate still passes — the four-slot batched-vs-serial test, the C7 boundary case, and the C8a prefix copy. The unit test pins the resolution against start + position, including a real compaction move and a release.