01 — mistakes we caught, wrong turns, and ruled-out ideas (do not repeat) #

2.7 mistakes we caught, wrong turns, and ruled-out ideas (do not repeat) #

A. our own wrong turns during planning #

  1. writing off ds4-recipe as "wrong model, no MoE tiering".
    • mistake: dismissed astrra.space/ds4-recipe at first sight because it targets DeepSeek V4.1 Flash on a 128 GB Strix Halo box instead of Qwen3.8-Flash-Next.
    • reality: once read (00, 06, 07, 13), it had the exact production measurements we needed for an SSD-streamed fine-grained MoE (~31% experts resident): MoCE cache-prior routing (−52…−73% misses), 12-tape offline trace replay against Belady's MIN ceiling, popularity EMA + near-pick credit eviction, non-evicting staged prefetch, 4×8192 chunked prefill, disk KV cache + SCR staleness (+0.015–0.040 nats/token), and the exactness/squeeze discipline.
    • rule: judge references by whether they share the bottleneck (SSD-streamed fine-grained MoE + long-context agent turns), not the model name.
  2. treating Strata's 93.6 tok/s Mac prefill as our pruned ceiling, and conflating chunk size with pruning.
    • mistake: cited Strata's 93.6 tok/s 4K prompt on a 32 GB M4 Mac mini as proof the laptop would be painfully slow, forgetting that number was the unpruned 512-expert model at default 1024-token chunks.
    • what we figured out: at top-10 of 512, a 1024-token chunk already touches almost every expert per layer (1 - (1 - 10/512)^1024 ≈ 1). so:
      • chunk size (1024 → 4096/8192) is what moves big cold prefills (~80–94 → ~200–300 tok/s) by amortizing one SSD pass over 4–8× more tokens.
      • pruning (512 → 224/256/320) is what moves the per-turn floor on short agent turns (~10 s → ~1.5–3 s) by shrinking non-resident expert bytes and raising the resident fraction.
  3. thinking ds4-recipe's 450 tok/s prefill meant kernel opts would make agent turns fast.
    • mistake: looked at ds4's 354–450 tok/s cold prefill and thought kernel work was the main path to usable speed.
    • reality: 450 tok/s is cold prefill on 32K chunks (4×8192 pipeline + hipBLASLt prepared GEMMs). on ds4's own 7-day telemetry, median prefill on 64–256 token turns is 19 tok/s (256–1K is 46 tok/s) because short turns are SSD-bound. what makes an agent fast in practice is 98.75–99.5% KV reuse (disk KV + SCR) plus cutting SSD bytes per turn (pruning + MoCE). kernel opts raise the cold-prompt ceiling; they barely move the per-turn floor.
  4. using ISTA's 256-expert Coder breaking on Chinese (Strata #438) as an argument against 224–256 experts for Japanese.
    • mistake: argued that since ISTA's 256-expert Coder broke on CJK, pruning to 224–256 would break JA.
    • what dawn caught: ISTA didn't target CJK at all. their calibration was code + agentic + EN + image, and their card explicitly says degradation outside the calibration set is expected.
    • decision: that failure only proves calibration decides what survives. put JA in the calibration mix and test 224/256/320 on a JA probe.
  5. worrying about pruned world knowledge, or trying to "heal" lost facts with teacher answers.
    • mistake: treated REAP-320's 25% fabrication rate on obscure facts as a reason to keep more experts, and first thought of standard teacher distillation to fix hallucinations.
    • trap ruled out: naive distillation on facts the pruned experts lost makes hallucination worse — it trains the student to give confident teacher-style answers it can no longer back up [src: 2405.05904, 2403.05612, 2509.04664].
    • decision: accept world-knowledge loss (the agent has web search), and do abstention-aware healing (R-Tuning style) — keep teacher answers only where the pruned model still knows them across multiple samples, and rewrite lost/unstable facts to "i don't know / i'm not sure / i haven't checked" in dawn's voice or a direct search call. measure both fabrication and over-refusal.
  6. thinking no teacher fits on one 80 GB Colab A100.
    • mistake: assumed teacher distillation needed the ~360 GB BF16 Swift checkpoint, which can't sit next to the student on one GPU.
    • decision: use unpruned Swift GSQ-RCO IQ3_S as an offline teacher via llama.cpp. its resident shard is 41.77 GB (fits in 80 GB VRAM with the 28.8 GB PLE shard on disk), it matches BF16 on benchmarks, and dumping top-k logits offline means teacher and student never share VRAM.
  7. assuming phase 3 training had to stream 512 experts from host RAM.
    • mistake: planned phase 3 around streaming ~125B experts over PCIe from 167 GB host RAM on Colab.
    • realization: training happens after pruning. at 224 experts, even int8 experts (~52.8 GB) + BF16 dense (~7.3 GB) fit in one 80 GB A100 with no streaming; if we dequantize the frozen Q2_0 GGUF experts directly in the forward pass (~16.5 GB), it fits with room to spare and gives free QAT against the deployment quant.
  8. dismissing small runtime wins ("prefix sliding won't win much, drop it") instead of squeezing.
    • mistake: wanted to throw out prefix sliding outright because 36/48 layers are GDN, the 12 QSA layers already cap KV at 2048 tokens, and Swift cuts thinking length by 63%.
    • correction (dawn): we're squeezing — anything that measures net-positive on speed without hurting evals stays in.
    • decision: rule out baking prefix sliding into training (hurts long-range agent recall and cache reuse), keep it as an untrained runtime toggle (default off) on the QSA indexer, and A/B test it.
  9. overestimating cache eviction policy, and considering instantaneous inverse-router eviction.
    • what we thought: maybe evicting the router's bottom-ranked experts per token beats recency/frequency by a lot.
    • what ds4-recipe doc 06 proved: eviction has a low online ceiling. Belady's oracle gets −56.3% misses because 1–128 token reuse is invisible at eviction time; across 12 production tapes, the best online score beat production by −4.0% (and pure LRU/LFU was +5…+46% worse), whereas MoCE cache-aware routing cut misses by −52…−73%.
    • decision: rule out instantaneous per-token bottom-k eviction (too noisy); test popularity EMA + near-pick credit / inverse-router EMA in the trace replayer for the ~few % it can legitimately win, and put the real effort into MoCE.
  10. evicting next-layer prefetch.
    • trap ruled out: prefetching predicted next-layer experts straight into the main cache lost 8.6% live in ds4 because wrong guesses evicted needed experts (and under MoCE, changing residency changes routing).
    • decision: only non-evicting staged prefetch (confident misses into a separate pinned shelf, adopted only on a real demand miss).
  11. framing the model as a SWE-style coding agent instead of a persistent persona agent.
    • mistake: initially designed the calibration mix and eval around repo coding (SWE-bench Verified).
    • correction (dawn): it's a long-lived persistent persona agent; its only "coding" is python-tool use (like albedo's kernel) and linux ops.
    • decision: drop repo coding from calibration and eval (freeing expert budget for persona, JA, ops, memory, vision) and optimize runtime for many small turns + regular compaction + long idle gaps.
  12. worrying that weight-baked persona hurts instruction following, vs the voice-override trap.
    • what we thought: Swift dropped IFBench 73.2 → 70.1, so baking a persona into weights might hurt instruction adherence further.
    • why baking it in actually helps: taking persona rules out of the system prompt avoids multi-turn attention drift (~8 rounds [src: 2402.10962]) and lowers simultaneous instruction count so memories, self-edited prompt rules, and compaction summaries get followed better [src: 2507.11538, IFScale].
    • the real trap to guard against: if every training trace is casual persona text, the model learns voice outranks format/register constraints (kaomoji leaking into JSON tool calls, config files, or messages written for others). fix: include explicit voice-override examples in training and test context-instruction following in the eval.
  13. treating Connectome folds and LCM compaction as different SCR problems, or re-rotating keys without recomputing values at junctions.
    • mistake: initially ignored Connectome, then briefly treated Connectome deep folds and LCM summary compaction as different architectures.
    • decision: both have the exact same [head + folded summaries + verbatim tail] shape and hit the exact same SCR path ([A B C] → [A B' C]). what changes local KV cost is just config: tail size, layout stability, and fold frequency.
    • KV trap from Anima's experiments: reusing trimmed cached values under re-rotated RoPE keys caused degenerate output loops ("worked worked worked") unless values/junctions were recomputed and the attention-sink head was pinned. added to our SCR test (§6.1.7).
  14. trusting F_NOCACHE=1 on macOS APFS without checking iostat -Id (disk0).
    • trap caught during bench/run_02: repeating pread over identical file offsets across benchmark reps on APFS can partially hit cache on rep 2+ even when the descriptor has F_NOCACHE=1 (or when identical buffer contents get deduplicated).
    • rule: for any local SSD benchmark on macOS, write unique incompressible blocks, read virgin slices per test, and verify iostat -Id (disk0 physical MiB delta == 100.0% of requested MiB).
  15. scraping blanket transcripts across all models for persona training / calibration.
    • clarified (dawn): don't dump arbitrary session transcripts into calibration or training. for phase 3 voice examples from albedo, filter only to gemini 3.8 flash and gemini 4 argon turns (plus APPEND_SYSTEM.md), while all assistant transcripts can be used to mine LLM-isms (mine_llm_isms_from_transcripts.py).
  16. hand-writing a tiny synthetic calibration file (~5K tokens) instead of assembling ~400K–525K tokens from real HF datasets.
    • mistake caught by dawn: wrote 8 synthetic conversations (~21 KB) for data/calibration_corpus.txt before checking HF. on a 512-expert MoE, a 5K-token corpus leaves many experts untouched or under-sampled; REAP-256-duo specifically warns that subsampling a corpus creates fake "domain specialization" and inflates apparent domain divergence by ~11× [src: REAP-256-duo card].
    • rule: always assemble ~400K–525K tokens from real open datasets (ukisai/Qwen3.8-27B-multi-turn-agent-sft, Lottolabs/terminal-bench-2.1-qwen3.8-27b-traces, logic65/whittle-distill-qwen3.8-27b, ChiTako/japanese-imatrix-calibration, hotchpotch/japanese-qa-reasoning-100k).
  17. decaying EMA scores only on resident experts (opus-auditor catch in replay_moe_cache.py).
    • bug caught by tangled-clod/claude-opus-5-5: near_pick_ema and inverse_router_ema initially decayed ema_score only for c in cache_set while adding credit to non-resident experts in top_n, letting non-resident experts pile up undecayed sums and stay pinned for ~99 tokens once admitted. fixed by decaying all n_experts every step.
  18. six Phase 1 calibration/eval traps caught by opus-methodology (audits/opus_audit_phase1_methodology.md):
    • missing --parse-special in llama-imatrix: without --parse-special, <|im_start|>, <|im_end|>, <think>, </think>, <tool_call> are tokenized as literal BPE pieces instead of the special tokens the router sees at termination and tool calls.
    • sequential concatenation before --chunks cutoff + char vs token balancing: concatenating all 20 Terminal-Bench docs before ukisai meant a 65K-token prefix gave ukisai and hotchpotch 0 tokens, and equal character counts gave Japanese ~2–3× more tokens than English. fixed by document-level round-robin across all 5 sources with per-round token budgeting (~1,600 JA chars ≈ 1,000 tokens matched to ~3,600 EN/code chars ≈ 1,000 tokens) and -c 2048.
    • task/shingle leakage across splits: Terminal-Bench (low/ vs medium/), ukisai (_run2.._run5 / __dsv4-trial-2 suffixes + the shared 2,834-char Terminus-2 preamble), and ChiTako (18.9% duplicate paragraphs) all leaked across row-level splits. fixed by group-splitting on base task ID / URL, stripping the repeated Terminus-2 preamble, using xhigh/.../prompt.txt, and verifying 20-char shingle disjointness.
    • Whittle double-templating: Whittle's prompt column already contains <|im_start|>system ... <|im_start|>assistant\n<think>\n; wrapping it again produced nested templates. fixed to use prompt + response directly and truncate long prompts from the left after the system header.
    • ranking by reap_l2_sum (count × REAP) alone: frequency-weighted sum prunes rare-token specialists first. added published REAP mean with empirical-Bayes shrinkage + reap_hybrid + offline top-10 substitution error scoring from the Q4RT trace.
  19. three Phase 1a v2 / 1b bugs caught by opus-review-1a-1b (audits/opus_audit_phase1a_v2.md):
    • count-weighted shrinkage prior (mu_layer = Σreap / Σcount) keeping dead experts: because heavy-hitter experts have higher gated norms, the count-weighted mean sits above 63–85% of per-expert means, scoring all 8 zero-hit experts (count == 0) in the top 15–37%! fixed in scripts/select_and_eval_prunes.py v3 to use a zero-prior damped LCB ((c / (c + m)) * max(0, mean - SE), scoring count == 0 strictly at 0.0 with a fail-closed assert zero_hit_kept == 0).
    • outlier max-normalization in reap_hybrid: a single extreme expert in Layer 47 (e143 mean 5.56 vs layer median 0.31) squashed max(shrink) normalization; fixed reap_hybrid to use outlier-proof rank-percentile averaging (0.5 * pct(reap_l2_sum) + 0.5 * pct(reap_shrink)).
    • undercounting misses in pruned trace replay when < k_used candidates survive in top-64: fixed scripts/replay_moe_cache.py and scripts/replay_moe_cache_fast.cpp to count max(0, k_used - len(chosen)) as forced misses, added deterministic lowest-ID tie-breaking, and added out-of-sample held-out trace verification (eval_trace_heldout.q4rt).
  20. matching strncmp(t->name, "ffn_moe_weighted-", 17) without checking t->op == GGML_OP_MUL (patch_imatrix.py):
    • bug: in llama.cpp (src/llama-graph.cpp), ggml_view_2d(ctx0, experts, ...) automatically names the 10 per-expert view slices "ffn_moe_weighted-{il} (view)" (GGML_OP_VIEW), which also starts with the 17 chars "ffn_moe_weighted-", tripping a fail-closed GGML_ASSERT(down_t->op == GGML_OP_MUL_MAT_ID) on the first ubatch. fixed by checking t->op == GGML_OP_MUL (and t->op == GGML_OP_MUL_MAT for ffn_moe_logits-).
  21. re-parsing Q4RT binary traces and running 120-config cache sweeps in pure Python:
    • mistake: unpacking 786,432 structs per config (run_phase1b_sweep.py) or building ~100M Python tuples + running min(candidates, key=lambda...) across 120 configs × 48 layers × 16,384 tokens timed out at 300s twice. fixed by writing scripts/replay_moe_cache_fast.cpp (g++ -O3 -fopenmp), cutting the 120-config 48-layer sweep to 24–47s.
  22. llama-perplexity special-token (parse_special = false) and --save-all-logits CLI traps:
    • trap 1: unlike llama-imatrix, llama-perplexity rejects --parse-special on the CLI AND calls common_tokenize(ctx, params.prompt, true) with parse_special = false by default in tools/perplexity/perplexity.cpp (lines 310, 473). always patch perplexity.cpp to common_tokenize(ctx, params.prompt, true, true) before evaluating chat/agent traces.
    • trap 2: --save-all-logits FNAME and --kl-divergence-base FNAME are aliases for the same params.logits_file option; passing both on the baseline run treats --kl-divergence-base as the filename. always assert p.returncode == 0 in eval wrappers so bad CLI flags fail closed immediately instead of recording -1.0.
  23. slicing 60 GB candidate GGUFs onto Colab /content overlayfs (~300 MB/s) instead of /dev/shm RAM tmpfs (83 GB):
    • mistake: writing each 64.8 GB K320 candidate to /content/pruned took 216.8s (3.6 min) of network-disk writes per candidate, and Colab reclaimed the 3-hour A100 VM mid-slice on candidate #2 (K320_reap_hybrid) right after K512_unpruned and K320_reap_l2_sum completed (when a Colab VM vanishes mid-job, job.poll() remains None until colab.instances() is checked). fixed by slicing temporary candidates into /dev/shm/pruned (83 GB RAM tmpfs on the 167 GB High-RAM A100) and saving results/phase1c_kld_partial.json locally.
  24. running Tier-2 generation evals at max_tokens=640 and temperature=0.0 on a thinking model (run_phase1c_tier2_eval.py):
    • mistake: with --reasoning-format deepseek, any prompt whose <think>...</think> trace exceeds 640 tokens leaves message.content == "" (confounding normal 700–1,200 token technical JA explanations on K512_unpruned with genuine <think> stall loops), and naive 5x 4-gram counting flags common English prepositional phrases ("in the kv cache"). fixed by setting max_tokens=2048 (temp=0.6, top_p=0.95, top_k=20, seed=42), explicitly tracking empty_content_stall + completion_tokens, and checking consecutive 20-char repeats + 8x 5-grams.
  25. Phase 1c KLD & Tier-2 rubric traps caught by opus-review-1c (audits/opus_audit_phase1c.md):
    • delta_ppl clamped-base bias & JA PPL non-monotonicity: --save-all-logits quantizes base log-probs to uint16 over [max_logit - 16, max_logit], so PPL(base) in KLD mode sits 0.6% / 4.0% / 6.5% below exact FP32 K512.ppl, and pruned models have up to 42% lower JA PPL (K256_shrink gets 16.2 PPL with the worst 0.520 KLD) because router renormalization shifts probability mass toward calibration domains. rule: rank strictly on KLD and Top-1 agreement, never quote JA PPL as a quality gain, and record exact_delta_ppl = ppl_q - base_fp32_ppl.
    • Terminal-Bench execution-contract footer in eval_ops_tools.txt: Lottolabs/terminal-bench-2.1 prompt.txt files end in a shared ~1.1k-char Execution contract: / Benchmark rules: block. scripts/assemble_corpus.py now strips \n\nExecution contract: on tb_eval_docs and places ukisai multi-turn traces first (while keeping tb_calib_docs unchanged so calibration_mix_v2.txt reproduces Phase 1a bit-identically); note that the Phase 1c ops_tools KLD numbers in results/phase1c_kld_results.json were measured on the pre-dedupe 4-chunk window.
    • Tier-2 rubric false negatives & truncation accounting: checking case-sensitive "WAL" when PERSONA_SYS mandates lowercase failed journal_mode=wal, and requiring katakana エキスパート/チャンク failed valid expert/chunk loanwords in Japanese engineering prose. results/phase1c_tier2_results.json (rubric_version: "v1_run_rescored_ja_trunc_only") is the 2048-token cap v1 run rescored with case-insensitive JA alias groups + curly-apostrophe abstention markers + completion_tokens >= 2048 truncation failure (4 of the pruned failures are 2048-cap truncations, and ops_1/ops_3 keep their v1 verdicts since v1 did not save raw tool args); scripts/run_phase1c_tier2_eval.py is updated to max_tokens=4096 with full raw output logging for Phase 3/4 re-runs.
  26. slop-miner traps caught by opus-review-slop-miner (audits/opus_audit_slop_miner_and_1c_v3.md, audits/opus_audit_slop_miner_v4.md):
    • Monroe pooled-background cancellation at c_hum == 0: any p_pooled = (c_llm + c_hum) / total_bg term in p_bg(w) makes α_w proportional to c_llm at c_hum == 0 and creates an artificial ceiling (δ ≈ log(1 + N / (w_pool * α_0))). Fixed by dropping the self-pooled term completely and using a pure per-order unigram-product + per-order uniform prior normalized over the candidate vocabulary (p_bg = 0.85 * p_indep_per_order + 0.15 * p_unif_per_order, with 0.0 self-pooled weight and sum_w α_w == α_0), so c_llm never enters α_w and all 23 c_hum == 0 rows have distinct, non-ceiling log-odds deltas ([0.626, 2.814]).
    • phantom n-grams across stripped spans & hyphenated identifiers: replacing inline code/URLs/paths with spaces (" ") or skipping digits/hyphens inside WORD_RE glued words across removed spans ("run \make check` then commit"->"run then commit", "hyd-eu-1"->"hyd eu", "break-check"->"break check"). Fixed by replacing every stripped span with a hard clause boundary ("\n") and splitting clauses on any non-[a-z'\s]` run.
    • curly apostrophes, super-gram stop-list leaks, mid-phrase "I" casing, (let pretokenizer merges, and cross-family floor: normalized \u2018/\u2019 -> "'" before tokenization, checked sub-tuples against PROTECTED_STOP_NGRAMS (including "i don't know", "i'm not sure", "haven't checked"), added mid-phrase "I" and (/[-prefixed variants in Qwen token mapping, required >= 2 model families each with >= max(5, 5% of c_llm) hits, read root-session transcript user rows while stripping > quote lines and skipping isSidechain subagent rows, and used the miner as a confirmation filter over curated rhetorical/signpost candidate lists (RHETORICAL_ALLOWLIST + SIGNPOST_PREFIXES minus EXCLUDED_WORK_AND_GENERIC_WORDS, dropping "think that's", "so far", and "as expected" from FTPO).
  27. applying single-token FTPO to structural speech acts ("do you want", "want me to", "let me check") instead of sentence-level SFT surgery (dawn):
    • trap: if the model's hidden state already intends to ask a servile closing permission question or narrate a tool call before emitting <tool_call>, suppressing the surface n-gram "do you want" or "let me check" via single-token FTPO rollback only forces synonym whack-a-mole ("would you like me to", "shall i", "allow me to inspect") and collateral-damages ordinary uses of "want" and "let".
    • rule: partition mined LLM-isms (results/banned_llm_isms.json) into phase3_sentence_surgery (25 pre-tool announcements, closing permission-solicitations, sycophantic openers — removed at the sentence/clause level in Phase 3 SFT data via scripts/prepare_phase3_sft_data.py) vs phase4_lexical_ftpo (7 surface rhetorical cliches -> 126 Qwen token tuples: "rather than", "right away", "end to end", "exactly what", "here's what", "what's actually", "short version" — safe for single-token FTPO in Phase 4).
  28. four Phase 3 prep blockers caught by opus-review-phase3-prep (audits/opus_audit_phase3_prep.md):
    • Q2_0 codec vs ggml-quants.c (B1): ggml-quants.c (quantize_row_q2_0_ref L74, dequantize_row_q2_0 L439) uses scale d = amax, zero-point 1 ((q - 1) * d -> {-d, 0, +d, +2d}), and packs 4 consecutive elements per byte (qs[j // 4] >> ((j % 4) * 2)), NOT (q - 2) * d with quarter-strided layout. Fixed in scripts/phase3_lora_and_transplant.py and verified in --self-test against a compiled C reference transcribed from ggml-quants.c (max_err == 0.0 across 4,096 floats + full 4-code GSQ blocks).
    • in-place span surgery vs " ".join(units) (B2) & kaomoji symbol preservation (B3): splitting replies on \n+ and rejoining with " " wiped newlines, lists, tables, and code fences (79% of Gemini reference turns have newlines, 21% have fences), and including U+2600..U+27BF in UNICODE_EMOJI_RE stripped kaomoji symbols (✧ ♡ ♥ ✿ ☆ ♪ in (◕‿◕✿)). Fixed in scripts/prepare_phase3_sft_data.py by performing in-place span deletion (absorbing only whitespace directly attached to removed emoji so 0/140 fenced Gemini reference turns and padded tables/kaomoji are altered) and narrowing UNICODE_EMOJI_RE to U+1F1E6..U+1FAFF + ✅❌⚠️✨⭐⏰⌛.
    • closing solicitation / opener / lowercasing precision (M1–M4) & R-Tuning / KD logit order (M5–M7): restricted closing solicitation stripping to genuine offers/permission questions (CLOSING_SOLICITATION_OFFER_RE, never stripping technical conditionals like "If you want lower prefill latency, set --ubatch 4096."), preserved hypotheses in compound pre-tool sentences ("Let me check X, since Y" -> keeps "Y"), blocked opener stripping before contrastive "but/except/though", protected README.md/Makefile/HOME/K=320/MemoryHigh during lowercasing, required word-boundary matching + template pool in classify_rtuning_abstention, and mandated running speech-act surgery before teacher-forced top-k logit extraction.
  29. four macOS Unified Memory / APFS bottlenecks in upstream llama.cpp (-cmoe) (scripts/patch_m5_lazy_moe.py, audits/opus_audit_phase0_2_m5_v3.md, audits/opus_audit_phase0_2_m5_v4.md):
  • --repack (CPU_AARCH64 weight repacking) on by default: src/llama-model.cpp (use_mmap_buffer) allocates an anonymous heap buffer in RAM to repack --cpu-moe expert tensors at load time, bypassing mmap and tripping the 64.0 MiB swapout watchdog before adding --no-repack / -nr ([unarchived pre-baseline smoke run]).
  • POSIX_MADV_WILLNEED on --cpu-moe experts + APFS sequential read-ahead (F_RDAhead) + QD=1 serial page faults: patched in scripts/patch_m5_lazy_moe.py. On K256_reap_l2_sum, unpatched llama-cli (chat mode on data/m5_val_prompts/llama_prompt_1264tok.txt with -n 64, ~64 generated tokens / 63 decode steps) ran at 1.40 tok/s decode and 14.30 / 21.70 tok/s prefill (-ub 1024 / -ub 2048, 105,851.73 / 85,725.73 MB whole-run disk0 read across load + prefill + decode), whereas patched llama-completion -no-cnv -t 8 decoded the 22-token prompt at 3.22 tok/s (33,428.34 MB whole-run disk0 across load + 22 prompt + 63 decode tokens) and prefilled the 5,600-byte (1,264-token) raw prompt at 38.11 / 63.27 tok/s (-ub 1024 / -ub 2048, hitting EOS immediately after prefill with 0 decode tokens so 46,808.58 / 28,272.94 MB measure load + prefill only; note the two runs also differ in llama-cli chat vs llama-completion -t 8 raw mode). --moe-cache-mib 1536 degrades decode to 0.28 tok/s, and llama.cpp -cmoe remains CPU-compute-bound because ggml-cpu computes expert matmuls on the CPU instead of on the M5 GPU.
  1. five 24 GB Mac / pruned-MoE bottlenecks in upstream Niko1221/Strata (scripts/patch_strata_metal.py, scripts/prepare_strata_k256.py, audits/opus_audit_phase0_2_m5.md, audits/opus_audit_phase0_2_m5_v2.md, audits/opus_audit_phase0_2_m5_v3.md, audits/opus_audit_phase0_2_m5_v4.md):
  • hardcoded NEXP = 512 across C++ (window.cpp, stream.cpp, prompt_chain.inc) and Metal shaders (router.metal, moe.metal, prompt.metal): patched to read n_expert dynamically (K = 256 / 320 / 224 / 512), run the prompt router GEMM over n_exp_ rows (ulong(t) * ulong(n_exp) stride in route_topk), and mask idx >= n_expert to -INFINITY (route_parity: 0 mismatches across 200 windows at K=512, pad=0 (2.16e-07 max rel err default math / 2.15e-07 fast math) and 200 windows at K=256 + MoCE=0.85 with random smap and 256..511 traps (3.06e-07 max rel err default math / 4.00e-07 fast math)).
  • ExpertStream::load_dense allocating 3.24 GB of anonymous RAM (cx.alloc): upstream load_dense already skipped the 1.28 GB BF16 hc_* originals and copied the remaining 3.24 GB (2.90 GB GPU dense + 0.34 GB CPU token_embd) into anonymous RAM. Patched load_dense to wrap <model>-metal.gguf (4.53 GB file map) zero-copy and mlock only the 2.90 GB of active GPU dense tensors while leaving 0.34 GB CPU-only token_embd.weight and 1.28 GB skipped BF16 hc_* unpinned via MADV_DONTNEED (~0.34 GB real resident RAM saving vs upstream's 3.24 GB anonymous copy).
  • PromptGemm (hold_region) carving prompt buffers out of the expert arena and MADV_DONTNEED not unmapping Darwin anonymous scratch: patched ExpertStream::region + Window::prompt_gemm_need_bytes() (verified equal to need at runtime in prompt_gemm_init) to map a separate 1.36 GB MAP_ANON buffer on demand without memsetting fresh zeroed pages, refresh slot.hot protection on the top k.n - (ne_ + 32) LRU slots on every prompt turn (in_prompt_), and immediately cx_->free + munmap the 1.36 GB scratch in prompt_region(false) after prefill (phys_footprint 4.73 -> 4.68 GB, resident RSS 7.41 -> 7.36 GB on the 1,024-token run; [hypothesis: only the ~0.052 GB CPU-faulted p.R+p.emb slice is faulted into host page tables before munmap]).
  • mixed IQ3_S/IQ4_NL expert slices (23.42 GiB, 3 fragmented size classes) vs uniform Q2_0 (15.82 GiB, 1 unified size class of 85 pages): built multi-threaded C++ converter convert_k256_q2_0 (vendor/strata/metal/src/convert_k256_q2_0.cpp), shrinking K256 experts from 23.42 GiB (1,550 slots / 1,440 filled across 3 classes in a 3.40 GB arena) to 15.82 GiB (2,441 slots in 1 unified class in the same 3.40 GB arena).
  • race-free MoCE cache-prior reranking (STRATA_METAL_MOCE_BONUS=0.85) in moe_route / router_route + ExpertStream::ahead & 3-domain complete-turn assistant quality measurement: wired smap residency into moe.metal / router.metal using a per-lane 16-bit register mask (scaled_mask + simd_broadcast, reading smap once at rank == 3 so background I/O threads cannot race with true_prob restoration) and updated ahead() to protect top-3 unperturbed candidates (ranks 0..2) and read smap. At a uniform 3.40 GB arena on the 43-token SQLite WAL prompt, MOCE=0.85 cuts misses by 51.5% vs Q2_0 (MOCE=0.0) (84.3 -> 40.9 misses/tok, 304.5 -> 127.2 MB/tok, 82.44% -> 91.48% hit rate) and 62.4% vs IQ3_S (MOCE=0.0) (108.9 -> 40.9 misses/tok), reaching 320.58 / 353.49 tok/s prefill (3,194.19 / 5,793.65 ms) and 9.51 -> 13.52 tok/s 43-token decode (16.11 ms wait + 56.29 ms GPU) in the earlier v3 pass (audits/opus_audit_phase0_2_m5_v3.md) and 229.99 / 268.23 tok/s prefill and 8.04–9.33 tok/s decode in the later results/phase0_phase2_m5_validation.json pass ([hypothesis: passive SoC thermal throttling + concurrent background CPU load]). Across 192 teacher-forced assistant tokens (n = 64 final assistant tokens × 1 complete untruncated turn/domain after n_prompt = 244 / 159 / 456 prompt tokens in data/m5_val_prompts/eval_tf_{ops_tools,chat_json,ja}.txt), isolated MOCE=0.85 vs Q2_0 (MOCE=0.0) adds only 0.0588 macro mean KLD (96.88% top-1 agreement: 100.0% on chat_json, 98.44% on ja, 92.19% on ops_tools), while unhealed IQ3_S/IQ4_NL -> Q2_0 expert conversion costs 0.2032 macro mean KLD (90.63% top-1: 100.0% on chat_json, 92.19% on ja, 79.69% on ops_tools; 0.2386 combined with MOCE=0.85), motivating Phase 3's LoRA healing against the deployed Q2_0 experts.
  1. Phase 3 local healing & transplant + multi-prompt generalization audit (scripts/heal_k256_q2_0.cpp, scripts/run_phase3_local_healing.py, results/phase3_local_healing_validation.json, results/phase3_22win_kld_audit_v1.json, results/phase3_22win_kld_audit_v2.json, results/phase3_9prompt_speed_audit_v1.json, results/phase3_9prompt_speed_audit_v2.json, audits/opus_audit_phase3_local_v1.md, audits/opus_audit_phase3_local_v2.md):
  • Signed-d importance-weighted Lloyd-Max Q2_0 expert quantization (heal_k256_q2_0 --quantize): GGML_TYPE_Q2_0 (type 42, 18 bytes per 64 weights) stores an IEEE fp16 scale d whose sign bit is unused by standard quantizers (d >= 0 -> {-d, 0, +d, +2d}), while all Strata Metal Q2_0 kernels (units.metal, moe.metal, moe_mm.metal, mma.metal) multiply (code - 1) * d linearly in signed half d. Unlocking d in R gives every 64-weight block a 1-bit polarity switch between {-|d|, 0, +|d|, +2|d|} (d > 0) and {-2|d|, -|d|, 0, +|d|} (d < 0) with 0 extra bits and 0 shader changes. audits/opus_audit_phase3_local_v1.md caught that weighting W_gate/W_up by blk.l.hc_ffn_norm.weight had -0.03 mean log-correlation with the actual expert input E[mixed_d^2] (because gemm_hc_up_64_64_4 mixes the 4 hyper-connection streams after hc_ffn_norm) and W_down had col_imp = nullptr. Upgraded --quantize (v2) to weight W_gate/W_up by tempered real activation RMS (E[mixed_d^2])^0.5 (beta = 0.50) and W_down by per-expert analytical fourth-root importance (sigma_g * sigma_u)^0.5 ~ (E[h_c^2])^0.25 (beta = 0.25 on W_down vs beta = 0.50 on W_gate/W_up) over the 1,022-token multi-domain calibration trace (data/m5_val_prompts/cal_disjoint_1024tok.txt, verified 0 16-gram overlap on scored spans across all 22 unseen windows and 3 canary windows; 18/22 unseen windows also have 0 full-window context overlap, while 4 windows share prompt context: <tools> boilerplate in cj20_tool_5turn_t3/cj20_t3_answer_tool/cj03_toolcall and another turn of persona conversation 3 in persona3_t4).
  • Real-trace per-layer alpha_l scale calibration + shared-expert ridge solver (heal_k256_q2_0 --calibrate-and-heal) & byte-level GGUF transplant: audits/opus_audit_phase3_local_v1.md showed that feeding synthetic Gaussian noise x = z * s into Part A overshot alpha_l on 22/48 layers (especially deep layers L24–L44 where real activations want 0.97–1.04 vs Gaussian 1.08–1.12). Upgraded --calibrate-and-heal (v2) to compute exact parameter-free real-trace MSE scales alpha_l = <y_ref, y_q2> / ||y_q2||_2^2 (gamma = 0.00, L0 = 1.308, L24 = 1.022, L47 = 1.121 vs unseen held-out oracle L0 = 1.337, L24 = 1.021, L47 = 1.134, reducing layers where alpha_l is worse than 1.0 on unseen data from 22/48 to 2/48 and unseen routed residual from 24.13% to 23.37%), solve ridge regression (46/48 layers full-rank 640/640; L0 rank 430, L2 rank 617; eta_shd = 0.15, ridge_rel = 6.0) on blk.0..47.ffn_down_shexp.weight (in-sample MoE MSE 19,061.40 -> 18,665.36 -> 18,460.89, -3.15% total: -2.08% from alpha_l + -1.10% from ridge after IQ4_NL requant; -0.9% LOBO CV, -0.46% on unseen dump), and transplant only the 48 modified ffn_down_shexp tensors into K256-q2_0-healed-metal.gguf (1175 / 1223 other tensors bit-identical), with 4-uint64 FNV-1a cache fingerprinting (yr_q2_bin.meta over d_bak + L0/L47 .mixed + L0/L24/L47 gate/up + t_cal), .quant.meta, and strict pread_exact/pwrite_exact I/O checks.
  • Deterministic smap_ synchronization in stream.cpp + full 22-window (2,686-token) and 9-prompt (v1 & v2) generalization audits (audits/opus_audit_phase3_local_v2.md): eliminated the background read-ahead CPU/GPU data race on smap_ in stream.cpp (publishing smap_ only when !x.pre or at claim time inside at_layer(), and evicting unchosen speculative read-aheads when moce_bonus_ > 0.0f). Across the full 22-window (2686-token) 4-stratum unseen KLD suite (results/phase3_22win_kld_audit_v2.json), v2 healed Q2_0 beats unhealed Q2_0 by -28.55% [95% CI -33.5%, -23.2%] pooled KLD at MOCE=0.0 (0.2029 -> 0.1450, 20/22 wins, +2.27 pt top-1: 85.22% -> 87.49%) and -26.98% [-32.0%, -22.1%] at MOCE=0.85 (0.2123 -> 0.1550, 21/22 wins, +1.53 pt top-1: 85.55% -> 87.08%). Comparing v2 vs v1 across all 22 windows gives -5.08% [-10.8%, -0.1%] at MOCE=0.0 (-3.26% at MOCE=0.85, -3.51% [-8.8%, +1.5%] n.s. on the 18 context-clean windows) and -16.53% [-28.3%, -4.4%] on isolated MoCE KLD (0.0389 -> 0.0325), with ops (0.1199 -> 0.1077) and chat_json (0.1243 -> 0.1130) improving, persona flat (0.2637 -> 0.2527), ja slightly higher vs v1 (0.1723 -> 0.1796, top-1 85.0% -> 83.7%, still -24.3% vs unhealed 0.2374), and cj20_tool_5turn_t3 moving from 0.1783 (v1) to 0.1444 (tied with unhealed 0.1440 at MOCE=0.0, 0.1417 vs 0.1534 at MOCE=0.85, aided by in-distribution <tools> context). On the 192-token 3-domain canary set, v2 macro KLD is 0.1563 at MOCE=0.0 (-23.08% vs unhealed 0.2032, worse than v1's 0.1320) and 0.1559 at MOCE=0.85 (-28.57% vs unhealed 0.2182, vs v1's 0.1365, with isolated MoCE -57.30%: 0.0616 -> 0.0263). Across the 9-prompt diverse speed suite on v2 (results/phase3_9prompt_speed_audit_v2.json), MOCE=0.85 is 9/9 bit-deterministic across reps with 92.96% mean combined hit rate (90.88%–95.60%, 33.8 mean misses/tok vs 85.12% at MOCE=0.0; v1 was 93.27%, with persona_458 moving 93.04% -> 91.26%), and 0.0 MiB swapouts across all runs.
  1. Phase 4 (#97) WFLD Layer-47 anti-slop + Strata Metal v3 (#98) runtime optimizations (scripts/solve_phase4_wfld_l47.cpp, scripts/run_phase4_colab_ftpo.py, results/phase4_healed_persona_ftpo_writers.bin, results/phase4_colab_validation.json, audits/opus_strata_metal_opts_v1.md, audits/opus_strata_metal_opts_v2.md, audits/opus_strata_metal_opts_v3.md):
  • 4-layer proxy (L40..47) intermediate-layer drift (v1) & unprefixed mined penalty (v2) + <|im_end|> special-token mass shift (v3.1) -> Special-Token-Orthogonal Per-Word-Null WFLD at Layer 47 (v3.2): v1 used a bag-of-n-grams proxy for L0..39 and summed L40..46 perturbations without passing them through the 7 intervening QSA/GDN/MoE layers (ops_tools KLD jumped to 0.276). v2 dumped exact post-Layer-47 states (R_post47, inj47, hh47, sgv47) from strata_window_cli, eliminating intermediate drift (47/48 tensors bit-identical to Phase 3, only updating blk.47.ffn_down_shexp.weight in Q8_0), but its mined_* loop penalized the 6 final words without checking tuple[:-1]. v3/v3.1 enforced strict prefix matching and per-word collateral covariance G_b, but audits/opus_strata_metal_opts_v3.md §6.1 caught that setting delta_cho = 0 without constraining <|im_end|> (248046) caused suppressed probability mass to shift into <|im_end|> (28 non-slop and 7/21 slop argmax flips to <|im_end|> on multi-turn chat). v3.2 solves a $15 \times 15$ output-space minimum-norm system in $\mathbb{R}^{2560}$ that boosts top-2 natural semantic alternatives $(w_{c_1}, w_{c_2})$ while enforcing $d_{\text{im_end}} \le 0$ and exact zero shift on <|im_start|>, <think>, </think>, <|endoftext|>, and common chat starters (||dW||/||W|| = 0.0978), passing all 7 bars (a)–(g) in §6.2 (0 flips to special tokens across all 4 traces, 56/56 real turn ends preserved, non-slop chat KLD = 0.00472 <= 0.005, ordinary Δlog p = +0.0208 >= -0.01, all 6 collateral words >= -0.042 with 13/142 = 9.15% <= 10.0% below -0.1, -69.7% / -70.0% / -84.1% slop reduction on dev separate-doc / same-doc / train, and 0.00319 3-domain real-engine macro KLD) and generalizing on a one-shot untouched 1,100-token blind chat test (§6.3: non-slop KLD = 0.0030, ordinary Δlog p = +0.012, 0 assistant-side turn damage, -54.4% weight-only slop reduction (0.554 -> 0.253) and -98.98% (0.554 -> 0.0056, 21/21 argmax clean) when stacked with the 0-overhead prefix-conditioned logit penalty (STRATA_ANTI_SLOP_PENALTY=6.0) in Window::finish).
  • Strata Metal v1–v3 runtime optimizations on 24 GB M5 Air (audits/opus_strata_metal_opts_v1..v3.md, scripts/strata_metal_complete.diff): (1) STRATA_METAL_MLOCK_DENSE=1 + STRATA_METAL_WIRE_ARENA=1 at 7.0 GB arena (0 pageins/decompressions, 0.0 MiB swapouts); (2) STRATA_METAL_SHEXP_OVERLAP=1 + STRATA_METAL_SPLIT_TURN=1 + STRATA_METAL_EXPERT_SPLIT=1 (resident-first 2-phase moe_gu/moe_down overlap with MTLSharedEvent SSD handshake, shrinking CB–kernel gap to < 1.0 ms and cutting decode from 58.72 ms/tok (17.03 tok/s) to 45.28 ms/tok (22.08 tok/s, +29.7%, 97.59% combined hit rate, 11.6 misses/tok, 5.08 ms SSD wait); 22.25 tok/s / 44.95 ms/tok with PF2=2, IO=12); (3) STRATA_METAL_PRED_Q8=1 (Q8_0 lookahead router saving 92 MB/tok DRAM reads), STRATA_METAL_ROUTER_BF16=1 (lossless BF16 router halving router reads 126 -> 63 MB/tok), STRATA_METAL_MOE_DN_TILE=8,2,4 (-0.25 ms/tok on M5), STRATA_METAL_PREWARM=1 (eliminates 282.8 ms T=3 speculative draft shader compile stall), and STRATA_HC_Q4 / STRATA_HC_Q5 (Q5_0 / Q4_0 hyper-connection GEMVs saving 240–315 MB wired RAM).
  1. Phase 5 Disk Session Save/Load (--save-session / --load-session), Suffix Cache Reuse (SCR, --scr-test) with 36-Layer GDN Staleness, and 190-Pair Real-VN Anti-Overfitting Audit (audits/opus_phase5_scr_and_persona_v1.md, audits/opus_phase5_scr_and_persona_v2.md, scripts/scr_state.hpp, scripts/window_session_scr.patch, results/phase5_scr_and_session_validation.json):
  • Disk Session Checkpoint (Window::session_save / session_load): Window::checkpoint_save only saved the 112.60 MiB running state (36 GDN gst_ + ghist_, pleHist_, tail_, dead_), omitting QSA kq/vq/ksc/vsc, pooled, and ss.ple_prev. session_save/load serializes all 118,069,248 fixed bytes + 14,208 B/cell (139.46 MiB at 1,983 cells) directly from/into MTLBuffer unified memory via 4-thread pread/pwrite with F_NOCACHE + fsync, loading in 28.93 ms cold NVMe (5.18 GB/s, 281x faster than 8,138 ms cold prefill) and 5.00 ms warm (1,627x faster) at 0.000000 KLD (100.00% top-1).
  • Suffix Cache Reuse (SCR) on 36-Layer GDN + 12-Layer QSA (Window::scr_relocate): in qsa.metal, only group 0 (dims 0..63, ksc[.][.][0]) of kq (int8) and pooled (f32) has NEOX RoPE applied (dims 64..255, ksc[1..3], vq, vsc move as exact bytes). Enforcing $\Delta \equiv 0 \pmod 4$ (padding B') relocates 1,535 tail tokens across all 12 QSA layers in 0.5 ms on CPU. Keeping the post-edit 36-layer GDN state (S_ABC, keep) gives +0.0038 nats/tok (C=511, where keep J=8 lowers PPL 7.49 -> 7.20) to +0.038 nats/tok (C=1535, 3.48x–3.79x prefill speedup) of isolated GDN staleness (keep - gdnex) and 3.20% repeated 4-grams (0% looping) over 128 greedy tokens, whereas dropping C's GDN history (lossy) explodes KLD to 1.41–2.33 (PPL 29.5–66.4). Across all 1,728 GDN value heads (36 x 48), median gate half-life is 18 tokens (76% decay within 100 tokens; 104 long-memory heads > 1,000 tokens).
  • Persona Overfitting & Template Audit (scripts/build_contrastive_persona_vn.py, scripts/pack_contrastive_persona_vn.py, data/real_vn_steins_gate_va11_120.json): (1) Ending <think> with </think>\n (198) instead of Qwen's </think>\n\n (271) caused 14/70 step0 special-token argmaxes; fixing </think>\n\n + the Window::finish step0 guard (p1 == 248069 && (p0 == 271 || p0 == 198)) reduced special-token step0 argmaxes to 0 / 190. (2) Full-row output.weight updates (2.9M–7.5M params) memorized training content nouns (train CE 2.83 -> 0.04, held-out CE 2.59 -> 2.82, blind chat KLD 0.075). Restricting to a rank-1 direction over 16 kaomoji/style tokens (2.6K params) on the 3-way split (113 train / 40 dev / 37 test, 17,352 whole-pair rows) achieved +0.196 nats/tok on test (blind chat KLD 0.0033, 0 flips), but test/train ratio was 0.38 < 0.70 because idx % 5 kaomoji assignment has $\ln 5 = 1.61$ nats of arbitrary glyph entropy that train fits and test cannot; per pre-registered anti-overfitting gate (>= 0.70), no head-row style edit is shipped.

#34. Unpruned K=512 ukisai GSQ-RCO IQ2_XS (scripts/build_k512_iq2_xs_metal.py, scripts/serve_k512_local.sh):

  • Rather than pruning 512 -> 256 experts and re-quantizing scalar Q2_0, streaming ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF's imatrix/GSQ-RCO vector-quantized IQ2_XS across all 512 experts (3.90 GB dense GGUF + 36.23 GB .experts, replacing the 3 IQ1_M layers blk.8,13,37 with IQ3_XXS's native IQ3_S/IQ2_XXS tensors and skipping the 28.80 GB duplicate PLE table) runs at 21.00 tok/s (47.61 ms/tok, 94.38% hit rate, 7.97 GB resident footprint, 0.0 MiB swap) on the 24 GB M5 Air with 100% of all 512 experts and ukisai's full imatrix quality; setting reasoning_effort|default('low') in tokenizer/chat_template.jinja keeps <think> concise (~50–100 tokens) without any hard token cutoff.
  1. Anti-Slop v4.1b (slop4 per-position stateful shaper, audits/opus_antislop_v4.md):

    • Layer B (linear WFLD on blk.47.ffn_down_shexp in IQ4_NL) was rejected by falsifier: ridge regression on 640-d z = silu(sgv)*hh leaked 0.28–0.68 (1–4 nats of bullet/# penalty) into code fences and <think> blocks.
    • Layer A (slop4::State, 12 B/pos in window.cpp, scanned via slop_scan on prompts/restores and relocated via scr_relocate) gates penalties strictly on PROSE && !TOOL && !FENCE (0 penalties and 0.000000 KLD across all 14,090 <think>, <tool_call>, and ``` code-fence rows; < 0.005 KLD on ordinary prose).
    • Empirical verification of the committed-prefix synonym treadmill vs clause-entry branch points (arXiv:2510.15061 Antislop, arXiv:2604.19139v2 Verbal Tics): penalizing at "If I'm being ___" leaves 64% of remaining mass inside candor siblings (real, totally, frank, blunt), whereas penalizing one token earlier at the branch point "If I'm ___" (being/honest) collapses meta-frame mass 0.161 -> 0.005 (-97%) so the model branches into content ("If I'm reading this right"). Includes slop4::hold() so penalizing a high-probability token never boosts <|im_end|>.
  2. M5 GPU Decode & Prefill Optimizations (audits/argon_k512_decode_opt.md, dense_mm.metal, prompt.metal, stream.cpp, window.cpp):

    • Decode (ntok=1): M5-tuned Window::gemv tiles (IQ4_XS output.weight/attn_qkv/attn_gate -> 32,2,4,1 at 130–138 GB/s, ssm_out/attn_out IQ4_XS -> 32,2,4,1 eliminating split-K barrier, IQ3_S 6144x2560 -> 4,2,4,1, shexp IQ4_XS 2560x640 -> 16,2,4,1 3.96x faster, pred Q8_0 -> 32,2,4,1), gr_down_q8_8_4 (158 GB/s), IQ2_XXS moe_gu gl=16, IQ3_S QUANT_CTAB=0 (118.4 GB/s), make_arena continue fix (3,726 -> 4,326 wired experts at 6.8 GB), and STRATA_METAL_PREFETCH_UNPERT=1 default (96.90% combined hit+prefetch rate, 21.00 -> 22.64 tok/s decode on unpruned K=512 IQ2_XS).
    • Prefill (prompt_chunk): Added gemm_f16_128_64_4 (-605 ms / -20% chunk GPU time), in_prompt_ single-piece pread + hold_region(false) pre_q_ drain + hot_cap = cls.n - (2*ne_ + 32) ping-pong, layer_rank_ hottest-first background prefetch (plus layer-0 kickoff in hold_region(true)), and STRATA_METAL_PROMPT_MOCE=320435.0 (Tail=32 exact + Body top-4 protected), cutting 244-token prefill from 12,493.9 ms -> 2,371.1 ms (5.27x faster, 95.31% top-1 agreement) and 1,024-token prefill from 7,642.9 ms -> 4,067.1 ms (+87.9% faster).
  3. Cross-Session Convergence with 01a11408 (regent@astaroth):

    • Ported 01a11408's include/strata/core/scr.hpp (scr_carry + automatic 0..3 newline pad insertion so any mid-context edit satisfies $\Delta \equiv 0 \pmod 4$ + <path>.pads persistence) into strata_window_cli.cpp --serve, and sent our chat_template.jinja default('low'), antislop_v4.inc, and layer_rank_ / PROMPT_MOCE findings back to 01a11408 via albedo send.
  4. Prefill Next-Layer Router Lookahead + Async PLE + PROMPT_CHUNK=2048 (audits/opus_prefill_opt.md):

    • In prompt_chunk (window.cpp), evaluating layer $l+1$'s router (dq(ffn_gate_inp) + gemm(p.mixh, p.Wr2, p.logits2) + route_topk) before layer $l$'s cut (~10 ms total GPU across 47 layers) and queueing predicted non-resident experts most-tokens-first before done_cv_.wait in stream.cpp cuts synchronous SSD misses on a 1,024-token chunk from 34.3% -> 9.5% (streamed-ahead 50.9% -> 75.5–79.0%, exposed SSD wait 1.4s -> 0.39–0.69s), while running ss.ple.table->gather_batch on a background thread joined at cut 0 hides ~150 ms of PLE SSD reads.
    • All 48 layers' STRATA_METAL_DUMP_MOE tensors (mixed, ids, wts, yr, hh, sgv, R_post47) are 100% byte-identical to baseline (0.000000 prefill KLD).
    • Setting STRATA_METAL_PROMPT_CHUNK=2048 (with STRATA_METAL_ARENA_GB pinned) cuts total experts read on a 2,048-token prompt by -43% (32.6K -> 18.3K) and boosts 100% bit-exact cold prefill to 335–346 tok/s (5,919 ms chunk for 2,048 tokens; 276.7–286.6 tok/s / 3,440–3,570 ms chunk on 1,024 tokens).
  5. Long-Context Decode Scaling Out to 64K (65,536) and 128K (131,072) on 24 GB M5 Air (K=512 Unpruned IQ2_XS):

    • Because 36/48 layers are GDN (112.2 MiB fixed recurrent state, 0 B/tok growth) and 12/48 layers are QSA (int8 KV + 128-dim 4-token pooled lightning indexer selecting top-2048 tokens), total session state is only 1.00 GiB at 64K (9.0 GB total resident RAM) and 1.84 GiB at 128K (10.3 GB total resident RAM), decoding at 21.13 tok/s (16K), 20.27 tok/s (32K), 20.35 tok/s (64K, 49.15 ms/tok), and 20.75 tok/s (128K, 48.20 ms/tok, 97.04% hit+prefetch rate).
  6. Three gpt-6.1-sol Audits (audits/sol_moe_cache_policy.md, audits/sol_spec_mtp_decode.md, audits/sol_persona_steering_arxiv.md):

    • Decode Prefetch (sol_moe_cache_policy.md): Verified io_main never publishes speculative pre slots into smap_, and tested STRATA_METAL_PREFETCH_UNPERT=0 vs 1 (call_2343482): UNPERT=0 halves read-ahead volume (144.5 -> 98.7 MB/tok) but increases synchronous misses (3.47% -> 5.18%, 22.43 -> 21.76 tok/s) because early cross-layer pred logits drift across layer $l+1$'s token mixer; kept UNPERT=1.
    • QSA-Boundary-Safe Suffix Speculative Decoding (sol_spec_mtp_decode.md, scripts/suffix_window.hpp): Caught an upstream qsa.metal (qsa_prep) rollback bug where tail[4][128] wraps modulo 4 across 4-cell block boundaries on partially rejected speculative windows (T > 1), corrupting pooled_ when remaining cells of the 4-cell block are re-run; capping speculative verification steps at T <= 4 - (pos & 3) (suffix_limit in suffix_window.hpp) eliminates all 736 unrestricted rollback failures with 0 overhead (20.62 -> 21.47 tok/s on cj20_tool_5turn_t3).
    • Post-Norm Persona Steering (sol_persona_steering_arxiv.md): Proved that any post-output_hc additive steering vector $h \leftarrow h + g \alpha v$ is mathematically equivalent to a gated vocabulary bias $g \alpha W_{\text{out}} v$, and verified on the 190-pair 113/40/37 split that ridge-whitened readout contrast fails the 0.70 generalization bar (dev/train = 0.514, test/train = 0.337); kept weights untouched.
  7. Native 49th-Layer Neural MTP (models/K512_iq2_xs/mtp-q2_0.gguf, 889 MB) + Unrestricted QSA tail_bak_ Rollback + [A B C] -> [A C D] Stripped-<think> SCR (scripts/build_mtp_q2_0.py, window.cpp, strata_window_cli.cpp):

    • Fetched all 31 mtp.* BF16 tensors (5.21 GB) from Qwen/Qwen3.8-Flash-Next via 24-worker HTTP Range requests and packed models/K512_iq2_xs/mtp-q2_0.gguf (889 MB, 10,629-token draft_vocab.bin, single-command-buffer mtp_chain.metal).
    • Added 24 KiB tail_bak_ snapshot/restore in Window::run / Window::commit (window.cpp) so T <= 4 speculative verification rolls back across 4-cell QSA block boundaries with zero restriction, and fixed Window::mtp_draft so --spec-min-p gates j = 0 as well as j >= 1 (Adaptive T in {1, 2, 3}).
    • 4-Domain Decode Benchmark (192 generated tokens, K=512 Unpruned IQ2_XS on 24 GB M5 Air):
      • Coding (C++20 lock-free SpscRing): 23.42 tok/s (T=1) -> 23.83 tok/s (SuffixDrafter) -> 28.49–29.44 tok/s (MTP + Suffix, +25.7%, 78.9–80.9% acceptance, 1.82–2.27 tok/round).
      • Reasoning (MoE N* crossover + GDN/QSA math): 22.38 tok/s (T=1) -> 25.04 tok/s (SuffixDrafter) -> 27.81–29.11 tok/s (MTP, +30.1%, 75.5–79.8% acceptance, 1.91–2.48 tok/round).
      • Agentic / Tool-Calling (cj20_tool_5turn_t3): 21.22 tok/s (T=1) -> 22.18 tok/s (SuffixDrafter) -> 27.76–27.95 tok/s (MTP + Suffix, +31.7%, 75.2–76.2% acceptance, 2.08–2.62 tok/round).
      • Persona Chat (lowercase nerd-girl + kaomoji): 21.03 tok/s (T=1) -> 21.21 tok/s (SuffixDrafter) -> 26.92–28.20 tok/s (MTP + Suffix, +28.0–34.1%, 1.72–2.10 tok/round).
    • Wired strata::core::scr_plan (min_span=32, seed=16, block=4) + bounded 2-slot MRU ck_ring into serve_loop so multi-turn chats that strip prior turns' <think>...</think> blocks ([A B C] -> [A C D], arXiv:2609.37725 Appendix B) automatically SCR-relocate the assistant reply C (cached_tokens: 414 / 446 = 92.8% reused on Turn 2!).
    • Slim MTP + AcceptMoE Draft Routing (router.metal, moe.metal, stream.cpp, window.cpp, call_2008636 & call_1884893):
      1. madvise(MADV_DONTNEED) on the 181 MB of source BF16 dense matrices in mtp-q2_0.gguf after quantizing them into mq8_ (90 MB Q8_0) + switching draft_vocab from 106,299 tokens (145 MB) to draft_vocab_en.bin (40,525 tokens, 55 MB) saves 271 MB of RAM, allowing STRATA_METAL_ARENA_GB=6.5 (9.9 GB resident RSS, -1.0 GB lower than baseline!).
      2. Capping mtp_encode's local QSA sliding window to 256 cells (4 chunks of 64 instead of 33) and running mtp_prefill only on the last 256 tokens of the final prompt chunk eliminates the multi-chunk mtp_prefill lag on 12K–32K prompts (314–330 tok/s end-to-end 2K prefill!).
      3. Setting STRATA_METAL_PREFETCH_DRAFT=0 in stream.cpp (ahead() only prefetches for root token t=0) and AcceptMoE (arXiv:2608.02989) draft routing in router.metal/moe.metal (sg > 0 protects r_prot = 1 and squares moce_scale) cuts speculative SSD traffic from 332 -> 177 MB/window (94.1% arena hit rate!) and pushes 32K-context decode to 30.37 tok/s (persona_chat), 33.44 tok/s (reasoning), 34.41 tok/s (coding), and 42.48 tok/s (prompt_2048tok_ops, 3.26 tok/round, 89.8% acceptance)!
    • Rounds 1–38 Optimization Blitz (bec6ab3, 4b0f233, d74a874, 89188ac, f15273e, f11c331, 059bf9c, 8bc5f03, 0548b0c, aeb4f92, 70aef90, 86255f2, c6fccd4, ceb1722, f73d5df, 2eb7dc0, f3384f6, 5dad020, 35f84cb, 75d719c, 4b7eddc, 0d39e8e, 3c41c24, e1737f7, a620276, a1f7083, 69531aa, 5dc9905, 2267f2a, 0958c6b, 9e4b432, eb15e06, aa55103, d94675c, 910b3ef, 9d9dc8a, 3404184, c005fa3, 6aeb30d):
      1. MTP STRATA_METAL_MTP_SINK=64 + mtp_wrote Valid-Run KV Fix + Batched 4-Row fc_hidden GEMV + slop4-Aware mtp_slop_pen (window.cpp, qsa.metal, mtp_chain.metal): Fixed multi-chunk MTP QSA KV zero-reads, preserved the 64-token system-prompt sink in MTP KV, batched fc_hidden (-20.9 MB/step), and synced slop4 state into the MTP draft head + wired SuffixDrafter::propose_after to extend partial neural drafts.
      2. M5 Multi-Token Decode & Prefill Kernels (gemv_cat, gr_norm_comb, gdn_step, gdn_scanr_4, qsa_prep, moe_gu/moe_down, mm_gu3, -1,247 Prefill Barriers, P_GSLOT L2 Fix): Fused ssm_alpha + ssm_beta (gemv_cat) and moe_combine_sh + gr_norm (gr_norm_comb, -83 dispatches & -48 barriers/window), optimized gdn_step (-12% at ntok=4, simd_broadcast) and gdn_scanr_4 (-8.4% at T=2048), parallelized qsa_prep, fixed the Apple M5 L2 cache-coherency race in EXPERT_SPLIT=1 via P_GSLOT (100% bitwise deterministic across all 48 layers), removed 1,247 redundant prefill barriers, and added M5 shape-aware gemm_f16 tiles + publish_ready gated on prompt.
      3. Round 4 Prefill + PLE + MTP + MoE Plan Fusions (gr_norm_h_moe, ple_moe_gate_norm, ple_query_gate_norm, mtp_rms_xprep, mtp_combine_gnorm, mm_place 32-subblock + mm_scan pad-clear, mm_gather_h): Folded moe_out into per-token 4-stream gr_norm_h_moe (-25%), fused ple_gnorm + ple_gate_norm (-42.5%, freeing 188.8 MB of prompt scratch RAM (2.35 -> 2.16 GB)), fused mtp_rms_xprep + mtp_combine_gnorm (mtp_ms down to 218 ms / 2.25 ms/draft), eliminated mm_clear and replaced mm_place's $O(1024^2)$ loop with 32-subblock nibble counting (12.9x faster mm_place, 5.89x faster MoE plan), and added FP16 mm_gather_h (1.70x faster gather, -40.2 ms/chunk).
      4. Round 5 Prefill + Decode + Scratch Aliasing + Deterministic Lookahead (gr_norm_h_moec, gr_apply_moec, gr_norm_h4, gr_norm rv[10], gr_norm_comb_4s, qsa_prep#qf16, -580 MB Scratch Aliasing, stream.cpp Lookahead Fix): Folded mm_combine into gr_norm_h_moec and gr_apply_moec (eliminating 47/48 mm_combine dispatches and 2.01 GB DRAM traffic/chunk), added per-token 4-stream gr_norm_h4 (-18.3%), skipped dead p.mixed FP32 writes in gemm_hc_up (-11.3%), cached gr_norm in registers + added 4-stream gr_norm_comb_4s, fused FP16 query output into qsa_prep#qf16 (1.98x faster, -8.91 ms/chunk), aliased disjoint GDN/QSA prompt scratch buffers (cutting MT=2048 transient prompt scratch from 2.16 GB -> 1.58 GB, -770 MB vs original 2.35 GB), and fixed stream.cpp at_layer to deterministically reclaim unused prompt lookahead slots.
      5. Round 6 QSA Flash + PLE + GDN + Register-Cached gr_norm_h* + -286 MB PLE/bo/yr Scratch Aliasing (qsa_flash, ple_moe_gate_norm_c, ple_conv_res, gdn_pre, gdn_post, gr_norm_h*): Cached half4 ksc/vsc scales in threadgroup memory and parallelized qsa_flash online softmax across all 4 simdgroups (8 lanes/head, -20.0% at T=1984: 34.36 -> 27.50 ms/layer, -26.6% at pos0=30000, T=96, memcmp == 0) + skipped qsa_topk when p.flash && pos0 + T <= 2051; eliminated gated DRAM writes/reads everywhere (-167.8 MB/chunk + -40 MB decode window RAM), added 1-TG-per-token ple_moe_gate_norm_c (-28.3%, eliminating the last mm_combine dispatch at Layer 0), and aliased PLE + bo/yr prompt scratch (cutting MT=2048 transient prompt scratch from 1.58 GB -> 1.29 GB, -1.06 GB / -45.1% vs original 2.35 GB); cached rv in registers across gr_norm_h* (-7.89 GB DRAM reads/chunk) and refactored gdn_pre (-16.1%, 0 barriers) and gdn_post (0 barriers) to 1 simdgroup per head.
      6. Round 7 Single-Pass simd_broadcast gdn_scanr_4 (-22.2%) + Cooperative-Register gemm_hc_up (-8.9%) + ple_hist_adv & -159 MB Scratch Aliasing (1.13 GB Scratch): Replaced TB=16 threadgroup staging and float4 kc[8] (32 VGPRs/lane) in gdn_scan_r with single-pass TB=0 scalar simd_broadcast (23.8 KiB -> 1 KiB TG memory, -256 barriers/layer, -127 wasted state passes, -22.2% / -55.6 ms across 36 layers at T=2048); added cooperative-destination register epilogue to gemm_hc_up (16 KiB -> 0 B TG memory, 64x64 < 1024, 64x128 >= 1024, -8.9% unbound / -14.0% bound); added ple_hist_adv GPU kernel and aliased p.normW = p.XS, p.HS = p.qkv, p.gsh = p.gqn, p.ush = p.gkn, p.hh = p.ygh, p.xqh = p.z (cutting MT=2048 transient prompt scratch from 1.29 GB -> 1.13 GB, -1.22 GB / -51.9% vs original 2.35 GB, and 0.708 GB at C=1024).
      7. Round 8 prompt_chunk sigs/cuts Shared-Expert + MoE-Plan Overlap + W8<IQ2_XS/IQ2_XXS> + gr_norm_h_stream_out + dq_f16/act_f16 half4 (audits/opus_r8_moemm.md, audits/sol_r8_grvec.md, audits/argon_r8_dq.md): Signaled sigs right after route_topk in Window::prompt_chunk so mm_hist, mm_scan, mm_place, mm_gather_h, and the 4 shared-expert GEMMs + swiglu_h run on the GPU concurrently beside the host/SSD at_cut(l) (5,165 -> 4,837 ms = -328 ms / 410.0 tok/s inside prompt_chunk!), folded W8<IQ2_XS> scale nibbles (-2.2..-3.2%) and W8<IQ2_XXS> popcount signs (-1.8..-2.3%), reduced gr_norm_h4 / gr_norm_h_moec stream-by-stream (-30 VGPRs/lane, -7.5 ms in kprof), and vectorized dq_f16 (half4 stores, barrier-free when TAB == 1) and act_f16 (float4 -> half4, 4x fewer threadgroups)!
      8. Round 9 qsa_flash 19,904 -> 7,616 B Threadgroup Memory (4 TGs/Core, -17.5% / -51.4 ms) + TAB > 1 Barrier Guards + IQ3_S Dead Sign-Load Removal + dn_body NE==1 (audits/opus_r9_qsa.md, audits/sol_r9_moemm.md, audits/argon_r9_quant.md): Shrunk qsa_flash TG memory from 19,904 B -> 7,616 B via 2-round 12-row s.store into sp[2][HPG*QCH], register-combined ((s0+s1)+s23)*0.0625f, and aliasing p (fp16) into dead sp[1] (unlocking 4 TGs/core on M5, cutting pg:qsa_flash x12 in kprof from 332.81 -> 281.41 ms = -51.41 ms / -15.4%, and -21.6% at pos0=30000), guarded load_tab + threadgroup_barrier with fmt<F>::TAB > 1 in mm_gu3, mm_down3, moe_gu, and moe_down, removed dead load_signs from load_tab<T_IQ3_S>, and added NE == 1 hoisting in dn_body!
      9. Round 10 qsa_flash 2.2x (-51.9% / -146.1 ms in kprof) Unconditional Hoisted rv/rk + vec4 Epilogue + gdn_pre float4 Conv + uint4 mm_gather_h + s_id-Free mm_place (audits/opus_r10_qsa.md, audits/sol_r10_gdn.md, audits/argon_r10_gather.md): Hoisted each lane's 8 cell row pointers and scales in qsa_flash rv and rk cooperative tensor fills and made the 64 byte/char4 loads unconditional (max(id, 0) with scale 0.0f on empty cells) + vec4 float4/half4 output epilogue, cutting pg:qsa_flash x12 in full-model kprof at T=1983 from 281.41 ms -> 135.36 ms (-146.05 ms / -51.9% vs Round 9, and -396.66 ms / 3.93x faster than Round 5's 532.02 ms!) and cutting total full-model kprof GPU time from 3,693.1 -> 3,548.4 ms (-144.7 ms); vectorized value-head float4 conv in gdn_pre (-2.0..-2.9%), upgraded mm_gather_h to 128-bit uint4, and eliminated threadgroup short s_id[1024] (2 KiB -> 0 B) in mm_place (-6.2% faster MoE plan)!
      10. Round 11 gemm_f16 Tile Selector (-18% / -234 ms) + qsa_attn_mm Hoisted rv/rk (-44.5%) + Branchless W8<T_IQ3_S> (-50.8% / 2.03x) & mm_gu3/mm_down3 Pointer/Block Hoisting (-57.4 ms in kprof) (audits/opus_r11_dense.md, audits/sol_r11_qsa.md, audits/argon_r11_moemm.md): Fixed Window::prompt_chunk's gemm_f16 tile selector (M > 1536 instead of M >= 2048) and added 4 new M5 tile instantiations (16_64_2, 32_64_2, 32_128_2, 64_32_2), ported Round 10's unconditional hoisted rk/rv fill to qsa_attn_mm (70.5 -> 39.2 ms, -44.5% / 1.80x faster), replaced divergent hh branches in W8<T_IQ3_S> with branchless byte-shift indexing (37.2 -> 18.3 ms/layer, -50.8% / 2.03x faster) + hoisted invariant row pointers (g0_l, u0_l, w0_l) and super-block boff in mm_gu3 and mm_down3 (-57.4 ms across 48 layers in kprof!), cutting total full-model kprof GPU time at T=1983 from 3,548.4 ms -> 3,397.5 ms (-150.9 ms vs Round 10, -828.2 ms vs Round 5 = 583.7 tok/s pure GPU prefill throughput on M5!) and cutting 10,180-token Albedo prompt prefill from 28.25 s -> 24.59 s (360.3 -> 413.9 tok/s, -3.66 s faster & -68.2 GB fewer SSD reads)!
      11. Round 12 qsa_topk 3.85x (-73.6%) + gemm_hc_up_reg 32_128_4 & Unrolled Epilogue (-6.9% / -7.8 ms/chunk) + T_Q6_K Hoisted mm_gu3 (-4.3% / -10.2 ms) + 3,375.4 ms GPU kprof (audits/opus_r12_moemm.md, audits/sol_r12_qsa.md, audits/argon_r12_gemm.md): Cached 4 order keys per thread in qsa_topk (nb <= 4096), replaced the tid == 0 256-bucket serial loop with a parallel simdgroup scan, and fused the (equal, greater) emit scan (4 -> 2 barriers per 1,024 blocks, 2.55x–3.85x faster qsa_topk: 1.861 -> 0.491 ms at T=1984, pos0=2048, 5.543 -> 2.157 ms at pos0=30000, memcmp == 0) + bounded qsa_flash's PV loop by c0 < nc; switched gemm_hc_up_reg to 32_128_4 with unrolled pointer-hoisted epilogue (-6.9% / -7.8 ms/chunk, 116.5 -> 108.8 ms); and enabled the pointer-hoisted fast path for T_Q6_K in mm_gu3 (-4.3% at T=1983, -6.3% at T=512), cutting full-model kprof GPU time at T=1983 to 3,375.4 ms (587.5 tok/s pure GPU throughput!).
      12. Round 13 gemm_hc_down + gemm_swiglu_up K4 Register Epilogues (-144 Dispatches, -985 MB DRAM Traffic/Chunk, 600.6 tok/s GPU kprof) + dq_f16 Exact-Simdgroup 2D Grid + SpecGate & mtp_draft PLE Prefetch (audits/opus_r13_hcdown.md, audits/sol_r13_specgate.md, audits/argon_r13_dqflat.md, audits/opus_colibri_scout.md): Fused gr_epi_down into gemm_hc_down's cooperative-tensor register epilogue (-96 dispatches, -96 barriers, -493.4 MB p.V DRAM writes/reads, -2.65 MB p.V scratch, -13..-27% faster at T <= 192, -4..-5% at T = 384..512) and fused swiglu_h into single-accumulator gemm_swiglu_up (-48 dispatches, -491.5 MB p.ush DRAM writes/reads, -11.5% / -3.2 ms/chunk at T=1024, -6..-12% at T=256..512, -2.8% at T=2048), cutting full-model kprof GPU time at T=1983 from 3,375.4 ms -> 3,301.86 ms (600.6 tok/s pure GPU prefill throughput on M5, -73.5 ms vs Round 12 and -923.8 ms vs Round 5, 2,298 -> 2,154 dispatches!); upgraded dq_f16 to hardware [[thread_position_in_grid]] with exact-simdgroup 2D threadgroup sizing (128/160/192/256 threads, 1.12x–1.27x faster on small/medium matrices) + T_Q5_1 instantiation; and added SpecGate (STRATA_SPEC_GATE=suffix|1) + mtp_draft tok PLE prefetch during the 2.1 ms MTP GPU run!
      13. Round 14 qsa_flash/qsa_attn_mm Byte-Offset & Zero-Scale Hoist (-8.7%..-9.7% across all pos0, 443.2 tok/s on 10.18K Albedo Prompt) + gdn_scanr_4 hk-Locality + Compacted moe_comb_slots roff + mm_gu3 IQ2_XS half4 HS Store (audits/opus_r14_moemm.md, audits/sol_r14_qsa.md, audits/argon_r14_gdn_gr.md): Hoisted per-cell byte offsets (cid * QKV + g) * QD) and zeroed ks_tg/vs_tg scale slots once per cell before the h = 0..3 head loop in qsa_flash and qsa_attn_mm (-8.7% at pos0=0: 10.33 -> 9.42 ms, -9.7% at pos0=2048: 21.80 -> 19.68 ms, -7.3% at pos0=8192: 26.82 -> 24.86 ms, -7.6% at pos0=30000: 32.62 -> 30.13 ms, cutting full-model kprof pg:qsa_flash x12 from 128.03 -> 121.93 ms = 4.36x faster than Round 5's 532.02 ms!); remapped gdn_scanr_4 to hk-grouped head scheduling (-2.18 ms across 36 layers at T=2048), compacted moe_comb_slots (roff[0..cnt-1] via simd_prefix_exclusive_sum) + unrolled stream loops in gr_norm_h4/gr_norm_h_moec (-2.4..-3.5% on gr_norm_h4, -0.9% on gr_norm_h_moec), and added half4 HS stores to mm_gu3 (T_IQ2_XS, -0.5..-0.6%), pushing prompt_albedo_11k (10,180 tokens) prefill to 443.22 tok/s (22.97 s vs 28.25 s, -5.28 s / +23.0% faster, with Chunk 1 hitting 501.5 tok/s / 4,084 ms end-to-end!).
      14. Round 15 prompt_chunk Router+Plan/Shared-Expert Wave Zipping (-190 Barriers) + Fused Wal+Wbe N=96 GEMM (-36 Dispatches) + QSA PSO Pre-Warm (-202 ms Chunk-1 Stall) + qsa_score_mm Causal-Future Tile Skip (-19.4%) + qsa_prep 320 B TG & 2 Barriers (audits/opus_r15_win.md, audits/sol_r15_qsa.md, audits/argon_r15_prep.md): Zipped router lookahead (dq(Wr2) + gemm(Wr2) + route_topk(ids2), -94 barriers) and zipped MoE Plan (mm_hist/mm_scan/mm_place) with Shared-Expert GEMMs (Wshg+Wsg/gemm_swiglu_up/Wshd, -96 barriers, 100% overlap of mm_scan/mm_place with shared-expert GEMMs); fused ssm_alpha + ssm_beta (N = 48 + 48 = 96 = 3 * 32, -36 dispatches, 2,154 -> 2,118 dispatches/chunk, -378 MB/chunk p.mixh reads); pre-warmed qsa_score_mm/qsa_topk PSOs in prompt_gemm_init (cutting Chunk 1 graph build from 203 ms -> 1 ms on 10.18K Albedo prompt!); added causal-future tile early-exit (uint(b0) >= max_bid) + dead[0..3] hoist to qsa_score_mm (-19.4% faster qsa_score_mm and -14.1% combined qsa_score_mm + qsa_topk at T=1984, pos0=2048) and removed the 3rd radix barrier in qsa_topk; shrank qsa_prep row[256] -> row[64] (1,088 -> 320 B TG memory) + cut single-call barriers from 4 -> 2 (-1.0..-1.9%), and unrolled route_topk for n_exp == 512 (-1.2..-2.2%), pushing pg:qsa_flash x12 in kprof down to 116.58 ms (9.71 ms/layer, 4.56x faster than Round 5's 532.02 ms!) and prompt_albedo_11k to 443.43 tok/s (22.96 s)!
      15. Round 16 Fused QSA Wk+Wv+Wik(+Wiq) Single GEMM (-24..-36 Dispatches, -5.6 MB Scratch -> 1.12 GB) + Decode Resident-MoE & Shared-Expert Wave Zip (-48 Barriers/Window) + gemv.metal TAB>1 Guards + ple_moe_gate_norm_c SIMD Compaction & ple_conv half4/t>=9 (audits/opus_r16_kvik.md, audits/sol_r16_dec.md, audits/argon_r16_ple.md): Fused QSA Wk (512), Wv (512), Wik (128), and Wiq (512) into a single N = 1152 (9 * 128) or N = 1664 (13 * 128) GEMM in p.Wk/p.kcur with zero-overhead function constant QSA_KV_STRIDE in qsa_prep + PSO pre-warming (eliminating 24–36 GEMM dispatches per chunk -> 2,094 dispatches, and freeing p.iqr/p.ikr/p.Wik = -5.6 MB prompt scratch -> 1.12 GB); zipped decode Window::layer's resident routed experts (phase = 1) with the shared expert inside shx (4 -> 3 waves, -48 barriers/window, starting moe_gu(phase=1) immediately at sigs and hiding swiglu_prep beside moe_down(phase=1)) + guarded load_tab + threadgroup_barrier on fmt<F>::TAB > 1 in gemv/gemv_cat/gemv_pair; and compacted ple_moe_gate_norm_c routed slots (-6.8..-10.3%) + vectorized ple_conv half4 weights with a branchless t >= 9 fast path (-4.7..-5.8% full PLE stage) + hoisted gdn_post row pointers (-0.9..-1.8%)!
      16. Round 17 mtp.metal float4 Vectorization (-21%..-45%) & Q8_0 gemv_catk + Barrier-Free simd_shuffle_down in gr_up_f/gr_up3/qsa_merge (-97 Barriers/Window, 1,024 -> 0 B qsa_merge TG) + gdn_step ib==0 Readout Pruning (-7.7%..-10.0%) & qsa_attn char4 K Loads (-6.7%..-11.4%) (audits/opus_r17_mtp.md, audits/sol_r17_shuf.md, audits/argon_r17_gdn.md): Vectorized mtp_prepare/mtp_xprep/mtp_rms_xprep/mtp_combine_gnorm with float4 loads/stores (-41..-45% on mtp_xprep, -21..-32% on mtp_rms_xprep, -9% on mtp_combine_gnorm) and added Q8_0 gemv_catk for mtp_encode's k_proj+v_proj and gate_proj+up_proj; replaced threadgroup s16 reductions in gr_up_f, gr_up3, and qsa_merge with barrier-free simd_shuffle_down (-512 B TG & -1 barrier in gr_up_f/gr_up3 = -97 barriers/window, and 1,024 -> 0 B TG & 0 barriers in qsa_merge) + added n == 2560 unroll fast path in gr_norm (12–40% faster); and pruned redundant ib = 1..3 red2 shared-memory reads + unrolled gdn_step (-7.7%..-10.0% faster multi-token verify, -3.8%..-4.4% faster 4-token commit) + vectorized qsa_attn K loads with char4 + hoisted qcur pointers (-6.7%..-11.4% faster across all pos0 & ntok).
      17. Round 18 Per-Stream gr_norm_comb 5.09x (-54%..-80%) + xprep float4 (-41.5%..-46.8%) & moe_combine_sh (-32%..-49%) + Q6_K moe_gu NE>=3 (-18%..-25%) + qsa_topk2/qsa_score (-5%..-9%) (audits/opus_r18_moe.md, audits/sol_r18_grcomb.md, audits/argon_r18_qsa2.md): Switched Window::gr_norm from 1-TG-per-token gr_norm_comb_4s (rv[4][10] register pressure + 100 plan loads/thread) to pointer-hoisted, unrolled 4-TG-per-token gr_norm_comb (grid = (4, n_tok), yep[0..9] hoisted, 54%–80% faster across ntok = 1..8: 22.67 -> 4.46 us = 5.09x at ntok=1); vectorized xprep (units.metal) with float4 stores (-41.5%..-46.8%: 19.8 -> 10.5 us), unrolled moe_combine_sh / moe_combine for topk == 10 (-45%..-49% at ntok=1, -32%..-33% at ntok=8), and fused qg/qu decode per entry in gu_body for UNIT == 32 (T_Q6_K -18%..-25% at NE >= 3); and ported register-cached keys[4] (nb <= 4096 with if (b >= nb) break;) + -7 radix barriers (20 -> 13 barriers = -84 barriers/window) to decode's qsa_topk2 (qsa2.metal) + hoisted/unrolled qsa_score (-4.9%..-7.6% faster combined qsa_score + qsa_topk2 across pos0 = 3000..30000).
      18. Round 19 gemv.metal NTOK/NR/KS Unroll (-5.4%..-29.4% at ntok>=2) + mtp_pick/mtp_slop_pen unroll_count(4) (-27..60%) + ple_comb_query_gate_norm & Load-Time window_init Prewarm (audits/opus_r19_prewarm.md, audits/sol_r19_sample.md, audits/argon_r19_gemv.md): Uncovered that because NTOK ([[function_constant(0)]]) is a Metal function constant rather than a C++ template parameter, the Apple Metal compiler was NOT unrolling for (int t = 0; t < NTOK; ++t) in gemv, gemv_pair, gemv_cat, gemv_catk, and gemv_x, forcing float acc[NR][8] to spill to thread-local stack memory on every u iteration whenever ntok >= 2; adding #pragma clang loop unroll(full) across NTOK/NR/KS loops makes Q8_0 GEMVs -5.4%..-23.7% faster, IQ3_S 32_2_4_1 GEMVs -7.7%..-20.8% faster, Q2_0 640->2560 -15.6%..-29.4% faster, and gemv_cat/gemv_catk -6.6%..-15.8% faster at ntok = 2, 4; unrolled mtp_pick (~27–39% faster) and mtp_slop_pen (~23–60% faster) + Q4_K/Q5_K mtp_embed (~27–28% faster) in mtp_chain.metal; and fused decode Layer 0's gr_norm_comb + ple_query_gate_norm into ple_comb_query_gate_norm (-12.5%..-22%, -1 dispatch, -1 barrier, -1 dead RMSNorm) + register-cached ple_gnorm (-29%..-52%) + pre-warmed t = 1..max_t and MTP encode/pick PSOs in window_init!
      19. Round 20 Window::mtp_load 707.8 MB Expert mlock (mtp_ms 398 -> 122.2 ms on 11K, 433 -> 202.4 ms on 2K, 40.57 tok/s Decode) + Decode qsa_attn (-20%..-40.5%) & qsa_merge (-9.6%..-12.7%) + gdn_scanr_4 32-Bit Base Offsets (186.15 ms in kprof) (audits/opus_r20_mtp_qsa.md, audits/sol_r20_gather.md, audits/argon_r20_gdn.md): Uncovered that Window::mtp_load (window.cpp) was not mlocking the 707.8 MB of T_Q2_0 routed experts in Wm_ (mtp-q2_0.gguf), allowing macOS to evict them during 1.12 GB prompt-scratch prefills and causing +220..305 ms of hard SSD page faults during mtp_draft; wiring (mlocking) the 707.8 MB of active MTP experts + MADV_DONTNEEDing 181.2 MB of unused pages cut mtp_ms from 397.88 -> 122.18 ms (3.26x faster!) on the 10,180-token Albedo prompt and from 432.82 -> 202.37 ms (2.14x faster!) on prompt_2048tok_ops, pushing decode to 40.57 tok/s (24.65 ms/tok, 3,130.15 ms); register-cached s0, s1 across softmax + unrolled V x4 in decode qsa_attn (-20.4%..-40.5% faster across pos0 = 512..8192) and shared ek[]/L_sh in qsa_merge (-9.6%..-12.7% faster); and replaced 64-bit per-lane pointers in gdn_scan_r with 32-bit base offsets (s0, k0, v0, ks, rs) + final-token dead-lookahead elimination, cutting pg:gdn_scanr_4 x36 in kprof at T=1983 to 186.15 ms (5.17 ms/layer, -61.4 ms / -24.8% vs Round 5)!
      20. Round 21 gr.metal GR_NTOK Unroll (-5.4%..-18.0% across 194 Dispatches/Window) + qsa2.metal (-36.8%..-40.6% qsa_attn2 & -8.3%..-14.6% qsa_merge2) + gdn_step 32-Bit st_off & moe_plan topk10 unroll_count(10) (-9.5%..-26.6%) (audits/opus_r21_gr.md, audits/sol_r21_qsa2.md, audits/argon_r21_gdn_moe.md): Unrolled GR_NTOK ([[function_constant(1)]]) and KS/UNIT/hc loops in gr_down_f (GR_NTOK <= 6, -5.4%..-7.0% on gr_down_q8_8_4, -13.4%..-14.2% on gr_down_q8_8_8) and gr_up_f (-1.6% at ntok=1, -6.0% at ntok=2, -14.7% at ntok=4, -18.0% at ntok=8) + moved gr_q2s/gr_p4s to program-scope constant (67/67 bitwise PASS); ported register-softmax, unrolled 8-cell V blocks, shared ek[]/L_sh, and barrier-free qsa_write_prepared to qsa_attn2 (-36.8%..-40.6%) and qsa_merge2 (-8.3%..-14.6%); and replaced gdn_step_t's 64-bit st_base with 32-bit st_off + float4 conv_w + constant gdn_q2s/gdn_p4s, and replaced moe.metal plan_body's runtime integer division i / int(a.topk) with topk10 ? ((i * 205) >> 11) : (i / int(a.topk)) + unroll_count(10) (-9.5%..-20.4% faster moe_route and -11.7%..-26.6% faster moe_plan at ntok = 2..8), setting all-time decode speed records across all 4 workloads!
      21. Round 22 gr_down_4_4 / _r2 NR=2 xn Reuse (-8.1% ntok=4, -16.8%..-26.1% ntok=5..8) + moe_down / moe_gu q2_combine (-7.7%..-22.5% at NE >= 2) + plan_body sg==0 SIMD Prefix Sum (-14.6%..-35.4%) + Barrier-Free swiglu_prep (audits/opus_r22_gr.md, audits/sol_r22_router_glue.md, audits/argon_r22_moe.md): Added template NR + direct float4 FMA (no xs array copy) to grd_acc / gr_down_f, using 128-thread gr_down_*_4_4 at ntok <= 4 (-1.9%..-8.1% faster) and gr_down_*_4_4_r2 / 8_8_r2 (NR = 2 rows/simdgroup sharing xn reads) at ntok >= 5 in gr_read and mtp_gr_read (-16.8%..-26.1% faster, 75/75 bitwise PASS); added q2_combine<F> (ushort2 pk[NR]; float d[NR]; = 4 registers for NR=2, loading xs[16] once per token t across all NR rows) in moe.metal, cutting moe_down_142_8_2_4 by -7.7%..-8.6% at NE = 2..4 (-16.1% on T_Q2_0 prep), Q2_0 moe_gu by -12.2%..-22.5%, and IQ2_XS moe_gu by -8.2%..-17.1%, plus sg == 0 simd_prefix_exclusive_sum in plan_body (-8.8%..-22.7% faster moe_route and -14.6%..-35.4% faster moe_plan); and converted swiglu_prep (glue.metal) to program-scope constant glue_q2s/glue_p4s + barrier-free simd_shuffle_down (64 -> 0 B TG mem, -48 barriers/window).
      22. Round 23 ksg popcount (-2.5%..-5.5% IQ2_XS moe_gu & -3.9% IQ3_XXS) + gemv_direct_x/gemv_q2_raw_combine (-3.4%..-5.4% Q8_0&-2.2%..-4.1% Q2_0R) + mm_gu3 sign4_iq2xs & Constant-Ptr Folding (776.7 ms IQ2_XS&45.9 ms IQ3_S, 3,293.32 ms/602.1 tok/sGPUkprof, 42.01 tok/s Decode) (audits/opus_r23_quant.md, audits/sol_r23_gemv.md, audits/argon_r23_moemm.md):** Computed ksg sign parity (i | ((popcount(i) & 1u) << 7)) for IQ2_XS (ctab == 1, -2.5%..-5.5%fastermoe_gu) and IQ3_XXS (TAB = 256, -3.9%) in quant.metal; added direct float4FMAgemv_direct_x (Q8_0 ntok=1&LPR=16, NR=1, ntok=2, -3.4%..-5.4%faster) andgemv_q2_raw_combine (Q2_0R ntok=1, -2.2%..-4.1%faster) ingemv.metal; and added 4-bit register sign4_iq2xs(j, hh)(zeropopcountonhh==0 lanes) + constant-offset pointer folding ((k & 255) >> 3/>> 2) + fmascale math inmoe_mm.metal, cutting W8<T_IQ2_XS> by **-4.19%..-4.35%** (pg:mm_gu3_22_16_32_32_4 x34 down to **776.71 ms/22.84 ms/layer, -25.5 ms/chunk**) and W8<T_IQ3_S> by **-8.5%atT=2048 (45.93 msforx2) and 1.82xatT=512 (18.73 -> 10.28 msacross the VGPR occupancy cliff!)**, pushing full-modelkprofGPU time atT=1983 down to **3,293.32 ms (602.1 tok/spure GPU prefill)** andprompt_2048tok_ops decode to **42.01 tok/s (23.80 ms/tok, 3,023.11 ms`)!
      23. Round 24 W8<T_Q6_K> Scalar uint OR & Address Folding (+12.8%..+24.9% at T=512..1024, +3.67% at T=2048) + qsa_flash 32-Bit ko[8]/vo[8] Offsets (-11.71% at pos0=8192) + Dead Sign-Table Removal (TAB = 1024/512) + gr_norm_h_moec k-Outer s[10] & gemm_hc_down * 0.25f (audits/opus_r24_prompt.md, audits/sol_r24_qsa.md, audits/argon_r24_moemm.md): Folded W8<T_Q6_K> and W4<T_Q6_K> address/shift math and replaced 4-wide uint4 unpack+shift+OR with a single scalar uint OR before float4(as_type<uchar4>(qa)), speeding up Q6_K mm_gu3 by +12.79% at T=512, +12.28%..+24.91% at T=1024, and +3.67% at T=2048 (36.82 -> 35.52 ms/layer); moved sign4_iq2xs to quant.metal, updated deq<T_IQ3_XXS/T_IQ2_XXS/T_IQ2_XS> to use raw s on the low 4 values and sign4_iq2xs(s, 1) on the high 4 values, eliminated ksg and load_signs across all formats, and shrunk fmt<T_IQ2_XS>::TAB (1024) and fmt<T_IQ2_XXS>::TAB (512); replaced 64-bit kr[8]/vr[8] pointers in qsa_flash with 32-bit uint ko[8]/vo[8] byte offsets (kq + (ko[j] + uint(k)) and vq[vo[r + j] + uint(d)]), cutting qsa_flash at T=1984, pos0=8192 by -11.71% (74.61 -> 65.87 ms for x3 = -35.0 ms across 12 layers!); and inverted gr_norm_h_moec's k/i loops (-0.5..-0.6% at T=2048) + replaced / float(e.hc) with * 0.25f in gemm_hc_down (-0.8..-1.1%).
      24. Round 25 mm_down3 (T_Q2_0, TK=32) 64-Element Block Stepping (k0 += 64, w0_l += 18) + W4<T_Q2_0>::run_q<2/6/10/14>(rdw, uint kl4) (-11.8% / -42.7 ms in kprof: 361.9 -> 319.2 ms) + gdn_post Gate/Gamma Preload (-0.7%) + mm_gu3 half4 HS Store Pointer Hoist (audits/opus_r25_gdn.md, audits/sol_r25_ple.md, audits/argon_r25_moemm.md): Unrolled mm_down3 (T_Q2_0, TK=32) by 2 TK=32 steps per 64-element T_Q2_0 block (k0 += 64, w0_l += 18) with W4<T_Q2_0>::run_q<2/6/10/14>(rdw, uint kl4) using (b + kl4)[OFF] immediate displacement, speeding up mm_down3 by +2.48% at T=1024 (5,020.1 -> 4,898.7 us), +1.85% at T=2048 (8,060.3 -> 7,914.2 us in isolated A/B, and 361.9 -> 319.20 ms = -42.70 ms / -11.8% across 48 layers in full-model kprof!) + hoisted ot_p/hs_p in mm_gu3's half4 coalesced HS store (+0.09..+0.91% across IQ2_XS/IQ2_XXS/IQ3_XXS/Q2_0); and preloaded gamma and precomputed the sigmoid gate before the 4 simd_sum reductions in gdn_post (-0.7%).
      25. Round 26 qsa_attn_mm 32-Bit ko[8]/vo[8] Offsets & float4 TG Loads (-23.3%..-29.9% across pos0=2048..30000) + gemv T_IQ4_XS packed_uint4 5-Register gemv_iq4_combine (-1.5%..-5.0% Cold DRAM on 2560->640, -1.0%..-2.7% at ntok=3..4 on 2560->6144..12288) + DW<T_IQ4_XS/T_IQ4_NL/T_Q6_K> Unroll & MM_V3_ONLY (audits/opus_r26_gemv.md, audits/sol_r26_qsa.md, audits/argon_r26_moemm.md): Replaced 64-bit kr[8]/vr[8] pointers in qsa_attn_mm with 32-bit uint ko[8]/vo[8] byte offsets + unroll hints + float4 threadgroup loads, cutting qsa_attn_mm (T=1984) by -23.31% at pos0=2048 (30.88 -> 23.68 ms), -29.93% at pos0=8192 (31.79 -> 22.28 ms), and -24.74% at pos0=30000 (37.32 -> 28.09 ms); added gemv_iq4xs_codes (packed_uint4 pk[NR] + float d[NR] = 5 registers/row instead of 36 for float q[32] + float4 sc, gated on LPR <= 16 || NTOK >= 4) + GEMV_AB_COLD=1 DRAM-streaming A/B benchmark in gemv.metal/gemv_parity.cpp, speeding up cold-DRAM IQ4_XS 2560->640 (ffn_gate/up_shexp x57/step) by -1.5%..-5.0% at ntok=1 and -2%..-3% at ntok=2..4, and 2560->6144/10240/12288 (attn_gate/qkv/q) by -1.0%..-2.7% at ntok=3..4; and added direct unrolled DW<T_IQ4_XS> and DW<T_IQ4_NL> + rd32 SWAR DW<T_Q6_K> (1.02x–1.07x faster dq_f16) in mma.metal + #ifndef MM_V3_ONLY (10x faster moe_mm_parity --v3 compile).
      26. Round 27 gdn_step_t 5 -> 4 Barriers/Token + dt/ssm_a Hoist + part[8]/red2[IB-1][S] Shrink (-3.0%..-7.9%) + qsa_merge2 / qsa_merge Parallel float2 (m, l) Loads & simd_max (-8%..-20%) + Window::gemv_tile n_out=512 {32,2,4,1} (-3.2%..-25.5%) & Q2_0 640->2560 {8,2,4,1} (-4%..-6%) (audits/opus_r27_gemv.md, audits/sol_r27_qsa.md, audits/argon_r27_gdn.md): Eliminated a 512-thread threadgroup_barrier per token (5 -> 4 barriers/token = -144 barriers per 4-token verify window across 36 GDN layers!) in gdn_step_t (gdn.metal) by writing raw cv to qs/ks/vs before barrier 1 and computing (rsq * qs[qk_idx]) * scale and (rsk * ks[qk_idx]) * scale after barrier 1, plus hoisting token-invariant dt[h] & ssm_a[h] and shrinking part[32] -> part[8] and red2[IB][S] -> red2[IB-1][S], winning all 16/16 production plane = 49152 configurations (-7.93% at ntok=4 verify, -7.82% at ntok=4 commit, -3.17% at ntok=1 verify, 88/88 bitwise); vectorized (m, l) loading in qsa_merge2 and qsa_merge (qsa2.metal, qsa.metal) as float2 ml = *(device const float2*)(pp + tid * STRIDE + QD) + simd_max(M) + threadgroup lk[k] readout, speeding up qsa_merge2 by -14%..-20% at ntok=1 and -8%..-14% at ntok=4 (80/80 bitwise); and fixed Window::gemv_tile (window.cpp) so 2560 -> 512 (attn_k & attn_v across 12 QSA layers) uses {32, 2, 4, 1} (preserving LPR=32 for 100% bitwise parity across all 13 formats, verified by gemv_parity's new same LPR: bitwise check) and Q2_0 640 -> 2560 at ntok >= 4 uses {8, 2, 4, 1} (-3.2%..-25.5% faster IQ4_XS 2560->512, -9.0%..-24.4% faster IQ3_S 2560->512, -7.6%..-13.9% faster Q6_K 2560->512, -4%..-6% faster Q2_0 640->2560).
      27. Round 28 mtp.fc_hidden [t][c] Stream Batching (-52.1% at n=2, -29.3% at n=3, -17.0% at n=4) + ple.metal tg_sum Barrier Halving (-1.5%..-5.8%) + sample.metal tg_argmax Unroll & router_route_8 (audits/opus_r28_mtp.md, audits/sol_r28_router.md, audits/argon_r28_ple.md): Batched mtp.fc_hidden.weight across all n * HC streams in [t][c] order (<= MAXT rows per gemv dispatch in window.cpp & mtp.metal), cutting the MTP hidden branch at n=2 (a=2 accepted drafts) by -52.1% (211.6 -> 101.4 us, 9 -> 3 dispatches), at n=3 (a=3 accepted drafts) by -29.3% (220.1 -> 155.7 us, 9 -> 5 dispatches), and at n=4 by -17.0% (228.9 -> 190.0 us) while matching n separate n=1 chains bit-for-bit; halved threadgroup_barriers in ple.metal (2 -> 1 in ple_gnorm, 4 -> 2 in ple_gate_norm, 6 -> 3 in ple_query_gate_norm & ple_comb_query_gate_norm, 8 -> 4 per stream in ple_moe_gate_norm & ple_moe_gate_norm_c) via ping-pong part[16], speeding up decode/verify PLE (T=1..7) by -1.5%..-3.7% and T=64 ple_moe_gate_norm_c by -1.5%..-5.8% (100% bitwise); and added #pragma clang loop unroll(full) to tg_argmax in sample.metal + router_route_8 (T_Q8_0) in router.metal.
      28. Round 29 mm_gu3<T_IQ2_S> Super-Block Scale Caching (-86.24 ms / -10.0% in kprof) & W8<T_IQ2_XXS> Byte Loads (-15.42 ms / -6.5% in kprof) + moe_route e512/res_mask (-20.6%..-26.4%) & deq<T_IQ2_S / T_IQ2_XXS> (-1.1%..-6.9%) (audits/opus_r29_prompt.md, audits/sol_r29_moemm.md, audits/argon_r29_moe.md): Cached per-row FP16 super-block scales dg[0..3] / du[0..3] across the 8 TK=32 steps per 256-element block in mm_gu3<T_IQ2_S> and shrank fmt<T_IQ2_S>::TAB from 2080 to 2048 (16,512 -> 16,384 B exact 16 KB TG memory), plus direct byte loads p[l] / p[l+2] in W8<T_IQ2_XXS> and fma in W8<T_IQ3_S>, cutting pg:mm_gu3_22 (34 T_IQ2_S layers) in kprof from 859.30 -> 773.06 ms (-86.24 ms / -10.0% faster, +5.48% at T=1024, +5.08% at T=2048 in isolated A/B) and pg:mm_gu3_16 (12 T_IQ2_XXS layers) from 237.28 -> 221.85 ms (-15.42 ms / -6.5% faster); eliminated dynamic register indexing v[ex / 32] = -INFINITY in both route_topk (-0.7%..-1.2%, 5.78 -> 5.305 ms in kprof) and moe_route (moe.metal), where adding e512 + preloaded 16-bit res_mask from smap cut moe_route (split=1) by -26.4% at ntok=1, -24.4% at ntok=2, and -20.6% at ntok=4; and optimized decode deq<T_IQ2_S> (ushort q2/s2 + pre-shifted qh + #pragma clang loop unroll(full), -1.1%..-6.9% faster across QUANT_CTAB=0/1), deq<T_IQ2_XXS> (-5.1%..-5.9% faster), and deq<T_IQ3_XXS> (-0.7%..-2.9% faster).
      29. Round 30 W4<T_Q2_0>::dec IEEE-754 FP16 0x6400 Exponent Trick (-10.5% mm_down3 at T=2048 / -37.02 ms in kprof) + gdn_pre & gdn_post Vector simd_sum(float4) Reductions (audits/opus_r30_moemm.md, audits/sol_r30_prompt.md, audits/argon_r30_gdn.md): Replaced FP32 int->float->half4 conversion in W4<T_Q2_0>::run_q with an IEEE-754 FP16 mantissa-OR decode (const uint2 c = (uint2(q | (q << 14), (q >> 4) | (q << 10)) & 0x00030003u) | 0x64006400u; return (as_type<half4>(c) - 1025.0h) * d;), proven 100% bit-for-bit identical across all 65,536 x 256 = 16,777,216 FP16 d x code-byte combinations, cutting mm_down3<T_Q2_0, 16, 64, 32, 2> by -10.5% at T=2048 (7,950 -> 7,113 us), -8.3% at T=1024, and -37.02 ms (312.39 -> 275.37 ms / -11.9%, down -23.9% from Round 24's 361.9 ms!) in full-model kprof; and replaced the 12 scalar simd_sum(float) trees in gdn_pre and 4 scalar simd_sum(float) trees in gdn_post with bitwise-identical vector simd_sum(float4) reductions (12 -> 3 and 4 -> 1) + cold t == T - 1 / a.f32 branch hoists (-0.66%..-1.05% on gdn_post).
      30. Round 31 mm_gu3<T_IQ2_S / T_IQ2_XXS / T_IQ2_XS / T_IQ3_S> tab + hh & Sign-Shift (-45.32 ms IQ2_S / -12.56 ms IQ2_XXS in kprof) + moe_down / gemv Q2_0R 0x6400 Half-Mantissa Decode (-7.6%..-16.9%) + dq_f16<T_IQ4_XS / T_IQ4_NL> Register Table Init & perm == 4 Contiguous Writes (audits/opus_r31_moemm.md, audits/sol_r31_moe.md, audits/argon_r31_dq.md): Hoisted invariant hh = (kl >> 2) & 1 and th = tab + hh across mm_gu3, replaced the divergent hh ? s >> 4 : s ternary in W8<T_IQ2_S> and W8<T_IQ3_S> with s >> (4 * hh), and extended 256-wide super-block scale caching dg[]/du[] to T_IQ2_XXS and T_IQ2_XS, speeding up mm_gu3<T_IQ2_S> by +8.17%..+12.79% (773.06 -> 727.74 ms in kprof, -131.56 ms / -15.3% faster than Round 28!), mm_gu3<T_IQ2_XXS> by +4.08%..+6.86% (221.85 -> 209.29 ms in kprof), mm_gu3<T_IQ2_XS> by +1.41%..+6.16%, and mm_gu3<T_IQ3_S> by +2.23%..+6.84% (100% bitwise PASS); replaced scalar float((pk >> ...) & 3) - 1.0f decode in q2_combine<T_Q2_0R> (moe.metal) and gemv_q2_raw_combine (gemv.metal) with packed IEEE-754 FP16 0x64006400u mantissa-OR decode, cutting decode moe_down (dn, all 48 layers!) by -7.6%..-16.9% (ratio 0.8311..0.9236 across NE=1..4 and ntok=1..8) and gemv Q2_0R (ntok=1) by -11.8% (100% bitwise PASS); and replaced the dynamic load_tab<T_IQ4_NL> loop and divergent constant-memory kvalues_iq4nl reads in dq_f16<T_IQ4_XS / T_IQ4_NL> (dense_mm.metal) with 4 bank-conflict-free register stores, guarded a.perm, and inverted ro = gid.y for perm == 4 (1.03x–1.05x on 320x10240p4, 100% bitwise PASS).
      31. Round 32 qsa_flash sp[4][HPG * QCH] 3 -> 1 Barrier & Pre-Added ko[j] + Loop Unrolling (-4.03%..-5.47% / -6.10 ms in kprof) + mm_gu3<T_IQ2_S> half4 Coalesced HS Store (-9.00 ms in kprof) + Decode deq 64-Bit uint2 Grid Loads & deq_iq_packed (audits/opus_r32_moemm.md, audits/sol_r32_qsa.md, audits/argon_r32_moe.md): Expanded sp[2][HPG * QCH] to sp[4][HPG * QCH] (13,760 < 16,384 bytes TG memory) in qsa_flash (qsa_prompt.metal) so all 4 simdgroups store S partials simultaneously across 1 threadgroup_barrier instead of 3 (-64 barriers/token at pos0 >= 2048), pre-added d0 + kl to ko[j] for constant-offset +0/+16/+32/+48 K loads, and unrolled fixed-trip loops, winning 7/7 pairs at every position on the quiet GPU (-4.31% at pos0=0, -5.47% at pos0=2048, -4.03% at pos0=8192, -0.80% at pos0=30000, and 130.50 -> 124.40 ms / -4.7% in full-model kprof!); enabled the half4 coalesced HS store epilogue in mm_gu3 (moe_mm.metal) for T_IQ2_S, cutting pg:mm_gu3_22 (34 layers) by another -9.00 ms (727.74 -> 718.74 ms in kprof, -140.56 ms / -16.4% faster than Round 28, and pushing T=2048 under 25,000 us to 24,986.6 us!); and replaced pair 32-bit grid loads in deq<T_IQ2_XXS / T_IQ2_XS / T_IQ2_S> (quant.metal) with 64-bit uint2 loads (4 -> 2 loads per unit), unrolled deq<T_IQ3_S> (-2.7%..-6.6% faster IQ3_S moe_gu), and added low-register 6-VGPR deq_iq_packed<F> + iq_combine (-9.4%..-15.8% at NR=1, NE=4 in moe.metal).
      32. Round 33 gdn_step_t Critical-Path kl Split + Pre-Barrier-3 state Commit (-1.1%..-4.1% Across 32/32 GDN Configs) + gr_norm Apply-Hoist & Pre-Offset Pointers (-6%..-11%) + GrdUnit<T_Q8_0, 2> Packed char4 Decode (-10%..-16%) (audits/opus_r33_gr.md, audits/sol_r33_qsa.md, audits/argon_r33_gdn.md): Split tot_q / rsq / ql / g / be / vj across Barrier 2 in gdn_step_t (gdn.metal) so only kl sits on the critical path before the 32-step kv FMA loop, moved hist_out writeback before Barrier 1 (killing hst live ranges across all 4 barriers on GDN_NTOK == 1), and issued the 32 global state stores before Barrier 3 to overlap the 64 KB/head state writeback with ib == 0's readout reduction (0.9589x..0.9895x on plane = 0, 0.9717x..0.9959x across all 16/16 production plane = 49152 configs, 88/88 bitwise PASS) + pre-offset moe_combine_sh & swiglu_prep (glue.metal, 24/24 bitwise PASS); hoisted if (a.apply) and pre-offset pointers in gr_norm (gr.metal, -6%..-11% at ntok=1..4, -2%..-3% at ntok=7..8) and added GrdUnit<T_Q8_0, 2> keeping Q8_0 codes packed in char4 q[2][8] (8 VGPRs/row vs 32 floats) for _r2 gr_down (-13%..-16% at ntok=7, -10%..-15% at ntok=8, 100% bitwise PASS), cutting pure GPU kprof at T=1983 to 3,153.46 ms (628.51 tok/s pure GPU prefill, 3,270.77 ms serialized sum!) and slashing mtp_ms to new all-time records across reasoning (323.71 ms, -12.4% vs 369.70 ms), persona_chat (383.38 ms, -9.8% vs 424.85 ms), and prompt_2048tok_ops (193.49 ms).
      33. Round 34 gemm_hc_down TN = 96 Tail-Tile Decomposition (-5.05%..-7.18%) + grd_acc ROLL = true at KS = 4 (-1.9%..-4.6% on _r2) + W8<T_IQ3_S>::block & Coalesced half4 HS Store + T_IQ2_XS sboff (audits/opus_r34_gr.md, audits/sol_r34_prompt.md, audits/argon_r34_moemm.md): Replaced gemm_hc_down<TM, 128, SG>'s 128-column tail tile (tg.x = 2, c0 = 256, where only 68 columns 256..323 exist) with a single 96-column (3 x 32) hc_down_tile<TM, 96, SG>, cutting gemm_hc_down on the quiet GPU by -5.05% at T=512 (ratio 0.9495), -6.60% at T=1024 (ratio 0.9340), -7.18% at T=1983 (1.1762 -> 1.0927 ms, ratio 0.9282), and -7.03% at T=2048 (1.2062 -> 1.1209 ms, ratio 0.9297, making gemm_hc_down faster than gemm_hc_up in full-model kprof for the first time ever: 125.998 ms vs 132.822 ms!); added ROLL = F == T_Q8_0 && NR == 2 && KS == 4 to grd_acc (gr.metal) so packed Q8_0 NR=2 (GrdUnit<T_Q8_0, 2>) keeps its token loop rolled at ntok = 4..6 (-4.6 / -3.1 / -1.9% faster 4_4_r2 at ntok = 4/5/6, -13.7% and -23.9% faster than 4_4 at ntok = 5, 6) while keeping window.cpp:1165 at ntok > 4 to preserve 324-threadgroup cold-DRAM latency hiding at ntok = 4; and added W8<T_IQ3_S>::block + F == T_IQ3_S half4 HS epilogue stores (+0.48% at T=2048) + T_IQ2_XS sboff reuse (-0.69% at T=1024, -1.25% at T=2048) in moe_mm.metal, pushing prompt_2048tok_ops to 432.78 tok/s (4,584.29 ms end-to-end prefill) and 42.27 tok/s (3,004.68 ms decode) and @10.2K decode to 25.04 tok/s (39.93 ms/tok, 1,237.88 ms)!
    1. Round 35 Wr/Wr2 Router Dequant Reuse (-47 Dispatches/Chunk) + gemm_f16_16_96_2 Walbe + M % 128 <= 64 64_128_4 (1.14x @ M = 306) + hc_up_tile Row-Tail Trim (-2.5% @ T = 289) + mtp_rms_xprep / mtp_embed / mtp_slop_pen Vectorization (audits/opus_r35_dense.md, audits/sol_r35_prompt.md, audits/argon_r35_mtp.md): Ping-ponged p.Wr and p.Wr2 by layer parity (l & 1) in Window::prompt_chunk (window.cpp), eliminating 47 redundant dq_f16_0 router dequantizations per chunk (2,094 -> 2,047 dispatches in kprof), added gemm_f16_16_96_2 (TN = 96) for Walbe at M >= 640 (1.16x / -13.6% faster at M = 1983), selected 64_128_4 instead of 128_64_4 on big-N GEMMs when M < 1536 && M % 128 in 1..64 (1.14x / -23.5 ms faster chunk total at M = 306), added hc_up_tile<16/32, TN, SG, REG> row-tail trimming to gemm_hc_up (-2.5% at T = 289), and vectorized/unrolled mtp_rms_xprep (-37.5% at n=2560, -35.8% at n=10240), mtp_embed (-15.8%..-31.6% across all 13 GGUF formats), mtp_slop_pen (-96.5% on mode=0, -45.7%..-70.6% on mode=1, -31.0%..-63.9% on mode=2), and router_route (-0.2%..-2.1%) (100% bitwise PASS across all suites!).
    2. Round 36 Decode predL_ / predL2_ Token-0 Fast Path When STRATA_METAL_PREFETCH_DRAFT <= 0 (-16.0% Hot / -6.8% Cold Across 93 Router GEMVs/Window) + Guarded TM > 32 Row-Tail Trimming in gemm_hc_down & gemm_swiglu_up (audits/opus_r36_decode.md, audits/sol_r36_prompt.md, audits/argon_r36_qsa.md): Added tile_ntok to Window::gemv (window.cpp) so when STRATA_METAL_PREFETCH_DRAFT <= 0 (= 0 in production, where ExpertStream::ahead only reads t = 0), the 93 router-guess GEMVs (47 predL_ + 46 predL2_, Q8_0 2560 -> 512) run at ntok = 1 on the window's tile (gemv_8_16_1_4_1#1), producing 100% bit-for-bit-identical t = 0 predictions while cutting their GPU time by -16.0% hot (21.78 -> 18.30 us) and -6.8% cold (23.91 -> 22.29 us); added if constexpr (TM > 32) row-tail (TM = 16 / 32) trimming to gemm_hc_down (-4.71% at T=1025, -3.73% at T=1041) and gemm_swiglu_up (-1.42% at T=1025, -1.14% at T=1041, -0.78% at T=1987) while preserving TM = 32 (SG = 2) kernels byte-for-byte (252/252 epi bitwise checks PASS); and confirmed qsa_prep / qsa_score_mm are at their M5 register/barrier optimum.
  8. Round 37 Multi-Chunk Prompt Hot-Expert LRU & Scratch Pool Preservation (stream.h, stream.cpp, window.cpp): Passed multi_chunk = (n > C) into prompt_region(true, n > C) -> hold_region(true, multi_chunk) so 1-chunk prompts (<= 2,048 tokens) run 100% identically (42.49 tok/s / 2,989.17 ms @2K decode — new all-time record, 88.9% accept), while multi-chunk prompts (> 2,048 tokens, e.g., prompt_albedo_11k.txt at 10,180 tokens) keep !x.hot scratch slots at lru_back(s) across chunks, mark pass == 1 demotions x.hot = false in victim() so pass == 0 reuses them instead of evicting more hot decode experts, and restrict allow_hot lookahead to distance == 1 — boosting @10.2K decode from 25.26 tok/s (1,227.18 ms, 45.2% accept, 2,713 SSD reads) to 34.01–34.74 tok/s (892.47–911.51 ms, +34.6–37.5% faster, 78.3% accept, 2,073 SSD reads / -23.6% fewer SSD reads) and @10.2K prefill from 423.48 tok/s (24,038.85 ms) to 458.34–472.89 tok/s (21,527–22,210 ms, -1.83s to -2.51s faster), with pure GPU T=1983 prefill (kprof) hitting 3,131.29 ms (633.28 tok/s — new all-time record).

  9. Round 38 hold_region Layer-0 allow_hot = false + Prompt/Decode bytes_read_/reads_ Split + mtp_prefill tail_cap Sync with STRATA_METAL_MTP_WIN (stream.h, stream.cpp, window.cpp, audits/astra_r38_qsa.md): Passed allow_hot = false to victim() during hold_region(true)'s initial Layer-0 prefetch (stream.cpp:1080) and fallback static lookahead (stream.cpp:890) so hold_region only fills the 251 !hot scratch slots of cls_[0] (IQ2_XXS) instead of evicting ~180 hot decode experts for tail Layer-0 experts that at_layer(0) cancels 5 ms later; split p_bytes_read_/p_reads_ from decode bytes_read_/reads_ and zeroed bytes_read_/reads_ after the 6.5 GB startup arena fill (fixing the 1-window **decode** report from 116,128.9 MB / 79,058 reads to the exact 375.2 MB / 258 reads; corrected afterwards: that line prints 480 distinct lookups = 48 layers x top_k 10, i.e. one decode step, not a prefill. prefill byte/read accounting is destroyed before the final report unless STRATA_METAL_TRACE=1, because strata_window_cli.cpp:1827 calls window_reset_stats(w) unconditionally and window.cpp:4137 -> stream.cpp:1177-1180 zeroes p_lookups_ p_hits_ p_pre_ p_misses_ p_streamed_ p_chunks_ p_bytes_read_ p_reads_. any prefill MB/read figure published without TRACE=1 is decode-only); synced Window::mtp_prefill's default tail_cap with STRATA_METAL_MTP_WIN clamped to [64, 2050]; and verified in qsa_parity --ab (audits/astra_r38_qsa.md) that caching keys[8] across qsa_topk2's 1024-thread barriers drops M5 occupancy (+3.5–12.9% slower at pos0 = 10180), pushing @10.2K prefill to 473.80 tok/s (21,485.94 ms — new all-time record!), @10.2K decode to 35.21–35.47 tok/s (873.92–880.47 ms, mtp_ms = 106.34 ms — new all-time record!), @2K prompt_2048tok_ops decode to 42.57 tok/s (2,982.99 ms, mtp_ms = 194.09 ms — new all-time record!), and reasoning decode to 37.20 tok/s (5,886.81 ms — new all-time record!).

  10. round 39 — selected strata v0.1.45 ports and checked local fixes: incremental prompt tokenization, bounded parser scans, content-keyed logit-bias caching, mac telemetry, thinking budgets/answer reservation, and PLE prefetch-key reuse (36381540). local review removed upstream's stale identity/length bias shortcut, capped prompt-cache entries at 16, and cached/released the telemetry device query. preserved the text-only image-omission switch. metal prompt load safety and the default moe parity tile are fixed; the local qsa_merge2 batching experiment is bitwise exact (40 edge cases), but only tests dispatch it; live engine decode remains on qsa_merge. rejected the cuda score butterfly (memcmp failure) and qsa_merge batching (repeatable nt=3 slowdown). active-chunk bounds were already present. standalone CLI smoke: 42.46 tok/s at 1,984 prompt tokens and 34.23 tok/s at 10,180; these are not paired speedup evidence. a later alternating before/after check produced identical token IDs in all five pairs, but substantial temporal drift in both builds left performance inconclusive (2k candidate decode time +1.5–1.7%; 10k mixed -3.6% to +2.8%). no zero-regression claim. serving now enables upstream answer_room by default (for thinking with output caps >=1,024 and no explicit budget), which can inject a wrap-up instruction; disable it explicitly for parity comparisons to prior service behavior. exact verification and known baseline failures: audits/r39_integration.md, audits/r39_serve.md, audits/r39_opus_review.md; evidence: results/round39/.

  11. Round 41 Bonsai-Style Automatic NVMe Prompt & Conversation .kvc Cache (--cache-disk / --cache-disk-size 16384) + 2-Prompt Shared-Prefix Learning & Pinning + MTP Sliding-Window Session Persistence (window.cpp, strata_window_cli.cpp): Ported the automatic NVMe .kvc disk tier and 2-prompt shared-prefix learning/pinning architecture from okami.mom/llama-cpp-bonsai (prompt-cache-disk, cache-tuning, prefix-learn, prefix-restore) into strata_window_cli --serve and window.cpp: (1) preserved the MTP draft layer's valid KV cell range (mtp_sink_, mtp_cell0_, mtp_cell1_) in scr::SessionHeader::flags across Window::session_save / Window::session_load and shifted mtp_cell0_/mtp_cell1_ in Window::scr_relocate so restored/relocated sessions retain full MTP draft context without re-prefilling; (2) added DiskCache (--cache-disk <dir> / STRATA_METAL_CACHE_DISK, --cache-disk-size <mib> / STRATA_METAL_CACHE_DISK_MIB, default 16384 MiB) with < 1 ms startup .kvc header indexing (scr::SessionHeader + n_ids * 4 token IDs + .pads + .pin without reading the 120 MiB+ KV payload), corrupt/partial file cleanup, post-DONE zero-latency auto-save (4-thread F_NOCACHE pwrite at 5.2 GB/s), and automatic obsolete-prefix pruning; (3) added observe_prefix (prompts-seen.bin) which detects when any 2 distinct prompts diverge after a shared prefix >= 32 tokens (lcp < cur.size() && lcp < s.size()), splits prefill at learn_at = lcp, saves a pinned .kvc + .pin checkpoint (up to 16 pins, up to 64K tokens, never evicted by normal conversation LRU), and ported Bonsai's prefix-share (8906fd8) as zero-copy GPU MTLBuffer delta .kvc storage (c0 = (parent_cells / 4) * 4 in SessionHeader + .parent sidecar + orphan pruning, writing only [c0, n_cells) QSA/MTP positional cells per child and reading [0, c0) directly from the parent .kvc via F_NOCACHE pread); and (4) wired find_best in serve_loop so cross-conversation switches, server restarts, and middle-edited prompts ([A B C] -> [A B' C D]) automatically save the dirty resident session to NVMe, restore the best matching .kvc session (pread at 5.2 GB/s), and run either exact continuation (cached_tokens = 104 / 130 on Conv B Turn 2, cached_tokens = 65 / 89 on brand-new Conv C Turn 1 from pinned system prefix) or scr_plan + scr_relocate splice directly over NVMe (cached_tokens = 149 / 184 on middle-edited Conv D Turn 2!). Also fixed SAVE / RESTORE wire responses (SAVED <tokens> <bytes> <ms> / RESTORED <tokens> <bytes> <ms>, 30/30 serve.test_slots unit tests OK). 40. Measured End-to-End Speed (32K Context, 6.5 GB Arena, Unpruned K=512 IQ2_XS, call_3817687, call_2483052, call_3102851, call_2898988, call_4110683, call_4178375, call_3948986, call_4500449, call_3444931, call_3116575, call_4373130, call_2769318, call_3781775, call_4410602, call_3511955, call_4432071, call_2918543, call_1514614, call_1173886, call_2456120, call_5024768, call_4223761, call_3543222, call_1624568, call_3678353, call_4256391, call_3418394 & call_3706633): prompt_2048tok_ops prefill 400.88–432.78 tok/s (4,584.29–4,949.14 ms end-to-end — all-time record 4,584.29 ms / 432.78 tok/s — -1,728 ms / +37.7% faster end-to-end, 3,131.29 ms / 633.28 tok/s pure GPU kprof — 3,267.43 ms serialized sum!, 2,047 dispatches (-47 dispatches/chunk!), 1.12 GB transient scratch!), decode 41.53–42.27 tok/s (23.66–24.08 ms/tok, 3,004.68–3,005.02 ms decode, mtp_ms = 190.62–197.99 ms / 1.63 ms/draft — new all-time MTP record 190.62 ms!); prompt_albedo_11k (10,180 tokens) prefill 421.49–446.59 tok/s (22.79–24.15 s vs 28.25 s, -5.46 s faster — all-time record; Chunk 0 = 4,613–4,719 ms, Chunk 1 = 501.5 tok/s / 4,084–4,572 ms, Chunk 2 = 4,362–4,736 ms, Chunk 3 = 4,610–4,911 ms, Chunk 4 = 4,685–4,898 ms, mtp_ms = 121.54–125.14 ms, and @10.2K decode hit 25.00–25.26 tok/s / 39.59–40.01 ms/tok / 1,227.18–1,240.21 ms — all-time record 25.26 tok/s at 10.2K context! vs 22.06 tok/s / 1,405.37 ms!); coding prefill 120.61–126.78 tok/s (2,413–2,537 ms, +32.2%), decode 35.46–37.1 tok/s (6,115.57–6,176.28 ms decode — all-time record 6,115.57 ms, mtp_ms = 374.22–380.29 ms); reasoning prefill 121.57–122.92 tok/s (2,416–2,458 ms, +33.3%), decode 34.92–36.74 tok/s (5,961.16–6,270.75 ms decode, mtp_ms = 322.48–370.78 ms — all-time record 322.48 ms!); persona_chat prefill 124.62–128.14 tok/s (2,255.32–2,443 ms, +38.5% — all-time record 2,255.32 ms / 128.14 tok/s!), decode 28.50–33.2 tok/s (6,628.00–7,685.17 ms decode, mtp_ms = 374.58–424.81 ms — all-time record 374.58 ms!).

  12. recurrent-state INT4 / stochastic-rounding quantization in Gated DeltaNet (2610.15139, 2610.15709):

    • trap: treating recurrent state like a once-written KV cache. with per-token dynamic scaling, $s_{t,c}=\max|\widehat S_{t,c}|/q_{\max}$ reads its own quantization-corrupted state: outliers increase scale, bulk resolution coarsens, and subsequent errors feed back into the state. scale-yoking to the FP trajectory removes numerical runaway but leaves substantial functional damage; bounded norm is not preserved language-model quality (2610.15139, sec_mechanism.tex, sec_whatworks.tex).
    • source corrections: stochastic rounding raises per-step RMS error by $\sqrt{2}$ (measured $1.42\times$), not $\sqrt{23}$; on GDN-0.8B at yoked INT4 it adds 6.71 nats excess NLL. the first GDN layer alone accounts for 63–78% of summed single-layer INT4 damage at H=2048; this is not a measured first-three-layer error fraction. protecting the first three recurrent layers is a separate matched-budget allocation result, identical to rescue selection/knapsack on their small GDN models, not established for Swift.
    • mechanism beyond runaway: small updates vanish on the coarse state grid and errors can accumulate coherently: $E_T=\sum_{\tau\le T}A_{\tau:T}R_\tau$, with off-diagonal error inner products in $|E_T|_F^2$. 2610.15709 preserves updates with a low-rank bypass (LRB) and bounded directional rounding (BDR), rather than assuming unbiased noise is safe. its GDN result uses 5.25 average / 6.00 peak state bits including bypass, not bare INT4; WikiText-103 8K PPL is 26.85 vs FP32 24.19 and naive INT4 79.76.
    • rule: keep Swift's existing high-precision recurrent path. our roughly 112.2 MiB fixed GDN state is small next to the 6.50 GB expert arena; do not trade it for unvalidated long-horizon drift. any future state compression must test natural-continuation KLD/NLL, long exposure, session/SCR and speculative rollback, not just NIAH or state norms. LRB/BDR and ReplaySSM are research candidates, not proof that arbitrary recurrent INT4 is safe.
  13. routing-frequency / entropy / phase-fragility proxy bit allocation without measured distortion curves (2610.13328, 2610.20826):

    • trap: using a cache-popularity rank, routing entropy, or phase-specific fragility as a substitute for $d_e(M)$ (Hessian-/activation-/decision-weighted error or direct ablation). DuoRate Proposition 1 cancels frequency only when the same positive $f_e$ weights distortion and read cost: $\arg\min_M f_e[d_e(M)+\lambda R(M)]=\arg\min_M[d_e(M)+\lambda R(M)]$. promotion is ordered by $\Delta d_e/\Delta R$, not popularity. Proposition 2 retains frequency under an unweighted storage constraint and through cross-domain mixtures $\pi_{d\mid e}=w_df_{d,e}/\sum_{d'}w_{d'}f_{d',e}$; it does not say frequency is useless for cache admission or SSD traffic.
    • negative result: on Qwen-3.6-35B-A3B, 2610.20826's phase/entropy/CV/Jaccard/top-1 proxy is neutral near 4 bpw and loses at the steep IQ2 floor: uniform KLD 0.034321, proxy 0.038142–0.038249, random 0.037728/0.037808 (random marginally better, within about 1 SE of proxy). proxy p99 KLD rises 0.437 -> 0.487–0.492; mandatory downgrades inject tail damage. this is a coding-domain, layerwise negative result, not a theorem against all mixed precision or a measurement of our IQ2 formats.
    • our boundary: the K512 header has 144 packed expert tensors, not 144 independently precision-assigned experts: 68 IQ2_S, 48 Q2_0, 24 IQ2_XXS, 4 IQ3_S; each tensor packs 512 experts with one type. ukisai's imatrix/GSQ-RCO artifact is not DuoRate's HQQ/GPTQ/MiLo/MoEQuant/BitsMoE baseline. its IQ codes are not prefix-readable StagQ planes; truncating their bytes is corruption, not Stage-2 read-tier selection.
    • rule: preserve the current quant assignment unless a format-/byte-/arena-matched ablation beats it on mean and tail KLD, top-1, persona/JA/ops/tool behavior and measured SSD/GPU time. cache ranks may optimize reuse, but must not silently become precision ranks. cache-conditioned SSD miss weights also differ from distortion weights, so Proposition 1 cannot be copied unchanged into stream.cpp.
  14. measurement traps found by re-measuring the chunk-0 MoCE default: the noise floor was measured on an arm that does not share the mechanism, "reproducible" was claimed from n=2, position-level KL exonerated a change that flipped a boolean, and prefill counters are invisible without STRATA_METAL_TRACE=1 (call_1381557, call_6642651, call_1274067, sol-moce-chunk0 d18939749fb6f294fb0c272a61cca5d1, audits/opus_prefill_opt.md:82-89,92-120):

    • trap 1 — a control can be deterministic for a reason its treatments do not share. repeat0 (arm 0) came out exactly 0.000000 mean KLD, 64/64, on both corpora and both cycles. that is structural, not luck: prompt_moce = 0 unpacks to p_prot = 0, p_bonus = 0 (prompt.metal:41-47), so the bias path at prompt.metal:51 never fires and routing never reads smap, making arm 0 residency-independent. every boosted arm reads smap[idx] != 0xFFFFFFFFu at :55, so its routing depends on which experts are instantaneously resident and therefore on I/O timing. a floor measured on the arm that ignores the mechanism bounds nothing for the arms that depend on it. the earlier 0.0012 / 95.3% floor (audits/opus_prefill_opt.md:83-87) was a property of that streamed configuration, not a constant, and must not be quoted as one.
    • trap 2 — n=2 agreeing is not reproducibility. ops incumbent KLD came out 0.055902 then 0.055018 (spread 0.000884, 1.6%), and that was written up as "reproducible ... tighter even than its chat value". a third run of the same config gave 0.076038 (+38%). boosted-arm KLD estimates on ops are not reproducible to better than ~0.02. require a third run before writing the word "reproducible".
    • trap 3 — heavy-tail KL and argmax damage are largely disjoint sets, so position-level KL alone can exonerate a bad change. the incumbent's highest-KL positions were The (0.873594) and plan (0.674194) — prose, argmax unchanged — with 49.65% of KL in the top 4 positions and 72.47% in the top 8. the actual argmax flips were elsewhere and small-KL, including true -> false at position 28 (KL 0.307667) and backtick -> () at position 24 (0.035575). one boolean inversion per 64 ops tokens is disqualifying on its own.
    • trap 4 — metric power is corpus-dependent, so a metric that is saturated on one corpus can be the deciding evidence on another. top-1 was 64/64 for every arm on chat (zero discriminating power at that KL magnitude) and 60/64, 61/64, 59/64 / 59/64, 58/64 on ops. audits/opus_prefill_opt.md's "base vs base top-1 95.3%" was likewise a chat-class result. never declare a metric uninformative from one corpus.
    • trap 5 — 64-token teacher-forced PPL cannot resolve a ~0.05-nat KL, and no standard error is derivable. KL = H(p,q) - H(p) connects teacher-forced NLL to KL only if the targets are p-samples; fixed corpus ids are not, so neither the identity nor a Monte-Carlo SE applies. cand_ppl 2.4035 vs ref_ppl 2.5027 at n=64 therefore corroborates and contradicts nothing. quote KLD (an exact per-position expectation over the vocab) as the resolving statistic and say PPL does not estimate full-vocab KL.
    • trap 6 — prefill byte/read accounting does not exist unless STRATA_METAL_TRACE=1. strata_window_cli.cpp:1827 calls window_reset_stats(w) unconditionally after prefill + the first head step, window.cpp:4137 prints only under PKPROF/TRACE and then always reset_stream_stats(), and stream.cpp:1177-1180 zeroes p_lookups_ p_hits_ p_pre_ p_misses_ p_streamed_ p_chunks_ p_bytes_read_ p_reads_. the shipped 375.2 MB / 258 reads in results/round39/*.log:25 reports 480 distinct lookups = 48 layers x top_k 10 = one decode step, not a prefill. also p_streamed_ is a count of submissions, not bytes (stream.cpp:893,931,1117); bytes is p_bytes_read_ (:1192).
    • trap 7 — never rank arms from a sequential sweep. two arms with the identical chunk-0 setting landed 10,075.62 ms and 12,115.41 ms (20% apart) purely from sweep position, via the arena fill read rate going 2.2 s -> 0.8 s as the SSD warmed. interleave (A/B/A/B), discard each arm's first run, report the median of paired ratios, and keep the paired-with-off comparisons separate from arm-vs-arm block comparisons (different arms are each paired against off, not against each other).
    • rule: a residency-biased routing knob must be quality-tested on ops, not just chat; an arm that reads residency is stochastic by construction and an arm that does not read it is bit-exact by construction; state arms in n/64 and name the corpus; and a rejected optimization is worth more when the report states all three of speed (real), mechanism (real), and quality cost (real, corpus-specific, unaffordable).

B. ruled out from the original blueprint #

details and primary-source checks in blueprint_review.md.

  • the blueprint's expert size math and mixed-quant expert tiers: claimed 224 experts × 48 layers at Q3_K is 14.2 GB (real: 22.7 GB, or 16.5 GB at Q2_0), and proposed mixing Q3_K hot experts with Q2_K cold experts inside one layer (impossible in GGUF: ffn_*_exps is one 3D tensor per layer with one quant type [src: ISTA card]).
  • stage 0 PLE 28.8 GB n-gram table pruning / vector editing: the table is hashed bigrams/trigrams with no 1-to-1 vocabulary mapping, stays on SSD (~16 row reads/token) with zero VRAM cost, and the cited "3.6% cliff" was from a letter-lipogram paper (banning the letter e in Gemma-3 4B) [src: 2505.20501].
  • runtime n-gram logit bans for slop: Gadsby-4B proved hard token bans lobotomize math/reasoning ("17×23 = sixteen forty nine"), and backtracking on Flash-Next requires restoring ~110 MiB GDN recurrent snapshots. replaced with Antislop FTPO training.
  • prefix sliding baked into training: breaks long-horizon agent recall and invalidates prefix/KV reuse across turns. kept only as an untrained runtime toggle (default off).
  • the blueprint's SCR GDN snapshot restore ("restore pre-edit 36-layer GDN snapshot + 4-token forward pass"): wipes the entire post-edit suffix (C) from 36 of 48 layers. replaced with CLM's rule: keep the post-edit GDN state and measure staleness.
  • stripping <think> from active context by default: Flash-Next's template preserves thinking by default, and Whittle measured post-success rewrite loops when reasoning was stripped [src: Whittle card].
  • <j_state> Jacobian-lens telemetry injected into the prefix: mutating the prefix on every turn invalidates position 0 and forces a full re-prefill.
  • the 7-term composite loss and misattributed papers: "DDCA" is RLVR advantage shaping [src: 2602.02099], "SBT-E" is SFT data construction [src: 2505.14604], "MoCCE" is training-free MoCE reranking whose own learned prior failed to beat training-free [src: 2412.00099], and the 0.25 PLE margin loss was taken from Whittle (which only needed it because its backbone wasn't pretrained with the PLE table, and needed margin 1.0–7.5, not 0.25 [src: Whittle card]).
  • custom GSQGroupQuantLinear QAT stack: wrong bitwidth (4-bit vs Q3_K/Q2_0), broken per-group scale broadcast, and unnecessary when freezing routed experts and transplanting dense tensors into ukisai's GSQ-RCO GGUF.