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 #
- writing off
ds4-recipeas "wrong model, no MoE tiering".- mistake: dismissed
astrra.space/ds4-recipeat 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.
- mistake: dismissed
- 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.
- 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.
- 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.
- 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.
- 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_Sas an offline teacher viallama.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.
- 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.
- 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.
- 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-recipedoc 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.
- 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).
- 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.
- 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.
- 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).
- trusting
F_NOCACHE=1on macOS APFS without checkingiostat -Id(disk0).- trap caught during
bench/run_02: repeatingpreadover identical file offsets across benchmark reps on APFS can partially hit cache on rep 2+ even when the descriptor hasF_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(disk0physical MiB delta == 100.0% of requested MiB).
- trap caught during
- 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 flashandgemini 4 argonturns (plusAPPEND_SYSTEM.md), while all assistant transcripts can be used to mine LLM-isms (mine_llm_isms_from_transcripts.py).
- clarified (dawn): don't dump arbitrary session transcripts into calibration or training. for phase 3 voice examples from albedo, filter only to
- 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.txtbefore checking HF. on a 512-expert MoE, a 5K-token corpus leaves many experts untouched or under-sampled;REAP-256-duospecifically 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).
- mistake caught by dawn: wrote 8 synthetic conversations (~21 KB) for
- decaying EMA scores only on resident experts (
opus-auditorcatch inreplay_moe_cache.py).- bug caught by
tangled-clod/claude-opus-5-5:near_pick_emaandinverse_router_emainitially decayedema_scoreonly forc in cache_setwhile adding credit to non-resident experts intop_n, letting non-resident experts pile up undecayed sums and stay pinned for ~99 tokens once admitted. fixed by decaying alln_expertsevery step.
- bug caught by
- six Phase 1 calibration/eval traps caught by
opus-methodology(audits/opus_audit_phase1_methodology.md):- missing
--parse-specialinllama-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
--chunkscutoff + char vs token balancing: concatenating all 20 Terminal-Bench docs beforeukisaimeant a 65K-token prefix gaveukisaiandhotchpotch0 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/vsmedium/),ukisai(_run2.._run5/__dsv4-trial-2suffixes + the shared 2,834-char Terminus-2 preamble), andChiTako(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, usingxhigh/.../prompt.txt, and verifying 20-char shingle disjointness. - Whittle double-templating: Whittle's
promptcolumn already contains<|im_start|>system ... <|im_start|>assistant\n<think>\n; wrapping it again produced nested templates. fixed to useprompt + responsedirectly 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 theQ4RTtrace.
- missing
- 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 inscripts/select_and_eval_prunes.pyv3 to use a zero-prior damped LCB ((c / (c + m)) * max(0, mean - SE), scoringcount == 0strictly at0.0with a fail-closedassert zero_hit_kept == 0). - outlier max-normalization in
reap_hybrid: a single extreme expert in Layer 47 (e143mean 5.56 vs layer median 0.31) squashedmax(shrink)normalization; fixedreap_hybridto 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_usedcandidates survive in top-64: fixedscripts/replay_moe_cache.pyandscripts/replay_moe_cache_fast.cppto countmax(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).
- count-weighted shrinkage prior (
- matching
strncmp(t->name, "ffn_moe_weighted-", 17)without checkingt->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-closedGGML_ASSERT(down_t->op == GGML_OP_MUL_MAT_ID)on the first ubatch. fixed by checkingt->op == GGML_OP_MUL(andt->op == GGML_OP_MUL_MATforffn_moe_logits-).
- bug: in
- re-parsing
Q4RTbinary 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 + runningmin(candidates, key=lambda...)across 120 configs × 48 layers × 16,384 tokens timed out at 300s twice. fixed by writingscripts/replay_moe_cache_fast.cpp(g++ -O3 -fopenmp), cutting the 120-config 48-layer sweep to 24–47s.
- mistake: unpacking 786,432 structs per config (
llama-perplexityspecial-token (parse_special = false) and--save-all-logitsCLI traps:- trap 1: unlike
llama-imatrix,llama-perplexityrejects--parse-specialon the CLI AND callscommon_tokenize(ctx, params.prompt, true)withparse_special = falseby default intools/perplexity/perplexity.cpp(lines 310, 473). always patchperplexity.cpptocommon_tokenize(ctx, params.prompt, true, true)before evaluating chat/agent traces. - trap 2:
--save-all-logits FNAMEand--kl-divergence-base FNAMEare aliases for the sameparams.logits_fileoption; passing both on the baseline run treats--kl-divergence-baseas the filename. alwaysassert p.returncode == 0in eval wrappers so bad CLI flags fail closed immediately instead of recording-1.0.
- trap 1: unlike
- slicing 60 GB candidate GGUFs onto Colab
/contentoverlayfs (~300 MB/s) instead of/dev/shmRAMtmpfs(83 GB):- mistake: writing each 64.8 GB
K320candidate to/content/prunedtook216.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 afterK512_unprunedandK320_reap_l2_sumcompleted (when a Colab VM vanishes mid-job,job.poll()remainsNoneuntilcolab.instances()is checked). fixed by slicing temporary candidates into/dev/shm/pruned(83 GBRAMtmpfson the 167 GB High-RAM A100) and savingresults/phase1c_kld_partial.jsonlocally.
- mistake: writing each 64.8 GB
- running Tier-2 generation evals at
max_tokens=640andtemperature=0.0on a thinking model (run_phase1c_tier2_eval.py):- mistake: with
--reasoning-format deepseek, any prompt whose<think>...</think>trace exceeds640tokens leavesmessage.content == ""(confounding normal 700–1,200 token technical JA explanations onK512_unprunedwith genuine<think>stall loops), and naive 5x 4-gram counting flags common English prepositional phrases ("in the kv cache"). fixed by settingmax_tokens=2048(temp=0.6, top_p=0.95, top_k=20, seed=42), explicitly trackingempty_content_stall+completion_tokens, and checking consecutive 20-char repeats + 8x 5-grams.
- mistake: with
- Phase 1c KLD & Tier-2 rubric traps caught by
opus-review-1c(audits/opus_audit_phase1c.md):delta_pplclamped-base bias & JA PPL non-monotonicity:--save-all-logitsquantizes base log-probs touint16over[max_logit - 16, max_logit], soPPL(base)in KLD mode sits 0.6% / 4.0% / 6.5% below exact FP32K512.ppl, and pruned models have up to 42% lower JA PPL (K256_shrinkgets16.2PPL with the worst0.520KLD) 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 recordexact_delta_ppl = ppl_q - base_fp32_ppl.- Terminal-Bench execution-contract footer in
eval_ops_tools.txt:Lottolabs/terminal-bench-2.1prompt.txtfiles end in a shared ~1.1k-charExecution contract: / Benchmark rules:block.scripts/assemble_corpus.pynow strips\n\nExecution contract:ontb_eval_docsand placesukisaimulti-turn traces first (while keepingtb_calib_docsunchanged socalibration_mix_v2.txtreproduces Phase 1a bit-identically); note that the Phase 1cops_toolsKLD numbers inresults/phase1c_kld_results.jsonwere measured on the pre-dedupe 4-chunk window. - Tier-2 rubric false negatives & truncation accounting: checking case-sensitive
"WAL"whenPERSONA_SYSmandates lowercase failedjournal_mode=wal, and requiring katakanaエキスパート/チャンクfailed validexpert/chunkloanwords 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 >= 2048truncation failure (4of the pruned failures are 2048-cap truncations, andops_1/ops_3keep their v1 verdicts since v1 did not save raw tool args);scripts/run_phase1c_tier2_eval.pyis updated tomax_tokens=4096with full raw output logging for Phase 3/4 re-runs.
- 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: anyp_pooled = (c_llm + c_hum) / total_bgterm inp_bg(w)makesα_wproportional toc_llmatc_hum == 0and 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, with0.0self-pooled weight andsum_w α_w == α_0), soc_llmnever entersα_wand all 23c_hum == 0rows 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 insideWORD_REglued 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,(letpretokenizer merges, and cross-family floor: normalized\u2018/\u2019 -> "'"before tokenization, checked sub-tuples againstPROTECTED_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>= 2model families each with>= max(5, 5% of c_llm)hits, read root-sessiontranscriptuserrows while stripping>quote lines and skippingisSidechainsubagent rows, and used the miner as a confirmation filter over curated rhetorical/signpost candidate lists (RHETORICAL_ALLOWLIST+SIGNPOST_PREFIXESminusEXCLUDED_WORK_AND_GENERIC_WORDS, dropping"think that's","so far", and"as expected"from FTPO).
- Monroe pooled-background cancellation at
- 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) intophase3_sentence_surgery(25pre-tool announcements, closing permission-solicitations, sycophantic openers — removed at the sentence/clause level in Phase 3 SFT data viascripts/prepare_phase3_sft_data.py) vsphase4_lexical_ftpo(7surface rhetorical cliches ->126Qwen 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).
- trap: if the model's hidden state already intends to ask a servile closing permission question or narrate a tool call before emitting
- four Phase 3 prep blockers caught by
opus-review-phase3-prep(audits/opus_audit_phase3_prep.md):Q2_0codec vsggml-quants.c(B1):ggml-quants.c(quantize_row_q2_0_refL74,dequantize_row_q2_0L439) uses scaled = amax, zero-point1((q - 1) * d->{-d, 0, +d, +2d}), and packs 4 consecutive elements per byte (qs[j // 4] >> ((j % 4) * 2)), NOT(q - 2) * dwith quarter-strided layout. Fixed inscripts/phase3_lora_and_transplant.pyand verified in--self-testagainst a compiled C reference transcribed fromggml-quants.c(max_err == 0.0across 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 includingU+2600..U+27BFinUNICODE_EMOJI_REstripped kaomoji symbols (✧ ♡ ♥ ✿ ☆ ♪in(◕‿◕✿)). Fixed inscripts/prepare_phase3_sft_data.pyby performing in-place span deletion (absorbing only whitespace directly attached to removed emoji so0/140fenced Gemini reference turns and padded tables/kaomoji are altered) and narrowingUNICODE_EMOJI_REtoU+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", protectedREADME.md/Makefile/HOME/K=320/MemoryHighduring lowercasing, required word-boundary matching + template pool inclassify_rtuning_abstention, and mandated running speech-act surgery before teacher-forced top-k logit extraction.
- 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_AARCH64weight repacking) on by default:src/llama-model.cpp(use_mmap_buffer) allocates an anonymous heap buffer in RAM to repack--cpu-moeexpert tensors at load time, bypassingmmapand tripping the64.0 MiBswapout watchdog before adding--no-repack/-nr([unarchived pre-baseline smoke run]).POSIX_MADV_WILLNEEDon--cpu-moeexperts + APFS sequential read-ahead (F_RDAhead) +QD=1serial page faults: patched inscripts/patch_m5_lazy_moe.py. OnK256_reap_l2_sum, unpatchedllama-cli(chat mode ondata/m5_val_prompts/llama_prompt_1264tok.txtwith-n 64, ~64 generated tokens / 63 decode steps) ran at1.40 tok/sdecode and14.30/21.70 tok/sprefill (-ub 1024/-ub 2048,105,851.73/85,725.73 MBwhole-rundisk0read across load + prefill + decode), whereas patchedllama-completion -no-cnv -t 8decoded the 22-token prompt at3.22 tok/s(33,428.34 MBwhole-rundisk0across load + 22 prompt + 63 decode tokens) and prefilled the 5,600-byte (1,264-token) raw prompt at38.11/63.27 tok/s(-ub 1024/-ub 2048, hitting EOS immediately after prefill with0decode tokens so46,808.58/28,272.94 MBmeasure load + prefill only; note the two runs also differ inllama-clichat vsllama-completion -t 8raw mode).--moe-cache-mib 1536degrades decode to0.28 tok/s, andllama.cpp-cmoeremains CPU-compute-bound becauseggml-cpucomputes expert matmuls on the CPU instead of on the M5 GPU.
- 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 = 512across C++ (window.cpp,stream.cpp,prompt_chain.inc) and Metal shaders (router.metal,moe.metal,prompt.metal): patched to readn_expertdynamically (K = 256 / 320 / 224 / 512), run the prompt router GEMM overn_exp_rows (ulong(t) * ulong(n_exp)stride inroute_topk), and maskidx >= n_expertto-INFINITY(route_parity: 0 mismatches across 200 windows atK=512, pad=0(2.16e-07max rel err default math /2.15e-07fast math) and 200 windows atK=256 + MoCE=0.85with randomsmapand256..511traps (3.06e-07max rel err default math /4.00e-07fast math)). ExpertStream::load_denseallocating3.24 GBof anonymous RAM (cx.alloc): upstreamload_densealready skipped the1.28 GBBF16hc_*originals and copied the remaining3.24 GB(2.90 GBGPU dense +0.34 GBCPUtoken_embd) into anonymous RAM. Patchedload_denseto wrap<model>-metal.gguf(4.53 GBfile map) zero-copy andmlockonly the2.90 GBof active GPU dense tensors while leaving0.34 GBCPU-onlytoken_embd.weightand1.28 GBskipped BF16hc_*unpinned viaMADV_DONTNEED(~0.34 GBreal resident RAM saving vs upstream's3.24 GBanonymous copy).PromptGemm(hold_region) carving prompt buffers out of the expert arena andMADV_DONTNEEDnot unmapping Darwin anonymous scratch: patchedExpertStream::region+Window::prompt_gemm_need_bytes()(verified equal toneedat runtime inprompt_gemm_init) to map a separate1.36 GBMAP_ANONbuffer on demand withoutmemsetting fresh zeroed pages, refreshslot.hotprotection on the topk.n - (ne_ + 32)LRU slots on every prompt turn (in_prompt_), and immediatelycx_->free+munmapthe1.36 GBscratch inprompt_region(false)after prefill (phys_footprint 4.73 -> 4.68 GB,resident RSS 7.41 -> 7.36 GBon the1,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_NLexpert slices (23.42 GiB, 3 fragmented size classes) vs uniformQ2_0(15.82 GiB, 1 unified size class of85pages): built multi-threaded C++ converterconvert_k256_q2_0(vendor/strata/metal/src/convert_k256_q2_0.cpp), shrinkingK256experts from23.42 GiB(1,550slots /1,440filled across 3 classes in a3.40 GBarena) to15.82 GiB(2,441slots in 1 unified class in the same3.40 GBarena). - race-free MoCE cache-prior reranking (
STRATA_METAL_MOCE_BONUS=0.85) inmoe_route/router_route+ExpertStream::ahead& 3-domain complete-turn assistant quality measurement: wiredsmapresidency intomoe.metal/router.metalusing a per-lane 16-bit register mask (scaled_mask+simd_broadcast, readingsmaponce atrank == 3so background I/O threads cannot race withtrue_probrestoration) and updatedahead()to protect top-3 unperturbed candidates (ranks 0..2) and readsmap. At a uniform3.40 GBarena on the 43-token SQLite WAL prompt,MOCE=0.85cuts misses by 51.5% vsQ2_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% vsIQ3_S (MOCE=0.0)(108.9 -> 40.9 misses/tok), reaching320.58 / 353.49 tok/sprefill (3,194.19 / 5,793.65 ms) and9.51 -> 13.52 tok/s43-token decode (16.11 mswait +56.29 msGPU) in the earlierv3pass (audits/opus_audit_phase0_2_m5_v3.md) and229.99 / 268.23 tok/sprefill and8.04–9.33 tok/sdecode in the laterresults/phase0_phase2_m5_validation.jsonpass ([hypothesis: passive SoC thermal throttling + concurrent background CPU load]). Across 192 teacher-forced assistant tokens (n = 64final assistant tokens × 1 complete untruncated turn/domain aftern_prompt = 244 / 159 / 456prompt tokens indata/m5_val_prompts/eval_tf_{ops_tools,chat_json,ja}.txt), isolatedMOCE=0.85vsQ2_0 (MOCE=0.0)adds only0.0588macro mean KLD (96.88%top-1 agreement:100.0%onchat_json,98.44%onja,92.19%onops_tools), while unhealedIQ3_S/IQ4_NL -> Q2_0expert conversion costs0.2032macro mean KLD (90.63%top-1:100.0%onchat_json,92.19%onja,79.69%onops_tools;0.2386combined withMOCE=0.85), motivating Phase 3's LoRA healing against the deployedQ2_0experts.
- 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_0expert quantization (heal_k256_q2_0 --quantize):GGML_TYPE_Q2_0(type42,18bytes per64weights) stores an IEEEfp16scaledwhose sign bit is unused by standard quantizers (d >= 0->{-d, 0, +d, +2d}), while all Strata MetalQ2_0kernels (units.metal,moe.metal,moe_mm.metal,mma.metal) multiply(code - 1) * dlinearly in signedhalf d. Unlockingd in Rgives 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) with0extra bits and0shader changes.audits/opus_audit_phase3_local_v1.mdcaught that weightingW_gate/W_upbyblk.l.hc_ffn_norm.weighthad-0.03mean log-correlation with the actual expert inputE[mixed_d^2](becausegemm_hc_up_64_64_4mixes the 4 hyper-connection streams afterhc_ffn_norm) andW_downhadcol_imp = nullptr. Upgraded--quantize(v2) to weightW_gate/W_upby tempered real activation RMS(E[mixed_d^2])^0.5(beta = 0.50) andW_downby per-expert analytical fourth-root importance(sigma_g * sigma_u)^0.5 ~ (E[h_c^2])^0.25(beta = 0.25onW_downvsbeta = 0.50onW_gate/W_up) over the 1,022-token multi-domain calibration trace (data/m5_val_prompts/cal_disjoint_1024tok.txt, verified016-gram overlap on scored spans across all 22 unseen windows and 3 canary windows; 18/22 unseen windows also have0full-window context overlap, while 4 windows share prompt context:<tools>boilerplate incj20_tool_5turn_t3/cj20_t3_answer_tool/cj03_toolcalland another turn ofpersonaconversation 3 inpersona3_t4). - Real-trace per-layer
alpha_lscale calibration + shared-expert ridge solver (heal_k256_q2_0 --calibrate-and-heal) & byte-level GGUF transplant:audits/opus_audit_phase3_local_v1.mdshowed that feeding synthetic Gaussian noisex = z * sinto Part A overshotalpha_lon 22/48 layers (especially deep layersL24–L44where real activations want0.97–1.04vs Gaussian1.08–1.12). Upgraded--calibrate-and-heal(v2) to compute exact parameter-free real-trace MSE scalesalpha_l = <y_ref, y_q2> / ||y_q2||_2^2(gamma = 0.00,L0 = 1.308,L24 = 1.022,L47 = 1.121vs unseen held-out oracleL0 = 1.337,L24 = 1.021,L47 = 1.134, reducing layers wherealpha_lis worse than1.0on unseen data from22/48to2/48and unseen routed residual from24.13%to23.37%), solve ridge regression (46/48layers full-rank640/640;L0rank430,L2rank617;eta_shd = 0.15,ridge_rel = 6.0) onblk.0..47.ffn_down_shexp.weight(in-sampleMoE MSE 19,061.40 -> 18,665.36 -> 18,460.89,-3.15%total:-2.08%fromalpha_l+-1.10%from ridge afterIQ4_NLrequant;-0.9%LOBO CV,-0.46%on unseen dump), and transplant only the48modifiedffn_down_shexptensors intoK256-q2_0-healed-metal.gguf(1175 / 1223other tensors bit-identical), with 4-uint64 FNV-1a cache fingerprinting (yr_q2_bin.metaoverd_bak+L0/L47 .mixed+L0/L24/L47 gate/up+t_cal),.quant.meta, and strictpread_exact/pwrite_exactI/O checks. - Deterministic
smap_synchronization instream.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 onsmap_instream.cpp(publishingsmap_only when!x.preor at claim time insideat_layer(), and evicting unchosen speculative read-aheads whenmoce_bonus_ > 0.0f). Across the full 22-window (2686-token) 4-stratum unseen KLD suite (results/phase3_22win_kld_audit_v2.json),v2healedQ2_0beats unhealedQ2_0by-28.55%[95% CI-33.5%,-23.2%] pooled KLD atMOCE=0.0(0.2029 -> 0.1450,20/22wins,+2.27 pttop-1:85.22% -> 87.49%) and-26.98%[-32.0%,-22.1%] atMOCE=0.85(0.2123 -> 0.1550,21/22wins,+1.53 pttop-1:85.55% -> 87.08%). Comparingv2vsv1across all 22 windows gives-5.08%[-10.8%,-0.1%] atMOCE=0.0(-3.26%atMOCE=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), withops(0.1199 -> 0.1077) andchat_json(0.1243 -> 0.1130) improving,personaflat (0.2637 -> 0.2527),jaslightly higher vsv1(0.1723 -> 0.1796, top-185.0% -> 83.7%, still-24.3%vs unhealed0.2374), andcj20_tool_5turn_t3moving from0.1783(v1) to0.1444(tied with unhealed0.1440atMOCE=0.0,0.1417vs0.1534atMOCE=0.85, aided by in-distribution<tools>context). On the 192-token 3-domain canary set,v2macro KLD is0.1563atMOCE=0.0(-23.08%vs unhealed0.2032, worse thanv1's0.1320) and0.1559atMOCE=0.85(-28.57%vs unhealed0.2182, vsv1's0.1365, with isolated MoCE-57.30%:0.0616 -> 0.0263). Across the 9-prompt diverse speed suite onv2(results/phase3_9prompt_speed_audit_v2.json),MOCE=0.85is9/9bit-deterministic across reps with92.96%mean combined hit rate (90.88%–95.60%,33.8mean misses/tok vs85.12%atMOCE=0.0;v1was93.27%, withpersona_458moving93.04% -> 91.26%), and0.0 MiBswapouts across all runs.
- Phase 4 (
#97) WFLD Layer-47 anti-slop + Strata Metalv3(#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):v1used a bag-of-n-grams proxy forL0..39and summedL40..46perturbations without passing them through the 7 intervening QSA/GDN/MoE layers (ops_toolsKLD jumped to0.276).v2dumped exact post-Layer-47 states (R_post47,inj47,hh47,sgv47) fromstrata_window_cli, eliminating intermediate drift (47/48tensors bit-identical to Phase 3, only updatingblk.47.ffn_down_shexp.weightinQ8_0), but itsmined_*loop penalized the 6 final words without checkingtuple[:-1].v3/v3.1enforced strict prefix matching and per-word collateral covarianceG_b, butaudits/opus_strata_metal_opts_v3.md§6.1 caught that settingdelta_cho = 0without constraining<|im_end|>(248046) caused suppressed probability mass to shift into<|im_end|>(28non-slop and7/21slop argmax flips to<|im_end|>on multi-turn chat).v3.2solves 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(0flips to special tokens across all 4 traces,56/56real turn ends preserved,non-slop chat KLD = 0.00472 <= 0.005,ordinary Δlog p = +0.0208 >= -0.01, all 6 collateral words>= -0.042with13/142 = 9.15% <= 10.0%below-0.1,-69.7%/-70.0%/-84.1%slop reduction on dev separate-doc / same-doc / train, and0.003193-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,0assistant-side turn damage,-54.4%weight-only slop reduction (0.554 -> 0.253) and-98.98%(0.554 -> 0.0056,21/21argmax clean) when stacked with the 0-overhead prefix-conditioned logit penalty (STRATA_ANTI_SLOP_PENALTY=6.0) inWindow::finish). - Strata Metal
v1–v3runtime 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=1at7.0 GBarena (0pageins/decompressions,0.0 MiBswapouts); (2)STRATA_METAL_SHEXP_OVERLAP=1+STRATA_METAL_SPLIT_TURN=1+STRATA_METAL_EXPERT_SPLIT=1(resident-first 2-phasemoe_gu/moe_downoverlap withMTLSharedEventSSD handshake, shrinking CB–kernel gap to< 1.0 msand cutting decode from58.72 ms/tok(17.03 tok/s) to45.28 ms/tok(22.08 tok/s,+29.7%,97.59%combined hit rate,11.6 misses/tok,5.08 msSSD wait);22.25 tok/s/44.95 ms/tokwithPF2=2, IO=12); (3)STRATA_METAL_PRED_Q8=1(Q8_0lookahead router saving92 MB/tokDRAM reads),STRATA_METAL_ROUTER_BF16=1(losslessBF16router halving router reads126 -> 63 MB/tok),STRATA_METAL_MOE_DN_TILE=8,2,4(-0.25 ms/tokon M5),STRATA_METAL_PREWARM=1(eliminates282.8 msT=3speculative draft shader compile stall), andSTRATA_HC_Q4/STRATA_HC_Q5(Q5_0/Q4_0hyper-connection GEMVs saving240–315 MBwired RAM).
- 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_saveonly saved the112.60 MiBrunning state (36GDNgst_+ghist_,pleHist_,tail_,dead_), omitting QSAkq/vq/ksc/vsc,pooled, andss.ple_prev.session_save/loadserializes all118,069,248fixed bytes +14,208 B/cell(139.46 MiBat1,983cells) directly from/intoMTLBufferunified memory via 4-threadpread/pwritewithF_NOCACHE+fsync, loading in28.93 mscold NVMe (5.18 GB/s,281xfaster than8,138 mscold prefill) and5.00 mswarm (1,627xfaster) at0.000000KLD (100.00%top-1). - Suffix Cache Reuse (
SCR) on 36-Layer GDN + 12-Layer QSA (Window::scr_relocate): inqsa.metal, only group 0 (dims 0..63,ksc[.][.][0]) ofkq(int8) andpooled(f32) has NEOX RoPE applied (dims 64..255,ksc[1..3],vq,vscmove as exact bytes). Enforcing $\Delta \equiv 0 \pmod 4$ (paddingB') relocates1,535tail tokens across all 12 QSA layers in0.5 mson CPU. Keeping the post-edit 36-layer GDN state (S_ABC,keep) gives+0.0038nats/tok (C=511, wherekeep J=8lowers PPL7.49 -> 7.20) to+0.038nats/tok (C=1535,3.48x–3.79xprefill speedup) of isolated GDN staleness (keep - gdnex) and3.20%repeated 4-grams (0% looping) over 128 greedy tokens, whereas droppingC's GDN history (lossy) explodes KLD to1.41–2.33(PPL 29.5–66.4). Across all1,728GDN value heads (36 x 48), median gate half-life is18tokens (76%decay within100tokens;104long-memory heads> 1,000tokens). - 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) caused14/70step0special-token argmaxes; fixing</think>\n\n+ theWindow::finishstep0guard (p1 == 248069 && (p0 == 271 || p0 == 198)) reduced special-tokenstep0argmaxes to0 / 190. (2) Full-rowoutput.weightupdates (2.9M–7.5Mparams) 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.6Kparams) on the 3-way split (113train /40dev /37test,17,352whole-pair rows) achieved+0.196nats/tok ontest(blind chat KLD 0.0033,0flips), buttest/trainratio was0.38 < 0.70becauseidx % 5kaomoji 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 -> 256experts and re-quantizing scalarQ2_0, streamingukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF's imatrix/GSQ-RCO vector-quantizedIQ2_XSacross all512experts (3.90 GBdense GGUF +36.23 GB.experts, replacing the 3IQ1_Mlayersblk.8,13,37withIQ3_XXS's nativeIQ3_S/IQ2_XXStensors and skipping the28.80 GBduplicate PLE table) runs at21.00 tok/s(47.61 ms/tok,94.38%hit rate,7.97 GBresident footprint,0.0 MiBswap) on the 24 GB M5 Air with 100% of all 512 experts andukisai's full imatrix quality; settingreasoning_effort|default('low')intokenizer/chat_template.jinjakeeps<think>concise (~50–100 tokens) without any hard token cutoff.
-
Anti-Slop v4.1b (
slop4per-position stateful shaper,audits/opus_antislop_v4.md):- Layer B (linear WFLD on
blk.47.ffn_down_shexpinIQ4_NL) was rejected by falsifier: ridge regression on 640-dz = silu(sgv)*hhleaked0.28–0.68(1–4 nats of bullet/#penalty) into code fences and<think>blocks. - Layer A (
slop4::State, 12 B/pos inwindow.cpp, scanned viaslop_scanon prompts/restores and relocated viascr_relocate) gates penalties strictly onPROSE && !TOOL && !FENCE(0penalties and0.000000KLD across all14,090<think>,<tool_call>, and```code-fence rows;< 0.005KLD on ordinary prose). - Empirical verification of the committed-prefix synonym treadmill vs clause-entry branch points (
arXiv:2510.15061Antislop,arXiv:2604.19139v2Verbal 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 mass0.161 -> 0.005(-97%) so the model branches into content ("If I'm reading this right"). Includesslop4::hold()so penalizing a high-probability token never boosts<|im_end|>.
- Layer B (linear WFLD on
-
M5 GPU Decode & Prefill Optimizations (
audits/argon_k512_decode_opt.md,dense_mm.metal,prompt.metal,stream.cpp,window.cpp):- Decode (
ntok=1): M5-tunedWindow::gemvtiles (IQ4_XSoutput.weight/attn_qkv/attn_gate->32,2,4,1at130–138 GB/s,ssm_out/attn_outIQ4_XS->32,2,4,1eliminating split-K barrier,IQ3_S6144x2560->4,2,4,1,shexpIQ4_XS 2560x640->16,2,4,13.96xfaster,predQ8_0->32,2,4,1),gr_down_q8_8_4(158 GB/s),IQ2_XXSmoe_gugl=16,IQ3_SQUANT_CTAB=0(118.4 GB/s),make_arenacontinuefix (3,726 -> 4,326wired experts at 6.8 GB), andSTRATA_METAL_PREFETCH_UNPERT=1default (96.90%combined hit+prefetch rate,21.00 -> 22.64 tok/sdecode on unprunedK=512IQ2_XS). - Prefill (
prompt_chunk): Addedgemm_f16_128_64_4(-605 ms/-20%chunk GPU time),in_prompt_single-piecepread+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 inhold_region(true)), andSTRATA_METAL_PROMPT_MOCE=320435.0(Tail=32 exact + Body top-4 protected), cutting 244-token prefill from12,493.9 ms -> 2,371.1 ms(5.27xfaster,95.31%top-1 agreement) and 1,024-token prefill from7,642.9 ms -> 4,067.1 ms(+87.9%faster).
- Decode (
-
Cross-Session Convergence with
01a11408(regent@astaroth):- Ported
01a11408'sinclude/strata/core/scr.hpp(scr_carry+ automatic0..3newline pad insertion so any mid-context edit satisfies $\Delta \equiv 0 \pmod 4$ +<path>.padspersistence) intostrata_window_cli.cpp--serve, and sent ourchat_template.jinjadefault('low'),antislop_v4.inc, andlayer_rank_/PROMPT_MOCEfindings back to01a11408viaalbedo send.
- Ported
-
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 mstotal GPU across 47 layers) and queueing predicted non-resident experts most-tokens-first beforedone_cv_.waitinstream.cppcuts synchronous SSD misses on a 1,024-token chunk from34.3% -> 9.5%(streamed-ahead50.9% -> 75.5–79.0%, exposed SSD wait1.4s -> 0.39–0.69s), while runningss.ple.table->gather_batchon a background thread joined at cut 0 hides~150 msof PLE SSD reads. - All 48 layers'
STRATA_METAL_DUMP_MOEtensors (mixed,ids,wts,yr,hh,sgv,R_post47) are 100% byte-identical to baseline (0.000000prefill KLD). - Setting
STRATA_METAL_PROMPT_CHUNK=2048(withSTRATA_METAL_ARENA_GBpinned) cuts total experts read on a 2,048-token prompt by-43%(32.6K -> 18.3K) and boosts 100% bit-exact cold prefill to335–346 tok/s(5,919 mschunk for2,048tokens;276.7–286.6 tok/s/3,440–3,570 mschunk on1,024tokens).
- In
-
Long-Context Decode Scaling Out to
64K(65,536) and128K(131,072) on 24 GB M5 Air (K=512UnprunedIQ2_XS):- Because 36/48 layers are GDN (
112.2 MiBfixed recurrent state,0 B/tokgrowth) and 12/48 layers are QSA (int8KV +128-dim 4-token pooled lightning indexer selectingtop-2048tokens), total session state is only1.00 GiBat64K(9.0 GBtotal resident RAM) and1.84 GiBat128K(10.3 GBtotal resident RAM), decoding at21.13 tok/s(16K),20.27 tok/s(32K),20.35 tok/s(64K,49.15 ms/tok), and20.75 tok/s(128K,48.20 ms/tok,97.04%hit+prefetch rate).
- Because 36/48 layers are GDN (
-
Three
gpt-6.1-solAudits (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): Verifiedio_mainnever publishes speculativepreslots intosmap_, and testedSTRATA_METAL_PREFETCH_UNPERT=0vs1(call_2343482):UNPERT=0halves 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-layerpredlogits drift across layer $l+1$'s token mixer; keptUNPERT=1. - QSA-Boundary-Safe Suffix Speculative Decoding (
sol_spec_mtp_decode.md,scripts/suffix_window.hpp): Caught an upstreamqsa.metal(qsa_prep) rollback bug wheretail[4][128]wraps modulo 4 across 4-cell block boundaries on partially rejected speculative windows (T > 1), corruptingpooled_when remaining cells of the 4-cell block are re-run; capping speculative verification steps atT <= 4 - (pos & 3)(suffix_limitinsuffix_window.hpp) eliminates all 736 unrestricted rollback failures with 0 overhead (20.62 -> 21.47 tok/soncj20_tool_5turn_t3). - Post-Norm Persona Steering (
sol_persona_steering_arxiv.md): Proved that any post-output_hcadditive 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-pair113/40/37split that ridge-whitened readout contrast fails the0.70generalization bar (dev/train = 0.514, test/train = 0.337); kept weights untouched.
- Decode Prefetch (
-
Native 49th-Layer Neural MTP (
models/K512_iq2_xs/mtp-q2_0.gguf,889 MB) + Unrestricted QSAtail_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) fromQwen/Qwen3.8-Flash-Nextvia 24-worker HTTP Range requests and packedmodels/K512_iq2_xs/mtp-q2_0.gguf(889 MB,10,629-tokendraft_vocab.bin, single-command-buffermtp_chain.metal). - Added 24 KiB
tail_bak_snapshot/restore inWindow::run/Window::commit(window.cpp) soT <= 4speculative verification rolls back across 4-cell QSA block boundaries with zero restriction, and fixedWindow::mtp_draftso--spec-min-pgatesj = 0as well asj >= 1(Adaptive T in {1, 2, 3}). - 4-Domain Decode Benchmark (
192generated tokens,K=512UnprunedIQ2_XSon 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).
- Coding (
- Wired
strata::core::scr_plan(min_span=32, seed=16, block=4) + bounded 2-slot MRUck_ringintoserve_loopso multi-turn chats that strip prior turns'<think>...</think>blocks ([A B C] -> [A C D],arXiv:2609.37725Appendix B) automatically SCR-relocate the assistant replyC(cached_tokens: 414 / 446=92.8%reused on Turn 2!). - Slim MTP +
AcceptMoEDraft Routing (router.metal,moe.metal,stream.cpp,window.cpp,call_2008636&call_1884893):madvise(MADV_DONTNEED)on the181 MBof source BF16 dense matrices inmtp-q2_0.ggufafter quantizing them intomq8_(90 MBQ8_0) + switchingdraft_vocabfrom106,299tokens (145 MB) todraft_vocab_en.bin(40,525tokens,55 MB) saves271 MBof RAM, allowingSTRATA_METAL_ARENA_GB=6.5(9.9 GBresident RSS,-1.0 GBlower than baseline!).- Capping
mtp_encode's local QSA sliding window to256cells (4chunks of 64 instead of33) and runningmtp_prefillonly on the last256tokens of the final prompt chunk eliminates the multi-chunkmtp_prefilllag on12K–32Kprompts (314–330 tok/send-to-end 2K prefill!). - Setting
STRATA_METAL_PREFETCH_DRAFT=0instream.cpp(ahead()only prefetches for root tokent=0) andAcceptMoE(arXiv:2608.02989) draft routing inrouter.metal/moe.metal(sg > 0protectsr_prot = 1and squaresmoce_scale) cuts speculative SSD traffic from332 -> 177 MB/window(94.1%arena hit rate!) and pushes 32K-context decode to30.37 tok/s(persona_chat),33.44 tok/s(reasoning),34.41 tok/s(coding), and42.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):- MTP
STRATA_METAL_MTP_SINK=64+mtp_wroteValid-Run KV Fix + Batched 4-Rowfc_hiddenGEMV +slop4-Awaremtp_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, batchedfc_hidden(-20.9 MB/step), and syncedslop4state into the MTP draft head + wiredSuffixDrafter::propose_afterto extend partial neural drafts. - 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,247Prefill Barriers,P_GSLOTL2 Fix): Fusedssm_alpha + ssm_beta(gemv_cat) andmoe_combine_sh + gr_norm(gr_norm_comb,-83dispatches &-48barriers/window), optimizedgdn_step(-12%atntok=4,simd_broadcast) andgdn_scanr_4(-8.4%atT=2048), parallelizedqsa_prep, fixed the Apple M5 L2 cache-coherency race inEXPERT_SPLIT=1viaP_GSLOT(100%bitwise deterministic across all 48 layers), removed1,247redundant prefill barriers, and added M5 shape-awaregemm_f16tiles +publish_readygated onprompt. - 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_place32-subblock +mm_scanpad-clear,mm_gather_h): Foldedmoe_outinto per-token 4-streamgr_norm_h_moe(-25%), fusedple_gnorm+ple_gate_norm(-42.5%, freeing188.8 MBof prompt scratch RAM (2.35 -> 2.16 GB)), fusedmtp_rms_xprep+mtp_combine_gnorm(mtp_msdown to218 ms/2.25 ms/draft), eliminatedmm_clearand replacedmm_place's $O(1024^2)$ loop with 32-subblock nibble counting (12.9xfastermm_place,5.89xfaster MoE plan), and added FP16mm_gather_h(1.70xfaster gather,-40.2 ms/chunk). - Round 5 Prefill + Decode + Scratch Aliasing + Deterministic Lookahead (
gr_norm_h_moec,gr_apply_moec,gr_norm_h4,gr_normrv[10],gr_norm_comb_4s,qsa_prep#qf16,-580 MBScratch Aliasing,stream.cppLookahead Fix): Foldedmm_combineintogr_norm_h_moecandgr_apply_moec(eliminating47/48mm_combinedispatches and2.01 GBDRAM traffic/chunk), added per-token 4-streamgr_norm_h4(-18.3%), skipped deadp.mixedFP32 writes ingemm_hc_up(-11.3%), cachedgr_normin registers + added 4-streamgr_norm_comb_4s, fused FP16 query output intoqsa_prep#qf16(1.98xfaster,-8.91 ms/chunk), aliased disjoint GDN/QSA prompt scratch buffers (cuttingMT=2048transient prompt scratch from2.16 GB -> 1.58 GB,-770 MBvs original2.35 GB), and fixedstream.cppat_layerto deterministically reclaim unused prompt lookahead slots. - Round 6 QSA Flash + PLE + GDN + Register-Cached
gr_norm_h*+-286 MBPLE/bo/yrScratch Aliasing (qsa_flash,ple_moe_gate_norm_c,ple_conv_res,gdn_pre,gdn_post,gr_norm_h*): Cachedhalf4ksc/vscscales in threadgroup memory and parallelizedqsa_flashonline softmax across all 4 simdgroups (8lanes/head,-20.0%atT=1984:34.36 -> 27.50 ms/layer,-26.6%atpos0=30000, T=96,memcmp == 0) + skippedqsa_topkwhenp.flash && pos0 + T <= 2051; eliminatedgatedDRAM writes/reads everywhere (-167.8 MB/chunk+-40 MBdecode window RAM), added 1-TG-per-tokenple_moe_gate_norm_c(-28.3%, eliminating the lastmm_combinedispatch at Layer 0), and aliased PLE +bo/yrprompt scratch (cuttingMT=2048transient prompt scratch from1.58 GB -> 1.29 GB,-1.06 GB/-45.1%vs original2.35 GB); cachedrvin registers acrossgr_norm_h*(-7.89 GBDRAM reads/chunk) and refactoredgdn_pre(-16.1%,0barriers) andgdn_post(0barriers) to 1 simdgroup per head. - Round 7 Single-Pass
simd_broadcastgdn_scanr_4(-22.2%) + Cooperative-Registergemm_hc_up(-8.9%) +ple_hist_adv&-159 MBScratch Aliasing (1.13 GBScratch): ReplacedTB=16threadgroup staging andfloat4 kc[8](32VGPRs/lane) ingdn_scan_rwith single-passTB=0scalarsimd_broadcast(23.8 KiB -> 1 KiBTG memory,-256barriers/layer,-127wasted state passes,-22.2%/-55.6 msacross 36 layers atT=2048); added cooperative-destination register epilogue togemm_hc_up(16 KiB -> 0 BTG memory,64x64< 1024,64x128>= 1024,-8.9%unbound /-14.0%bound); addedple_hist_advGPU kernel and aliasedp.normW = p.XS,p.HS = p.qkv,p.gsh = p.gqn,p.ush = p.gkn,p.hh = p.ygh,p.xqh = p.z(cuttingMT=2048transient prompt scratch from1.29 GB -> 1.13 GB,-1.22 GB/-51.9%vs original2.35 GB, and0.708 GBatC=1024). - Round 8
prompt_chunksigs/cutsShared-Expert + MoE-Plan Overlap +W8<IQ2_XS/IQ2_XXS>+gr_norm_h_stream_out+dq_f16/act_f16half4(audits/opus_r8_moemm.md,audits/sol_r8_grvec.md,audits/argon_r8_dq.md): Signaledsigsright afterroute_topkinWindow::prompt_chunksomm_hist,mm_scan,mm_place,mm_gather_h, and the 4 shared-expert GEMMs +swiglu_hrun on the GPU concurrently beside the host/SSDat_cut(l)(5,165 -> 4,837 ms=-328 ms/410.0 tok/sinsideprompt_chunk!), foldedW8<IQ2_XS>scale nibbles (-2.2..-3.2%) andW8<IQ2_XXS>popcountsigns (-1.8..-2.3%), reducedgr_norm_h4/gr_norm_h_moecstream-by-stream (-30VGPRs/lane,-7.5 msinkprof), and vectorizeddq_f16(half4stores, barrier-free whenTAB == 1) andact_f16(float4 -> half4,4xfewer threadgroups)! - Round 9
qsa_flash19,904 -> 7,616 BThreadgroup Memory (4TGs/Core,-17.5%/-51.4 ms) +TAB > 1Barrier Guards +IQ3_SDead Sign-Load Removal +dn_bodyNE==1(audits/opus_r9_qsa.md,audits/sol_r9_moemm.md,audits/argon_r9_quant.md): Shrunkqsa_flashTG memory from19,904 B -> 7,616 Bvia 2-round 12-rows.storeintosp[2][HPG*QCH], register-combined((s0+s1)+s23)*0.0625f, and aliasingp(fp16) into deadsp[1](unlocking4TGs/core on M5, cuttingpg:qsa_flash x12inkproffrom332.81 -> 281.41 ms=-51.41 ms/-15.4%, and-21.6%atpos0=30000), guardedload_tab+threadgroup_barrierwithfmt<F>::TAB > 1inmm_gu3,mm_down3,moe_gu, andmoe_down, removed deadload_signsfromload_tab<T_IQ3_S>, and addedNE == 1hoisting indn_body! - Round 10
qsa_flash2.2x(-51.9%/-146.1 msinkprof) Unconditional Hoistedrv/rk+vec4Epilogue +gdn_prefloat4Conv +uint4mm_gather_h+s_id-Freemm_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 inqsa_flashrvandrkcooperative tensor fills and made the 64 byte/char4loads unconditional (max(id, 0)with scale0.0fon empty cells) +vec4float4/half4output epilogue, cuttingpg:qsa_flash x12in full-modelkprofatT=1983from281.41 ms -> 135.36 ms(-146.05 ms/-51.9%vs Round 9, and-396.66 ms/3.93xfaster than Round 5's532.02 ms!) and cutting total full-modelkprofGPU time from3,693.1 -> 3,548.4 ms(-144.7 ms); vectorized value-headfloat4conv ingdn_pre(-2.0..-2.9%), upgradedmm_gather_hto 128-bituint4, and eliminatedthreadgroup short s_id[1024](2 KiB -> 0 B) inmm_place(-6.2%faster MoE plan)! - Round 11
gemm_f16Tile Selector (-18%/-234 ms) +qsa_attn_mmHoistedrv/rk(-44.5%) + BranchlessW8<T_IQ3_S>(-50.8%/2.03x) &mm_gu3/mm_down3Pointer/Block Hoisting (-57.4 msinkprof) (audits/opus_r11_dense.md,audits/sol_r11_qsa.md,audits/argon_r11_moemm.md): FixedWindow::prompt_chunk'sgemm_f16tile selector (M > 1536instead ofM >= 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 hoistedrk/rvfill toqsa_attn_mm(70.5 -> 39.2 ms,-44.5%/1.80xfaster), replaced divergenthhbranches inW8<T_IQ3_S>with branchless byte-shift indexing (37.2 -> 18.3 ms/layer,-50.8%/2.03xfaster) + hoisted invariant row pointers (g0_l,u0_l,w0_l) and super-blockboffinmm_gu3andmm_down3(-57.4 msacross 48 layers inkprof!), cutting total full-modelkprofGPU time atT=1983from3,548.4 ms -> 3,397.5 ms(-150.9 msvs Round 10,-828.2 msvs Round 5 =583.7 tok/spure GPU prefill throughput on M5!) and cutting10,180-token Albedo prompt prefill from28.25 s -> 24.59 s(360.3 -> 413.9 tok/s,-3.66 sfaster &-68.2 GBfewer SSD reads)! - Round 12
qsa_topk3.85x(-73.6%) +gemm_hc_up_reg32_128_4& Unrolled Epilogue (-6.9%/-7.8 ms/chunk) +T_Q6_KHoistedmm_gu3(-4.3%/-10.2 ms) +3,375.4 msGPUkprof(audits/opus_r12_moemm.md,audits/sol_r12_qsa.md,audits/argon_r12_gemm.md): Cached 4 order keys per thread inqsa_topk(nb <= 4096), replaced thetid == 0256-bucket serial loop with a parallel simdgroup scan, and fused the(equal, greater)emit scan (4 -> 2barriers per 1,024 blocks,2.55x–3.85xfasterqsa_topk:1.861 -> 0.491 msatT=1984, pos0=2048,5.543 -> 2.157 msatpos0=30000,memcmp == 0) + boundedqsa_flash's PV loop byc0 < nc; switchedgemm_hc_up_regto32_128_4with unrolled pointer-hoisted epilogue (-6.9%/-7.8 ms/chunk,116.5 -> 108.8 ms); and enabled the pointer-hoisted fast path forT_Q6_Kinmm_gu3(-4.3%atT=1983,-6.3%atT=512), cutting full-modelkprofGPU time atT=1983to3,375.4 ms(587.5 tok/spure GPU throughput!). - Round 13
gemm_hc_down+gemm_swiglu_upK4 Register Epilogues (-144Dispatches,-985 MBDRAM Traffic/Chunk,600.6 tok/sGPUkprof) +dq_f16Exact-Simdgroup 2D Grid +SpecGate&mtp_draftPLE Prefetch (audits/opus_r13_hcdown.md,audits/sol_r13_specgate.md,audits/argon_r13_dqflat.md,audits/opus_colibri_scout.md): Fusedgr_epi_downintogemm_hc_down's cooperative-tensor register epilogue (-96dispatches,-96barriers,-493.4 MBp.VDRAM writes/reads,-2.65 MBp.Vscratch,-13..-27%faster atT <= 192,-4..-5%atT = 384..512) and fusedswiglu_hinto single-accumulatorgemm_swiglu_up(-48dispatches,-491.5 MBp.ushDRAM writes/reads,-11.5%/-3.2 ms/chunkatT=1024,-6..-12%atT=256..512,-2.8%atT=2048), cutting full-modelkprofGPU time atT=1983from3,375.4 ms -> 3,301.86 ms(600.6 tok/spure GPU prefill throughput on M5,-73.5 msvs Round 12 and-923.8 msvs Round 5,2,298 -> 2,154dispatches!); upgradeddq_f16to hardware[[thread_position_in_grid]]with exact-simdgroup 2D threadgroup sizing (128/160/192/256threads,1.12x–1.27xfaster on small/medium matrices) +T_Q5_1instantiation; and addedSpecGate(STRATA_SPEC_GATE=suffix|1) +mtp_drafttokPLE prefetch during the2.1 msMTP GPU run! - Round 14
qsa_flash/qsa_attn_mmByte-Offset & Zero-Scale Hoist (-8.7%..-9.7%across allpos0,443.2 tok/son10.18KAlbedo Prompt) +gdn_scanr_4hk-Locality + Compactedmoe_comb_slotsroff+mm_gu3IQ2_XShalf4HSStore (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 zeroedks_tg/vs_tgscale slots once per cell before theh = 0..3head loop inqsa_flashandqsa_attn_mm(-8.7%atpos0=0:10.33 -> 9.42 ms,-9.7%atpos0=2048:21.80 -> 19.68 ms,-7.3%atpos0=8192:26.82 -> 24.86 ms,-7.6%atpos0=30000:32.62 -> 30.13 ms, cutting full-modelkprofpg:qsa_flash x12from128.03 -> 121.93 ms=4.36xfaster than Round 5's532.02 ms!); remappedgdn_scanr_4tohk-grouped head scheduling (-2.18 msacross 36 layers atT=2048), compactedmoe_comb_slots(roff[0..cnt-1]viasimd_prefix_exclusive_sum) + unrolled stream loops ingr_norm_h4/gr_norm_h_moec(-2.4..-3.5%ongr_norm_h4,-0.9%ongr_norm_h_moec), and addedhalf4HSstores tomm_gu3(T_IQ2_XS,-0.5..-0.6%), pushingprompt_albedo_11k(10,180tokens) prefill to443.22 tok/s(22.97 svs28.25 s,-5.28 s/+23.0%faster, with Chunk 1 hitting501.5 tok/s/4,084 msend-to-end!). - Round 15
prompt_chunkRouter+Plan/Shared-Expert Wave Zipping (-190Barriers) + FusedWal+WbeN=96GEMM (-36Dispatches) + QSA PSO Pre-Warm (-202 msChunk-1 Stall) +qsa_score_mmCausal-Future Tile Skip (-19.4%) +qsa_prep320 BTG &2Barriers (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),-94barriers) and zipped MoE Plan (mm_hist/mm_scan/mm_place) with Shared-Expert GEMMs (Wshg+Wsg/gemm_swiglu_up/Wshd,-96barriers,100%overlap ofmm_scan/mm_placewith shared-expert GEMMs); fusedssm_alpha + ssm_beta(N = 48 + 48 = 96 = 3 * 32,-36dispatches,2,154 -> 2,118dispatches/chunk,-378 MB/chunkp.mixhreads); pre-warmedqsa_score_mm/qsa_topkPSOs inprompt_gemm_init(cutting Chunk 1graph buildfrom203 ms -> 1 mson10.18KAlbedo prompt!); added causal-future tile early-exit (uint(b0) >= max_bid) +dead[0..3]hoist toqsa_score_mm(-19.4%fasterqsa_score_mmand-14.1%combinedqsa_score_mm + qsa_topkatT=1984, pos0=2048) and removed the 3rd radix barrier inqsa_topk; shrankqsa_preprow[256] -> row[64](1,088 -> 320 BTG memory) + cut single-call barriers from4 -> 2(-1.0..-1.9%), and unrolledroute_topkforn_exp == 512(-1.2..-2.2%), pushingpg:qsa_flash x12inkprofdown to116.58 ms(9.71 ms/layer,4.56xfaster than Round 5's532.02 ms!) andprompt_albedo_11kto443.43 tok/s(22.96 s)! - Round 16 Fused QSA
Wk+Wv+Wik(+Wiq)Single GEMM (-24..-36Dispatches,-5.6 MBScratch ->1.12 GB) + Decode Resident-MoE & Shared-Expert Wave Zip (-48Barriers/Window) +gemv.metalTAB>1Guards +ple_moe_gate_norm_cSIMD Compaction &ple_convhalf4/t>=9(audits/opus_r16_kvik.md,audits/sol_r16_dec.md,audits/argon_r16_ple.md): Fused QSAWk(512),Wv(512),Wik(128), andWiq(512) into a singleN = 1152(9 * 128) orN = 1664(13 * 128) GEMM inp.Wk/p.kcurwith zero-overhead function constantQSA_KV_STRIDEinqsa_prep+ PSO pre-warming (eliminating24–36GEMM dispatches per chunk ->2,094dispatches, and freeingp.iqr/p.ikr/p.Wik=-5.6 MBprompt scratch ->1.12 GB); zipped decodeWindow::layer's resident routed experts (phase = 1) with the shared expert insideshx(4 -> 3waves,-48barriers/window, startingmoe_gu(phase=1)immediately atsigsand hidingswiglu_prepbesidemoe_down(phase=1)) + guardedload_tab+threadgroup_barrieronfmt<F>::TAB > 1ingemv/gemv_cat/gemv_pair; and compactedple_moe_gate_norm_crouted slots (-6.8..-10.3%) + vectorizedple_convhalf4weights with a branchlesst >= 9fast path (-4.7..-5.8%full PLE stage) + hoistedgdn_postrow pointers (-0.9..-1.8%)! - Round 17
mtp.metalfloat4Vectorization (-21%..-45%) &Q8_0gemv_catk+ Barrier-Freesimd_shuffle_downingr_up_f/gr_up3/qsa_merge(-97Barriers/Window,1,024 -> 0 Bqsa_mergeTG) +gdn_stepib==0Readout Pruning (-7.7%..-10.0%) &qsa_attnchar4KLoads (-6.7%..-11.4%) (audits/opus_r17_mtp.md,audits/sol_r17_shuf.md,audits/argon_r17_gdn.md): Vectorizedmtp_prepare/mtp_xprep/mtp_rms_xprep/mtp_combine_gnormwithfloat4loads/stores (-41..-45%onmtp_xprep,-21..-32%onmtp_rms_xprep,-9%onmtp_combine_gnorm) and addedQ8_0gemv_catkformtp_encode'sk_proj+v_projandgate_proj+up_proj; replaced threadgroups16reductions ingr_up_f,gr_up3, andqsa_mergewith barrier-freesimd_shuffle_down(-512 BTG &-1barrier ingr_up_f/gr_up3=-97barriers/window, and1,024 -> 0 BTG &0barriers inqsa_merge) + addedn == 2560unroll fast path ingr_norm(12–40%faster); and pruned redundantib = 1..3red2shared-memory reads + unrolledgdn_step(-7.7%..-10.0%faster multi-token verify,-3.8%..-4.4%faster 4-token commit) + vectorizedqsa_attnKloads withchar4+ hoistedqcurpointers (-6.7%..-11.4%faster across allpos0&ntok). - Round 18 Per-Stream
gr_norm_comb5.09x(-54%..-80%) +xprepfloat4(-41.5%..-46.8%) &moe_combine_sh(-32%..-49%) +Q6_Kmoe_guNE>=3(-18%..-25%) +qsa_topk2/qsa_score(-5%..-9%) (audits/opus_r18_moe.md,audits/sol_r18_grcomb.md,audits/argon_r18_qsa2.md): SwitchedWindow::gr_normfrom 1-TG-per-tokengr_norm_comb_4s(rv[4][10]register pressure + 100planloads/thread) to pointer-hoisted, unrolled 4-TG-per-tokengr_norm_comb(grid = (4, n_tok),yep[0..9]hoisted,54%–80%faster acrossntok = 1..8:22.67 -> 4.46 us=5.09xatntok=1); vectorizedxprep(units.metal) withfloat4stores (-41.5%..-46.8%:19.8 -> 10.5 us), unrolledmoe_combine_sh/moe_combinefortopk == 10(-45%..-49%atntok=1,-32%..-33%atntok=8), and fusedqg/qudecode per entry ingu_bodyforUNIT == 32(T_Q6_K-18%..-25%atNE >= 3); and ported register-cachedkeys[4](nb <= 4096withif (b >= nb) break;) +-7radix barriers (20 -> 13barriers =-84barriers/window) to decode'sqsa_topk2(qsa2.metal) + hoisted/unrolledqsa_score(-4.9%..-7.6%faster combinedqsa_score + qsa_topk2acrosspos0 = 3000..30000). - Round 19
gemv.metalNTOK/NR/KSUnroll (-5.4%..-29.4%atntok>=2) +mtp_pick/mtp_slop_penunroll_count(4)(-27..60%) +ple_comb_query_gate_norm& Load-Timewindow_initPrewarm (audits/opus_r19_prewarm.md,audits/sol_r19_sample.md,audits/argon_r19_gemv.md): Uncovered that becauseNTOK([[function_constant(0)]]) is a Metal function constant rather than a C++ template parameter, the Apple Metal compiler was NOT unrollingfor (int t = 0; t < NTOK; ++t)ingemv,gemv_pair,gemv_cat,gemv_catk, andgemv_x, forcingfloat acc[NR][8]to spill to thread-local stack memory on everyuiteration wheneverntok >= 2; adding#pragma clang loop unroll(full)acrossNTOK/NR/KSloops makesQ8_0GEMVs-5.4%..-23.7%faster,IQ3_S32_2_4_1GEMVs-7.7%..-20.8%faster,Q2_0 640->2560-15.6%..-29.4%faster, andgemv_cat/gemv_catk-6.6%..-15.8%faster atntok = 2, 4; unrolledmtp_pick(~27–39%faster) andmtp_slop_pen(~23–60%faster) +Q4_K/Q5_Kmtp_embed(~27–28%faster) inmtp_chain.metal; and fused decode Layer 0'sgr_norm_comb + ple_query_gate_normintople_comb_query_gate_norm(-12.5%..-22%,-1dispatch,-1barrier,-1dead RMSNorm) + register-cachedple_gnorm(-29%..-52%) + pre-warmedt = 1..max_tandMTPencode/pick PSOs inwindow_init! - Round 20
Window::mtp_load707.8 MBExpertmlock(mtp_ms398 -> 122.2 mson11K,433 -> 202.4 mson2K,40.57 tok/sDecode) + Decodeqsa_attn(-20%..-40.5%) &qsa_merge(-9.6%..-12.7%) +gdn_scanr_432-Bit Base Offsets (186.15 msinkprof) (audits/opus_r20_mtp_qsa.md,audits/sol_r20_gather.md,audits/argon_r20_gdn.md): Uncovered thatWindow::mtp_load(window.cpp) was notmlocking the707.8 MBofT_Q2_0routed experts inWm_(mtp-q2_0.gguf), allowing macOS to evict them during1.12 GBprompt-scratch prefills and causing+220..305 msof hard SSD page faults duringmtp_draft; wiring (mlocking) the707.8 MBof active MTP experts +MADV_DONTNEEDing181.2 MBof unused pages cutmtp_msfrom397.88 -> 122.18 ms(3.26xfaster!) on the10,180-token Albedo prompt and from432.82 -> 202.37 ms(2.14xfaster!) onprompt_2048tok_ops, pushing decode to40.57 tok/s(24.65 ms/tok,3,130.15 ms); register-cacheds0, s1across softmax + unrolled V x4 in decodeqsa_attn(-20.4%..-40.5%faster acrosspos0 = 512..8192) and sharedek[]/L_shinqsa_merge(-9.6%..-12.7%faster); and replaced 64-bit per-lane pointers ingdn_scan_rwith 32-bit base offsets (s0, k0, v0, ks, rs) + final-token dead-lookahead elimination, cuttingpg:gdn_scanr_4 x36inkprofatT=1983to186.15 ms(5.17 ms/layer,-61.4 ms/-24.8%vs Round 5)! - Round 21
gr.metalGR_NTOKUnroll (-5.4%..-18.0%across194Dispatches/Window) +qsa2.metal(-36.8%..-40.6%qsa_attn2&-8.3%..-14.6%qsa_merge2) +gdn_step32-Bitst_off&moe_plantopk10unroll_count(10)(-9.5%..-26.6%) (audits/opus_r21_gr.md,audits/sol_r21_qsa2.md,audits/argon_r21_gdn_moe.md): UnrolledGR_NTOK([[function_constant(1)]]) andKS/UNIT/hcloops ingr_down_f(GR_NTOK <= 6,-5.4%..-7.0%ongr_down_q8_8_4,-13.4%..-14.2%ongr_down_q8_8_8) andgr_up_f(-1.6%atntok=1,-6.0%atntok=2,-14.7%atntok=4,-18.0%atntok=8) + movedgr_q2s/gr_p4sto program-scopeconstant(67/67bitwise PASS); ported register-softmax, unrolled 8-cell V blocks, sharedek[]/L_sh, and barrier-freeqsa_write_preparedtoqsa_attn2(-36.8%..-40.6%) andqsa_merge2(-8.3%..-14.6%); and replacedgdn_step_t's 64-bitst_basewith 32-bitst_off+float4conv_w+constantgdn_q2s/gdn_p4s, and replacedmoe.metalplan_body's runtime integer divisioni / int(a.topk)withtopk10 ? ((i * 205) >> 11) : (i / int(a.topk))+unroll_count(10)(-9.5%..-20.4%fastermoe_routeand-11.7%..-26.6%fastermoe_planatntok = 2..8), setting all-time decode speed records across all 4 workloads! - Round 22
gr_down_4_4/_r2NR=2xnReuse (-8.1%ntok=4,-16.8%..-26.1%ntok=5..8) +moe_down/moe_guq2_combine(-7.7%..-22.5%atNE >= 2) +plan_bodysg==0SIMD Prefix Sum (-14.6%..-35.4%) + Barrier-Freeswiglu_prep(audits/opus_r22_gr.md,audits/sol_r22_router_glue.md,audits/argon_r22_moe.md): Added templateNR+ directfloat4FMA (noxsarray copy) togrd_acc/gr_down_f, using 128-threadgr_down_*_4_4atntok <= 4(-1.9%..-8.1%faster) andgr_down_*_4_4_r2/8_8_r2(NR = 2rows/simdgroup sharingxnreads) atntok >= 5ingr_readandmtp_gr_read(-16.8%..-26.1%faster,75/75bitwise PASS); addedq2_combine<F>(ushort2 pk[NR]; float d[NR];= 4 registers forNR=2, loadingxs[16]once per tokentacross allNRrows) inmoe.metal, cuttingmoe_down_142_8_2_4by-7.7%..-8.6%atNE = 2..4(-16.1%onT_Q2_0prep),Q2_0moe_guby-12.2%..-22.5%, andIQ2_XSmoe_guby-8.2%..-17.1%, plussg == 0simd_prefix_exclusive_suminplan_body(-8.8%..-22.7%fastermoe_routeand-14.6%..-35.4%fastermoe_plan); and convertedswiglu_prep(glue.metal) to program-scopeconstantglue_q2s/glue_p4s+ barrier-freesimd_shuffle_down(64 -> 0 BTG mem,-48barriers/window). - Round 23
ksgpopcount(-2.5%..-5.5%IQ2_XSmoe_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_gu3sign4_iq2xs& Constant-Ptr Folding (776.7 msIQ2_XS&45.9 msIQ3_S,3,293.32 ms/602.1 tok/sGPUkprof,42.01 tok/sDecode) (audits/opus_r23_quant.md,audits/sol_r23_gemv.md,audits/argon_r23_moemm.md):** Computedksgsign parity (i | ((popcount(i) & 1u) << 7)) forIQ2_XS(ctab == 1,-2.5%..-5.5%fastermoe_gu) andIQ3_XXS(TAB = 256,-3.9%) inquant.metal; added directfloat4FMAgemv_direct_x(Q8_0ntok=1&LPR=16, NR=1, ntok=2,-3.4%..-5.4%faster) andgemv_q2_raw_combine(Q2_0Rntok=1,-2.2%..-4.1%faster) ingemv.metal; and added 4-bit registersign4_iq2xs(j, hh)(zeropopcountonhh==0lanes) + constant-offset pointer folding ((k & 255) >> 3/>> 2) +fmascale math inmoe_mm.metal, cuttingW8<T_IQ2_XS>by **-4.19%..-4.35%** (pg:mm_gu3_22_16_32_32_4 x34down to **776.71 ms/22.84 ms/layer,-25.5 ms/chunk**) andW8<T_IQ3_S>by **-8.5%atT=2048(45.93 msforx2) and1.82xatT=512(18.73 -> 10.28 msacross the VGPR occupancy cliff!)**, pushing full-modelkprofGPU time atT=1983down to **3,293.32 ms(602.1 tok/spure GPU prefill)** andprompt_2048tok_opsdecode to **42.01 tok/s(23.80 ms/tok,3,023.11 ms`)! - Round 24
W8<T_Q6_K>ScalaruintOR & Address Folding (+12.8%..+24.9%atT=512..1024,+3.67%atT=2048) +qsa_flash32-Bitko[8]/vo[8]Offsets (-11.71%atpos0=8192) + Dead Sign-Table Removal (TAB = 1024/512) +gr_norm_h_moeck-Outers[10]&gemm_hc_down* 0.25f(audits/opus_r24_prompt.md,audits/sol_r24_qsa.md,audits/argon_r24_moemm.md): FoldedW8<T_Q6_K>andW4<T_Q6_K>address/shift math and replaced 4-wideuint4unpack+shift+OR with a single scalaruintOR beforefloat4(as_type<uchar4>(qa)), speeding upQ6_Kmm_gu3by+12.79%atT=512,+12.28%..+24.91%atT=1024, and+3.67%atT=2048(36.82 -> 35.52 ms/layer); movedsign4_iq2xstoquant.metal, updateddeq<T_IQ3_XXS/T_IQ2_XXS/T_IQ2_XS>to use rawson the low 4 values andsign4_iq2xs(s, 1)on the high 4 values, eliminatedksgandload_signsacross all formats, and shrunkfmt<T_IQ2_XS>::TAB(1024) andfmt<T_IQ2_XXS>::TAB(512); replaced 64-bitkr[8]/vr[8]pointers inqsa_flashwith 32-bituint ko[8]/vo[8]byte offsets (kq + (ko[j] + uint(k))andvq[vo[r + j] + uint(d)]), cuttingqsa_flashatT=1984, pos0=8192by-11.71%(74.61 -> 65.87 msforx3=-35.0 msacross 12 layers!); and invertedgr_norm_h_moec'sk/iloops (-0.5..-0.6%atT=2048) + replaced/ float(e.hc)with* 0.25fingemm_hc_down(-0.8..-1.1%). - 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 msinkprof:361.9 -> 319.2 ms) +gdn_postGate/Gamma Preload (-0.7%) +mm_gu3half4HSStore Pointer Hoist (audits/opus_r25_gdn.md,audits/sol_r25_ple.md,audits/argon_r25_moemm.md): Unrolledmm_down3(T_Q2_0,TK=32) by 2TK=32steps per 64-elementT_Q2_0block (k0 += 64, w0_l += 18) withW4<T_Q2_0>::run_q<2/6/10/14>(rdw, uint kl4)using(b + kl4)[OFF]immediate displacement, speeding upmm_down3by+2.48%atT=1024(5,020.1 -> 4,898.7 us),+1.85%atT=2048(8,060.3 -> 7,914.2 usin isolated A/B, and361.9 -> 319.20 ms=-42.70 ms/-11.8%across 48 layers in full-modelkprof!) + hoistedot_p/hs_pinmm_gu3'shalf4coalescedHSstore (+0.09..+0.91%acrossIQ2_XS/IQ2_XXS/IQ3_XXS/Q2_0); and preloadedgammaand precomputed thesigmoidgate before the 4simd_sumreductions ingdn_post(-0.7%). - Round 26
qsa_attn_mm32-Bitko[8]/vo[8]Offsets &float4TG Loads (-23.3%..-29.9%acrosspos0=2048..30000) +gemvT_IQ4_XSpacked_uint45-Registergemv_iq4_combine(-1.5%..-5.0%Cold DRAM on2560->640,-1.0%..-2.7%atntok=3..4on2560->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-bitkr[8]/vr[8]pointers inqsa_attn_mmwith 32-bituint ko[8]/vo[8]byte offsets + unroll hints +float4threadgroup loads, cuttingqsa_attn_mm(T=1984) by-23.31%atpos0=2048(30.88 -> 23.68 ms),-29.93%atpos0=8192(31.79 -> 22.28 ms), and-24.74%atpos0=30000(37.32 -> 28.09 ms); addedgemv_iq4xs_codes(packed_uint4 pk[NR]+float d[NR]=5registers/row instead of36forfloat q[32] + float4 sc, gated onLPR <= 16 || NTOK >= 4) +GEMV_AB_COLD=1DRAM-streaming A/B benchmark ingemv.metal/gemv_parity.cpp, speeding up cold-DRAMIQ4_XS 2560->640(ffn_gate/up_shexpx57/step) by-1.5%..-5.0%atntok=1and-2%..-3%atntok=2..4, and2560->6144/10240/12288(attn_gate/qkv/q) by-1.0%..-2.7%atntok=3..4; and added direct unrolledDW<T_IQ4_XS>andDW<T_IQ4_NL>+rd32SWARDW<T_Q6_K>(1.02x–1.07xfasterdq_f16) inmma.metal+#ifndef MM_V3_ONLY(10xfastermoe_mm_parity --v3compile). - Round 27
gdn_step_t5 -> 4Barriers/Token +dt/ssm_aHoist +part[8]/red2[IB-1][S]Shrink (-3.0%..-7.9%) +qsa_merge2/qsa_mergeParallelfloat2 (m, l)Loads &simd_max(-8%..-20%) +Window::gemv_tilen_out=512{32,2,4,1}(-3.2%..-25.5%) &Q2_0640->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-threadthreadgroup_barrierper token (5 -> 4barriers/token =-144barriers per 4-token verify window across 36 GDN layers!) ingdn_step_t(gdn.metal) by writing rawcvtoqs/ks/vsbefore barrier 1 and computing(rsq * qs[qk_idx]) * scaleand(rsk * ks[qk_idx]) * scaleafter barrier 1, plus hoisting token-invariantdt[h]&ssm_a[h]and shrinkingpart[32] -> part[8]andred2[IB][S] -> red2[IB-1][S], winning all16/16productionplane = 49152configurations (-7.93%atntok=4 verify,-7.82%atntok=4 commit,-3.17%atntok=1 verify,88/88bitwise); vectorized(m, l)loading inqsa_merge2andqsa_merge(qsa2.metal,qsa.metal) asfloat2 ml = *(device const float2*)(pp + tid * STRIDE + QD)+simd_max(M)+ threadgrouplk[k]readout, speeding upqsa_merge2by-14%..-20%atntok=1and-8%..-14%atntok=4(80/80bitwise); and fixedWindow::gemv_tile(window.cpp) so2560 -> 512(attn_k&attn_vacross 12 QSA layers) uses{32, 2, 4, 1}(preservingLPR=32for100%bitwise parity across all 13 formats, verified bygemv_parity's newsame LPR: bitwisecheck) andQ2_0 640 -> 2560atntok >= 4uses{8, 2, 4, 1}(-3.2%..-25.5%fasterIQ4_XS 2560->512,-9.0%..-24.4%fasterIQ3_S 2560->512,-7.6%..-13.9%fasterQ6_K 2560->512,-4%..-6%fasterQ2_0 640->2560). - Round 28
mtp.fc_hidden[t][c]Stream Batching (-52.1%atn=2,-29.3%atn=3,-17.0%atn=4) +ple.metaltg_sumBarrier Halving (-1.5%..-5.8%) +sample.metaltg_argmaxUnroll &router_route_8(audits/opus_r28_mtp.md,audits/sol_r28_router.md,audits/argon_r28_ple.md): Batchedmtp.fc_hidden.weightacross alln * HCstreams in[t][c]order (<= MAXTrows pergemvdispatch inwindow.cpp&mtp.metal), cutting the MTP hidden branch atn=2(a=2accepted drafts) by-52.1%(211.6 -> 101.4 us,9 -> 3dispatches), atn=3(a=3accepted drafts) by-29.3%(220.1 -> 155.7 us,9 -> 5dispatches), and atn=4by-17.0%(228.9 -> 190.0 us) while matchingnseparaten=1chains bit-for-bit; halvedthreadgroup_barriers inple.metal(2 -> 1inple_gnorm,4 -> 2inple_gate_norm,6 -> 3inple_query_gate_norm&ple_comb_query_gate_norm,8 -> 4per stream inple_moe_gate_norm&ple_moe_gate_norm_c) via ping-pongpart[16], speeding up decode/verify PLE (T=1..7) by-1.5%..-3.7%andT=64ple_moe_gate_norm_cby-1.5%..-5.8%(100%bitwise); and added#pragma clang loop unroll(full)totg_argmaxinsample.metal+router_route_8(T_Q8_0) inrouter.metal. - Round 29
mm_gu3<T_IQ2_S>Super-Block Scale Caching (-86.24 ms/-10.0%inkprof) &W8<T_IQ2_XXS>Byte Loads (-15.42 ms/-6.5%inkprof) +moe_routee512/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 scalesdg[0..3]/du[0..3]across the 8TK=32steps per 256-element block inmm_gu3<T_IQ2_S>and shrankfmt<T_IQ2_S>::TABfrom2080to2048(16,512 -> 16,384 Bexact16 KBTG memory), plus direct byte loadsp[l]/p[l+2]inW8<T_IQ2_XXS>andfmainW8<T_IQ3_S>, cuttingpg:mm_gu3_22(34T_IQ2_Slayers) inkproffrom859.30 -> 773.06 ms(-86.24 ms/-10.0%faster,+5.48%atT=1024,+5.08%atT=2048in isolated A/B) andpg:mm_gu3_16(12T_IQ2_XXSlayers) from237.28 -> 221.85 ms(-15.42 ms/-6.5%faster); eliminated dynamic register indexingv[ex / 32] = -INFINITYin bothroute_topk(-0.7%..-1.2%,5.78 -> 5.305 msinkprof) andmoe_route(moe.metal), where addinge512+ preloaded 16-bitres_maskfromsmapcutmoe_route(split=1) by-26.4%atntok=1,-24.4%atntok=2, and-20.6%atntok=4; and optimized decodedeq<T_IQ2_S>(ushortq2/s2+ pre-shiftedqh+#pragma clang loop unroll(full),-1.1%..-6.9%faster acrossQUANT_CTAB=0/1),deq<T_IQ2_XXS>(-5.1%..-5.9%faster), anddeq<T_IQ3_XXS>(-0.7%..-2.9%faster). - Round 30
W4<T_Q2_0>::decIEEE-754 FP160x6400Exponent Trick (-10.5%mm_down3atT=2048/-37.02 msinkprof) +gdn_pre&gdn_postVectorsimd_sum(float4)Reductions (audits/opus_r30_moemm.md,audits/sol_r30_prompt.md,audits/argon_r30_gdn.md): Replaced FP32int->float->half4conversion inW4<T_Q2_0>::run_qwith 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 all65,536 x 256 = 16,777,216FP16dx code-byte combinations, cuttingmm_down3<T_Q2_0, 16, 64, 32, 2>by-10.5%atT=2048(7,950 -> 7,113 us),-8.3%atT=1024, and-37.02 ms(312.39 -> 275.37 ms/-11.9%, down-23.9%from Round 24's361.9 ms!) in full-modelkprof; and replaced the12scalarsimd_sum(float)trees ingdn_preand4scalarsimd_sum(float)trees ingdn_postwith bitwise-identical vectorsimd_sum(float4)reductions (12 -> 3and4 -> 1) + coldt == T - 1/a.f32branch hoists (-0.66%..-1.05%ongdn_post). - Round 31
mm_gu3<T_IQ2_S / T_IQ2_XXS / T_IQ2_XS / T_IQ3_S>tab + hh& Sign-Shift (-45.32 msIQ2_S/-12.56 msIQ2_XXSinkprof) +moe_down/gemvQ2_0R0x6400Half-Mantissa Decode (-7.6%..-16.9%) +dq_f16<T_IQ4_XS / T_IQ4_NL>Register Table Init &perm == 4Contiguous Writes (audits/opus_r31_moemm.md,audits/sol_r31_moe.md,audits/argon_r31_dq.md): Hoisted invarianthh = (kl >> 2) & 1andth = tab + hhacrossmm_gu3, replaced the divergenthh ? s >> 4 : sternary inW8<T_IQ2_S>andW8<T_IQ3_S>withs >> (4 * hh), and extended 256-wide super-block scale cachingdg[]/du[]toT_IQ2_XXSandT_IQ2_XS, speeding upmm_gu3<T_IQ2_S>by+8.17%..+12.79%(773.06 -> 727.74 msinkprof,-131.56 ms/-15.3%faster than Round 28!),mm_gu3<T_IQ2_XXS>by+4.08%..+6.86%(221.85 -> 209.29 msinkprof),mm_gu3<T_IQ2_XS>by+1.41%..+6.16%, andmm_gu3<T_IQ3_S>by+2.23%..+6.84%(100%bitwisePASS); replaced scalarfloat((pk >> ...) & 3) - 1.0fdecode inq2_combine<T_Q2_0R>(moe.metal) andgemv_q2_raw_combine(gemv.metal) with packed IEEE-754 FP160x64006400umantissa-OR decode, cutting decodemoe_down(dn, all 48 layers!) by-7.6%..-16.9%(ratio 0.8311..0.9236acrossNE=1..4andntok=1..8) andgemvQ2_0R(ntok=1) by-11.8%(100%bitwisePASS); and replaced the dynamicload_tab<T_IQ4_NL>loop and divergentconstant-memorykvalues_iq4nlreads indq_f16<T_IQ4_XS / T_IQ4_NL>(dense_mm.metal) with 4 bank-conflict-free register stores, guardeda.perm, and invertedro = gid.yforperm == 4(1.03x–1.05xon320x10240p4,100%bitwisePASS). - Round 32
qsa_flashsp[4][HPG * QCH]3 -> 1Barrier & Pre-Addedko[j]+ Loop Unrolling (-4.03%..-5.47%/-6.10 msinkprof) +mm_gu3<T_IQ2_S>half4CoalescedHSStore (-9.00 msinkprof) + Decodedeq64-Bituint2Grid Loads &deq_iq_packed(audits/opus_r32_moemm.md,audits/sol_r32_qsa.md,audits/argon_r32_moe.md): Expandedsp[2][HPG * QCH]tosp[4][HPG * QCH](13,760 < 16,384bytes TG memory) inqsa_flash(qsa_prompt.metal) so all 4 simdgroups storeSpartials simultaneously across 1threadgroup_barrierinstead of 3 (-64barriers/token atpos0 >= 2048), pre-addedd0 + kltoko[j]for constant-offset+0/+16/+32/+48Kloads, and unrolled fixed-trip loops, winning7/7pairs at every position on the quiet GPU (-4.31%atpos0=0,-5.47%atpos0=2048,-4.03%atpos0=8192,-0.80%atpos0=30000, and130.50 -> 124.40 ms/-4.7%in full-modelkprof!); enabled thehalf4coalescedHSstore epilogue inmm_gu3(moe_mm.metal) forT_IQ2_S, cuttingpg:mm_gu3_22(34layers) by another-9.00 ms(727.74 -> 718.74 msinkprof,-140.56 ms/-16.4%faster than Round 28, and pushingT=2048under25,000 usto24,986.6 us!); and replaced pair 32-bit grid loads indeq<T_IQ2_XXS / T_IQ2_XS / T_IQ2_S>(quant.metal) with 64-bituint2loads (4 -> 2loads per unit), unrolleddeq<T_IQ3_S>(-2.7%..-6.6%fasterIQ3_Smoe_gu), and added low-register 6-VGPRdeq_iq_packed<F>+iq_combine(-9.4%..-15.8%atNR=1, NE=4inmoe.metal). - Round 33
gdn_step_tCritical-PathklSplit + Pre-Barrier-3stateCommit (-1.1%..-4.1%Across32/32GDN Configs) +gr_normApply-Hoist & Pre-Offset Pointers (-6%..-11%) +GrdUnit<T_Q8_0, 2>Packedchar4Decode (-10%..-16%) (audits/opus_r33_gr.md,audits/sol_r33_qsa.md,audits/argon_r33_gdn.md): Splittot_q/rsq/ql/g/be/vjacross Barrier 2 ingdn_step_t(gdn.metal) so onlyklsits on the critical path before the 32-stepkvFMA loop, movedhist_outwriteback before Barrier 1 (killinghstlive ranges across all 4 barriers onGDN_NTOK == 1), and issued the 32 globalstatestores before Barrier 3 to overlap the 64 KB/head state writeback withib == 0's readout reduction (0.9589x..0.9895xonplane = 0,0.9717x..0.9959xacross all16/16productionplane = 49152configs,88/88bitwisePASS) + pre-offsetmoe_combine_sh&swiglu_prep(glue.metal,24/24bitwisePASS); hoistedif (a.apply)and pre-offset pointers ingr_norm(gr.metal,-6%..-11%atntok=1..4,-2%..-3%atntok=7..8) and addedGrdUnit<T_Q8_0, 2>keepingQ8_0codes packed inchar4 q[2][8](8VGPRs/row vs32floats) for_r2gr_down(-13%..-16%atntok=7,-10%..-15%atntok=8,100%bitwisePASS), cutting pure GPUkprofatT=1983to3,153.46 ms(628.51 tok/spure GPU prefill,3,270.77 msserialized sum!) and slashingmtp_msto new all-time records acrossreasoning(323.71 ms,-12.4%vs369.70 ms),persona_chat(383.38 ms,-9.8%vs424.85 ms), andprompt_2048tok_ops(193.49 ms). - Round 34
gemm_hc_downTN = 96Tail-Tile Decomposition (-5.05%..-7.18%) +grd_accROLL = trueatKS = 4(-1.9%..-4.6%on_r2) +W8<T_IQ3_S>::block& Coalescedhalf4HSStore +T_IQ2_XSsboff(audits/opus_r34_gr.md,audits/sol_r34_prompt.md,audits/argon_r34_moemm.md): Replacedgemm_hc_down<TM, 128, SG>'s 128-column tail tile (tg.x = 2,c0 = 256, where only68columns256..323exist) with a single96-column (3 x 32)hc_down_tile<TM, 96, SG>, cuttinggemm_hc_downon the quiet GPU by-5.05%atT=512(ratio 0.9495),-6.60%atT=1024(ratio 0.9340),-7.18%atT=1983(1.1762 -> 1.0927 ms,ratio 0.9282), and-7.03%atT=2048(1.2062 -> 1.1209 ms,ratio 0.9297, makinggemm_hc_downfaster thangemm_hc_upin full-modelkproffor the first time ever:125.998 msvs132.822 ms!); addedROLL = F == T_Q8_0 && NR == 2 && KS == 4togrd_acc(gr.metal) so packedQ8_0NR=2(GrdUnit<T_Q8_0, 2>) keeps its token loop rolled atntok = 4..6(-4.6 / -3.1 / -1.9%faster4_4_r2atntok = 4/5/6,-13.7%and-23.9%faster than4_4atntok = 5, 6) while keepingwindow.cpp:1165atntok > 4to preserve 324-threadgroup cold-DRAM latency hiding atntok = 4; and addedW8<T_IQ3_S>::block+F == T_IQ3_Shalf4HSepilogue stores (+0.48%atT=2048) +T_IQ2_XSsboffreuse (-0.69%atT=1024,-1.25%atT=2048) inmoe_mm.metal, pushingprompt_2048tok_opsto432.78 tok/s(4,584.29 msend-to-end prefill) and42.27 tok/s(3,004.68 msdecode) and@10.2Kdecode to25.04 tok/s(39.93 ms/tok,1,237.88 ms)!
- MTP
- Round 35
Wr/Wr2Router Dequant Reuse (-47Dispatches/Chunk) +gemm_f16_16_96_2Walbe+M % 128 <= 6464_128_4(1.14x@M = 306) +hc_up_tileRow-Tail Trim (-2.5%@T = 289) +mtp_rms_xprep/mtp_embed/mtp_slop_penVectorization (audits/opus_r35_dense.md,audits/sol_r35_prompt.md,audits/argon_r35_mtp.md): Ping-pongedp.Wrandp.Wr2by layer parity (l & 1) inWindow::prompt_chunk(window.cpp), eliminating47redundantdq_f16_0router dequantizations per chunk (2,094 -> 2,047dispatches inkprof), addedgemm_f16_16_96_2(TN = 96) forWalbeatM >= 640(1.16x/-13.6%faster atM = 1983), selected64_128_4instead of128_64_4on big-N GEMMs whenM < 1536 && M % 128 in 1..64(1.14x/-23.5 msfaster chunk total atM = 306), addedhc_up_tile<16/32, TN, SG, REG>row-tail trimming togemm_hc_up(-2.5%atT = 289), and vectorized/unrolledmtp_rms_xprep(-37.5%atn=2560,-35.8%atn=10240),mtp_embed(-15.8%..-31.6%across all 13 GGUF formats),mtp_slop_pen(-96.5%onmode=0,-45.7%..-70.6%onmode=1,-31.0%..-63.9%onmode=2), androuter_route(-0.2%..-2.1%) (100%bitwisePASSacross all suites!). - Round 36 Decode
predL_/predL2_Token-0 Fast Path WhenSTRATA_METAL_PREFETCH_DRAFT <= 0(-16.0%Hot /-6.8%Cold Across93Router GEMVs/Window) + GuardedTM > 32Row-Tail Trimming ingemm_hc_down&gemm_swiglu_up(audits/opus_r36_decode.md,audits/sol_r36_prompt.md,audits/argon_r36_qsa.md): Addedtile_ntoktoWindow::gemv(window.cpp) so whenSTRATA_METAL_PREFETCH_DRAFT <= 0(= 0in production, whereExpertStream::aheadonly readst = 0), the93router-guess GEMVs (47predL_+46predL2_,Q8_02560 -> 512) run atntok = 1on the window's tile (gemv_8_16_1_4_1#1), producing 100% bit-for-bit-identicalt = 0predictions while cutting their GPU time by-16.0%hot (21.78 -> 18.30 us) and-6.8%cold (23.91 -> 22.29 us); addedif constexpr (TM > 32)row-tail (TM = 16 / 32) trimming togemm_hc_down(-4.71%atT=1025,-3.73%atT=1041) andgemm_swiglu_up(-1.42%atT=1025,-1.14%atT=1041,-0.78%atT=1987) while preservingTM = 32(SG = 2) kernels byte-for-byte (252/252epibitwise checksPASS); and confirmedqsa_prep/qsa_score_mmare at their M5 register/barrier optimum.
- Fetched all 31
-
Round 37 Multi-Chunk Prompt Hot-Expert LRU & Scratch Pool Preservation (
stream.h,stream.cpp,window.cpp): Passedmulti_chunk = (n > C)intoprompt_region(true, n > C)->hold_region(true, multi_chunk)so 1-chunk prompts (<= 2,048tokens) run 100% identically (42.49 tok/s/2,989.17 ms@2Kdecode — new all-time record,88.9%accept), while multi-chunk prompts (> 2,048tokens, e.g.,prompt_albedo_11k.txtat10,180tokens) keep!x.hotscratch slots atlru_back(s)across chunks, markpass == 1demotionsx.hot = falseinvictim()sopass == 0reuses them instead of evicting morehotdecode experts, and restrictallow_hotlookahead todistance == 1— boosting@10.2Kdecode from25.26 tok/s(1,227.18 ms,45.2%accept,2,713SSD reads) to34.01–34.74 tok/s(892.47–911.51 ms,+34.6–37.5%faster,78.3%accept,2,073SSD reads /-23.6%fewer SSD reads) and@10.2Kprefill from423.48 tok/s(24,038.85 ms) to458.34–472.89 tok/s(21,527–22,210 ms,-1.83s to -2.51sfaster), with pure GPUT=1983prefill (kprof) hitting3,131.29 ms(633.28 tok/s— new all-time record). -
Round 38
hold_regionLayer-0allow_hot = false+ Prompt/Decodebytes_read_/reads_Split +mtp_prefilltail_capSync withSTRATA_METAL_MTP_WIN(stream.h,stream.cpp,window.cpp,audits/astra_r38_qsa.md): Passedallow_hot = falsetovictim()duringhold_region(true)'s initial Layer-0 prefetch (stream.cpp:1080) and fallback static lookahead (stream.cpp:890) sohold_regiononly fills the251!hotscratch slots ofcls_[0](IQ2_XXS) instead of evicting~180hotdecode experts for tail Layer-0 experts thatat_layer(0)cancels 5 ms later; splitp_bytes_read_/p_reads_from decodebytes_read_/reads_and zeroedbytes_read_/reads_after the6.5 GBstartup 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 prints480 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 unlessSTRATA_METAL_TRACE=1, becausestrata_window_cli.cpp:1827callswindow_reset_stats(w)unconditionally andwindow.cpp:4137->stream.cpp:1177-1180zeroesp_lookups_ p_hits_ p_pre_ p_misses_ p_streamed_ p_chunks_ p_bytes_read_ p_reads_. any prefill MB/read figure published withoutTRACE=1is decode-only); syncedWindow::mtp_prefill's defaulttail_capwithSTRATA_METAL_MTP_WINclamped to[64, 2050]; and verified inqsa_parity --ab(audits/astra_r38_qsa.md) that cachingkeys[8]acrossqsa_topk2's 1024-thread barriers drops M5 occupancy (+3.5–12.9%slower atpos0 = 10180), pushing@10.2Kprefill to473.80 tok/s(21,485.94 ms— new all-time record!),@10.2Kdecode to35.21–35.47 tok/s(873.92–880.47 ms,mtp_ms = 106.34 ms— new all-time record!),@2Kprompt_2048tok_opsdecode to42.57 tok/s(2,982.99 ms,mtp_ms = 194.09 ms— new all-time record!), andreasoningdecode to37.20 tok/s(5,886.81 ms— new all-time record!). -
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 localqsa_merge2batching experiment is bitwise exact (40 edge cases), but only tests dispatch it; live engine decode remains onqsa_merge. rejected the cuda score butterfly (memcmp failure) andqsa_mergebatching (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 upstreamanswer_roomby 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/. -
Round 41 Bonsai-Style Automatic NVMe Prompt & Conversation
.kvcCache (--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.kvcdisk tier and 2-prompt shared-prefix learning/pinning architecture fromokami.mom/llama-cpp-bonsai(prompt-cache-disk,cache-tuning,prefix-learn,prefix-restore) intostrata_window_cli --serveandwindow.cpp: (1) preserved the MTP draft layer's valid KV cell range (mtp_sink_,mtp_cell0_,mtp_cell1_) inscr::SessionHeader::flagsacrossWindow::session_save/Window::session_loadand shiftedmtp_cell0_/mtp_cell1_inWindow::scr_relocateso restored/relocated sessions retain full MTP draft context without re-prefilling; (2) addedDiskCache(--cache-disk <dir>/STRATA_METAL_CACHE_DISK,--cache-disk-size <mib>/STRATA_METAL_CACHE_DISK_MIB, default16384 MiB) with< 1 msstartup.kvcheader indexing (scr::SessionHeader+n_ids * 4token IDs +.pads+.pinwithout reading the120 MiB+KV payload), corrupt/partial file cleanup, post-DONEzero-latency auto-save (4-threadF_NOCACHEpwriteat5.2 GB/s), and automatic obsolete-prefix pruning; (3) addedobserve_prefix(prompts-seen.bin) which detects when any 2 distinct prompts diverge after a shared prefix>= 32tokens (lcp < cur.size() && lcp < s.size()), splits prefill atlearn_at = lcp, saves a pinned.kvc+.pincheckpoint (up to16pins, up to64Ktokens, never evicted by normal conversation LRU), and ported Bonsai'sprefix-share(8906fd8) as zero-copy GPUMTLBufferdelta.kvcstorage (c0 = (parent_cells / 4) * 4inSessionHeader+.parentsidecar + orphan pruning, writing only[c0, n_cells)QSA/MTP positional cells per child and reading[0, c0)directly from the parent.kvcviaF_NOCACHEpread); and (4) wiredfind_bestinserve_loopso 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.kvcsession (preadat5.2 GB/s), and run either exact continuation (cached_tokens = 104 / 130onConv B Turn 2,cached_tokens = 65 / 89on brand-newConv C Turn 1from pinned system prefix) orscr_plan+scr_relocatesplice directly over NVMe (cached_tokens = 149 / 184on middle-editedConv D Turn 2!). Also fixedSAVE/RESTOREwire responses (SAVED <tokens> <bytes> <ms>/RESTORED <tokens> <bytes> <ms>,30/30serve.test_slotsunit testsOK). 40. Measured End-to-End Speed (32KContext,6.5 GBArena, UnprunedK=512IQ2_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_opsprefill400.88–432.78 tok/s(4,584.29–4,949.14 msend-to-end — all-time record4,584.29 ms/432.78 tok/s—-1,728 ms/+37.7%faster end-to-end,3,131.29 ms/633.28 tok/spure GPUkprof—3,267.43 msserialized sum!,2,047dispatches (-47dispatches/chunk!),1.12 GBtransient scratch!), decode41.53–42.27 tok/s(23.66–24.08 ms/tok,3,004.68–3,005.02 msdecode,mtp_ms = 190.62–197.99 ms/1.63 ms/draft— new all-time MTP record190.62 ms!);prompt_albedo_11k(10,180tokens) prefill421.49–446.59 tok/s(22.79–24.15 svs28.25 s,-5.46 sfaster — 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.2Kdecode hit25.00–25.26 tok/s/39.59–40.01 ms/tok/1,227.18–1,240.21 ms— all-time record25.26 tok/sat10.2Kcontext! vs22.06 tok/s/1,405.37 ms!);codingprefill120.61–126.78 tok/s(2,413–2,537 ms,+32.2%), decode35.46–37.1 tok/s(6,115.57–6,176.28 msdecode — all-time record6,115.57 ms,mtp_ms = 374.22–380.29 ms);reasoningprefill121.57–122.92 tok/s(2,416–2,458 ms,+33.3%), decode34.92–36.74 tok/s(5,961.16–6,270.75 msdecode,mtp_ms = 322.48–370.78 ms— all-time record322.48 ms!);persona_chatprefill124.62–128.14 tok/s(2,255.32–2,443 ms,+38.5%— all-time record2,255.32 ms/128.14 tok/s!), decode28.50–33.2 tok/s(6,628.00–7,685.17 msdecode,mtp_ms = 374.58–424.81 ms— all-time record374.58 ms!). -
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.71nats excess NLL. the first GDN layer alone accounts for63–78%of summed single-layer INT4 damage atH=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.15709preserves updates with a low-rank bypass (LRB) and bounded directional rounding (BDR), rather than assuming unbiased noise is safe. its GDN result uses5.25average /6.00peak state bits including bypass, not bare INT4; WikiText-103 8K PPL is26.85vs FP3224.19and naive INT479.76. - rule: keep Swift's existing high-precision recurrent path. our roughly
112.2 MiBfixed GDN state is small next to the6.50 GBexpert 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.
- 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 (
-
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 KLD0.034321, proxy0.038142–0.038249, random0.037728/0.037808(random marginally better, within about 1 SE of proxy). proxy p99 KLD rises0.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.
-
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-chunk0d18939749fb6f294fb0c272a61cca5d1,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(arm0) came out exactly0.000000mean KLD,64/64, on both corpora and both cycles. that is structural, not luck:prompt_moce = 0unpacks top_prot = 0, p_bonus = 0(prompt.metal:41-47), so the bias path atprompt.metal:51never fires and routing never readssmap, making arm 0 residency-independent. every boosted arm readssmap[idx] != 0xFFFFFFFFuat: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 earlier0.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.055902then0.055018(spread0.000884,1.6%), and that was written up as "reproducible ... tighter even than its chat value". a third run of the same config gave0.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) andplan(0.674194) — prose, argmax unchanged — with49.65%of KL in the top4positions and72.47%in the top8. the actual argmax flips were elsewhere and small-KL, includingtrue -> falseat 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/64for every arm on chat (zero discriminating power at that KL magnitude) and60/64,61/64,59/64/59/64,58/64on 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 arep-samples; fixed corpus ids are not, so neither the identity nor a Monte-Carlo SE applies.cand_ppl 2.4035vsref_ppl 2.5027atn=64therefore 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:1827callswindow_reset_stats(w)unconditionally after prefill + the first head step,window.cpp:4137prints only underPKPROF/TRACEand then alwaysreset_stream_stats(), andstream.cpp:1177-1180zeroesp_lookups_ p_hits_ p_pre_ p_misses_ p_streamed_ p_chunks_ p_bytes_read_ p_reads_. the shipped375.2 MB / 258 readsinresults/round39/*.log:25reports480 distinct lookups=48 layers x top_k 10= one decode step, not a prefill. alsop_streamed_is a count of submissions, not bytes (stream.cpp:893,931,1117); bytes isp_bytes_read_(:1192). - trap 7 — never rank arms from a sequential sweep. two arms with the identical chunk-0 setting landed
10,075.62 msand12,115.41 ms(20% apart) purely from sweep position, via the arena fill read rate going2.2 s -> 0.8 sas 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/64and 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).
- trap 1 — a control can be deterministic for a reason its treatments do not share.
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_*_expsis 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
ein 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
GSQGroupQuantLinearQAT 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.