License: MIT (the ggml code, the font and the projection vector keep their own licenses)

Engine 0.1.16; setup requires it (the Coder needs the model's own expert count, PR #54 )
The GSQ-RCO Coder (256 of 512 experts): setup, profile, speeds (PR #54 )
2026-11-07 17:38:20 -10:00 · 2026-09-28 14:29:00 +02:00 · 2026-09-28 14:27:45 +02:00 · 2026-09-28 14:20:58 +02:00 · 2026-09-28 13:45:47 +02:00 · 2026-09-28 13:45:39 +02:00
35 changed files with 1984 additions and 183 deletions
@@ -8,7 +8,7 @@
 # written against sm_120 (RTX 5000 / Blackwell) and silently building for an older arch would produce a
 # binary that runs and is wrong about its own performance, which is the worst of both outcomes.
 cmake_minimum_required(VERSION 3.24)
-project(strata VERSION 0.1.13 LANGUAGES CXX)
+project(strata VERSION 0.1.16 LANGUAGES CXX)
 add_compile_definitions(STRATA_VERSION="${PROJECT_VERSION}")
 # issue #31: release builds keep a PDB beside the exe (not shipped in it), so a minidump from a stalled engine can
 # be read back to the exact source line
@@ -514,13 +514,16 @@ if(STRATA_NATIVE_EXPERTS)
  endif()
  # part of strata_kernels_cpu (the pool calls these and they call its CPU probe: one library, so GNU ld's
  # one-pass static linking has no cycle to resolve)
-  target_sources(strata_kernels_cpu PRIVATE src/kernels/cpu/native_expert.cpp src/kernels/cpu/iq_avx512.cpp)
+  target_sources(strata_kernels_cpu PRIVATE src/kernels/cpu/native_expert.cpp src/kernels/cpu/iq_avx512.cpp
+                                           src/kernels/cpu/iq_avx2.cpp)
  target_include_directories(strata_kernels_cpu PRIVATE third_party/ggml)
  if(MSVC)
    set_source_files_properties(src/kernels/cpu/iq_avx512.cpp PROPERTIES COMPILE_OPTIONS "/arch:AVX512")
+    set_source_files_properties(src/kernels/cpu/iq_avx2.cpp PROPERTIES COMPILE_OPTIONS "/arch:AVX2")
  else()
    set_source_files_properties(src/kernels/cpu/iq_avx512.cpp PROPERTIES COMPILE_OPTIONS
                                "-mavx512f;-mavx512bw;-mavx512vl;-mavx512dq;-mavx512vnni;-mavx512vbmi;-mf16c;-mfma")
+    set_source_files_properties(src/kernels/cpu/iq_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2;-mfma;-mf16c")
  endif()
  target_link_libraries(strata_kernels_cpu PUBLIC ggml-cpu ggml-base)
  target_compile_definitions(strata_kernels_cpu PUBLIC STRATA_NATIVE_EXPERTS=1)
@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (c) 2026 Niko1221 and the Strata contributors
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
@@ -25,13 +25,15 @@ Measured on an RTX 5070 (12 GB), a Ryzen 5 7600 and 64 GB of RAM:

 | Size | Writes answers (short chat) | Writes answers (128K context) | Reads your prompt |
 | --- | ---: | ---: | ---: |
-| **Q2_0** | 95 tokens/s | 65 tokens/s | 539 tokens/s |
-| **IQ2_XS** | 78 tokens/s | 52 tokens/s | 463 tokens/s |
-| **IQ3_XXS** | 66 tokens/s | 45 tokens/s | 410 tokens/s |
-| **IQ3_S** | 54 tokens/s | 42 tokens/s | 374 tokens/s |
+| **Q2_0** | 90 tokens/s | 67 tokens/s | 1,310 tokens/s |
+| **IQ2_XS** | 74 tokens/s | 60 tokens/s | 1,240 tokens/s |
+| **IQ3_XXS** | 62 tokens/s | 46 tokens/s | 1,110 tokens/s |
+| **IQ3_S** | 52 tokens/s | 41 tokens/s | 1,070 tokens/s |
+| **Coder** (IQ1_M) | 51 tokens/s | 44 tokens/s | 1,300 tokens/s |

 - **Writes answers** = how fast the reply appears (tokens per second).
- **Reads your prompt** = how fast it takes in what you send (long documents, code, chat history).
+- **Reads your prompt** = how fast it takes in what you send (long documents, code, chat history), measured on a
+  32K-token prompt; a 4K prompt reads at 740-1,000 tokens/s. A 32K prompt takes about 25 seconds with Q2_0.

 A card with more VRAM is faster, because more of the model fits on the GPU: an RTX 3090 (24 GB) should do roughly
 100-140 tokens per second. All measurements, long-context numbers and estimates for other cards are in the
@@ -57,11 +59,17 @@ lower the RAM needed.
 **The version:**

 - **Qwen3.8-Flash-Next** - the original.
+- **[Coder](https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF)** - ISTA-DASLab's coding
+  version: half of the experts removed, keeping the ones that code, tool use and images need (91% of the full model's
+  SWE-bench Verified score, 99% of LiveCodeBench, by its authors). One size (IQ1_M: its experts stored like IQ3_S):
+  shard 1 is **29.6 GB**, so it fits a PC with **32 GB of RAM**, runs 262K context on 64 GB, and reads long prompts
+  the fastest of all. Weaker outside coding.
 - **[Swift 1.5](https://huggingface.co/ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF)** - a fine-tune by UkisAI
  that thinks much shorter before answering, so you get the answer sooner, with about the same quality. Same speed per
  token, and about the same RAM as the same size of the original (no IQ3_S). Its own license applies (see its page).

-Not sure? Take **IQ2_XS**. You can add another one later with `START-HERE.bat --setup`.
+Not sure? Take **IQ2_XS** - or the **Coder** if you mainly write code, or have 32-48 GB of RAM. You can add another
+one later with `START-HERE.bat --setup`.

 ## Install

@@ -83,9 +91,17 @@ App). Everything else - Python, the engine, the model - is set up for you.
 Then it downloads everything (the model is ~70 GB, so the first time takes a while - you can stop and it picks up
 where it left off) and **starts the model**. Your browser opens the Strata app at `http://127.0.0.1:8080`.

+> **While the model starts, your PC can be slow or stop responding for 1-3 minutes** (longest the first time): Strata
+> loads 35-55 GB into your RAM and locks part of it for the graphics card. That's normal - wait, and don't close the
+> window. The window tells you what it is doing.
+
 **Next time**, just double-click `START-HERE.bat` again: it starts right away, nothing is downloaded twice. Close its
 window to stop the model.

+**Updating:** download the new version and unzip it anywhere (or `git pull`), then run `START-HERE.bat` in it. The
+model files are kept in a `Strata-data` folder next to your Strata folder, so a new copy finds them and sets itself up
+the same way - nothing big is downloaded again.
+
 **Linux:** run `./setup.sh` - same questions, same result.

 ## Using it
@@ -112,8 +128,8 @@ window to stop the model.
 ## Something went wrong?

 **My PC froze, or got very slow, the first time Strata started.**
-That's normal the first time. Strata loads 35-55 GB into your RAM, locks part of it for the graphics card, and works
-out how much of the model fits on your GPU. The mouse can freeze for a few minutes. **Wait, and don't close the
+That's normal while it starts, most of all the first time. Strata loads 35-55 GB into your RAM, locks part of it for
+the graphics card, and works out how much of the model fits on your GPU. The mouse can freeze for a few minutes. **Wait, and don't close the
 window.** The next starts are much faster. Still frozen after 10 minutes? Restart the PC, close other programs
 (browsers use a lot of RAM) and try again. If it keeps happening, pick a smaller size (Q2_0 or IQ2_XS).

@@ -164,3 +180,10 @@ The full story is in the [paper](docs/paper/Strata-Paper.pdf) and the [details](
 - Built with parts of [llama.cpp / ggml](https://github.com/ggml-org/llama.cpp) (MIT). Ideas from
  [Splash](https://github.com/incoai/splash), [ninfer](https://github.com/Neroued/ninfer) and
  [HyperQwen](https://github.com/syv-ai/HyperQwen). More in the [details](docs/DETAILS.md#credits-and-licenses).
+
+## License
+
+Strata is open source under the [MIT License](LICENSE). A few parts carry their own licenses: `third_party/ggml`
+(MIT, llama.cpp / ggml), the web app's font (SIL Open Font License 1.1) and the experimental speed projection's
+vector in `data/experimental-speed-projection` (Qwen Community License 1.0, from the model's activations). The
+models are not part of this repository; each model's own license applies to its files.
@@ -0,0 +1,44 @@
+# The GSQ-RCO Coder (256 of 512 experts) on Strata
+
+RTX 5070 12 GB, Ryzen 5 7600, 64 GB DDR5-5200, Windows 10. Engine 0.1.15 + PR #54 (the expert count from the model
+file), the settings setup writes (`--prefill auto`, 8-bit KV above 4K, KV streaming from 64K, MTP, greedy, 256 tokens),
+`data/expert-profile-coder.bin`, the same code-agent prompts as `2026-09-28-speed-0114`.
+
+| | 1K | 4K | 32K | 64K | 128K | 262K |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: |
+| Prompt, tokens/s | 599 | 1,152 | 1,298 | 1,350 | 1,266 | 1,034 |
+| Output, tokens/s | 53.3 | 50.6 | 53.3 | 50.8 | 44.0 | 42.8 |
+| Draft acceptance | 0.720 | 0.707 | 0.785 | 0.696 | 0.619 | 0.684 |
+| Experts in VRAM | 2,570 | 2,615 | 2,418 | 2,382 | 2,326 | 2,208 |
+
+Every run in [`matrix.json`](matrix.json). The 128K and 262K answers were read: coherent (they list and discuss the
+files of the prompt).
+
+## Why it is fast
+
+The kept experts are stored like IQ3_S's (IQ2_S / IQ3_XXS / IQ3_S / IQ4_XS gate and up, IQ4_NL / Q2_0 down; "IQ1_M" is
+the release's name for 1.89 bits per *original* parameter), so a token costs about what IQ3_S costs where the expert is
+on the CPU. But there are half as many experts: the 23 GB of them pin completely (`cudaHostRegister` of the whole
+arena), a 12 GB card holds ~20% of them instead of ~10%, and the prompt path streams half as many.
+
+## The expert profile
+
+`data/expert-profile-coder.bin` (48 x 256, all 12,288 pairs ranked): the shipped 48 x 512 ranking, each pair mapped to
+its kept expert's new index through section 2 of the release's `tensor-allocation/*.rco-allocation.txt` (the kept
+experts in their original order; gate, up and down agree on which were kept), pruned pairs dropped. 5,492 of the
+shipped top 8,000 survive (as PR #54 found). With 2,537 slots on the 12 GB card, 72% of the expert reads of a 4K code
+prompt hit the GPU - a wrong index mapping would give about the slot share, 21%.
+
+## Images
+
+The Coder ships the original's vision encoder (the same file). With `--vision` on the GPU it described a Strata
+Monitor screenshot correctly, windows, labels and the numbers of the dashboard and the requests table (1,079 prompt
+tokens with the image, 2,012 generated at ~50 tokens/s).
+
+## The experimental speed projection
+
+The vector was made from the full model's activations; the Coder has the same residual stream. On the Coder it loads
+(41 layers steered), costs no measurable speed (47.7-55.7 tokens/s with it on or off), and both settings write correct
+answers to three coding prompts. What it is for - fewer refusals - was not measured. (Two runs of the same prompt with
+it off also differ after ~130-200 characters in serve mode, so this check cannot show how far it moves ordinary text;
+the teacher-forced method of `2026-09-27-esp` needs the token-by-token prompt path, which native packs do not run.)
@@ -0,0 +1,86 @@
+[
+ {
+  "model": "CODER",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 18.0,
+  "prefill_tok_s": 599.05,
+  "ttft_s": 2.54,
+  "decode_tok_s": 53.33,
+  "decoded": 256,
+  "spec_accept": 0.72,
+  "tokens_per_round": 2.86,
+  "vram_slots": 2570
+ },
+ {
+  "model": "CODER",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 19.6,
+  "prefill_tok_s": 1152.17,
+  "ttft_s": 3.94,
+  "decode_tok_s": 50.6,
+  "decoded": 256,
+  "spec_accept": 0.707,
+  "tokens_per_round": 2.88,
+  "vram_slots": 2615
+ },
+ {
+  "model": "CODER",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 40.7,
+  "prefill_tok_s": 1297.8,
+  "ttft_s": 25.17,
+  "decode_tok_s": 53.34,
+  "decoded": 256,
+  "spec_accept": 0.785,
+  "tokens_per_round": 3.02,
+  "vram_slots": 2418
+ },
+ {
+  "model": "CODER",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 75.5,
+  "prefill_tok_s": 1350.09,
+  "ttft_s": 49.01,
+  "decode_tok_s": 50.84,
+  "decoded": 256,
+  "spec_accept": 0.696,
+  "tokens_per_round": 2.61,
+  "vram_slots": 2382
+ },
+ {
+  "model": "CODER",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 121.5,
+  "prefill_tok_s": 1266.31,
+  "ttft_s": 102.95,
+  "decode_tok_s": 44.03,
+  "decoded": 256,
+  "spec_accept": 0.619,
+  "tokens_per_round": 2.33,
+  "vram_slots": 2326
+ },
+ {
+  "model": "CODER",
+  "tier": "262k",
+  "prompt_tokens": 259943,
+  "exit": 0,
+  "wall_s": 280.7,
+  "prefill_tok_s": 1033.91,
+  "ttft_s": 252.92,
+  "decode_tok_s": 42.8,
+  "decoded": 256,
+  "spec_accept": 0.684,
+  "tokens_per_round": 2.6,
+  "vram_slots": 2208
+ }
+]
@@ -0,0 +1,32 @@
+# Speed with engine 0.1.14
+
+RTX 5070 12 GB (PCIe 5.0 x16), Ryzen 5 7600, 64 GB DDR5-5200, Windows 10. The ready-made 0.1.14 engine, one-shot runs
+with what setup writes: `--prefill auto`, the shipped expert profile, `--expert-cache auto`, 8-bit KV above 4K, KV
+streaming (`--kv-resident 32768`) from 64K, MTP (`--spec 4 --spec-min-p 0.5`), greedy, 256 generated tokens. One
+code-agent prompt per length (`bench/prompts/*/…-code-agent.ids`, the same prompts as `2026-09-24-final`).
+
+The tables: [`matrix.md`](matrix.md) (prompt and output speed, draft acceptance, experts in VRAM); every run in
+[`matrix.json`](matrix.json).
+
+- **IQ2_XS** was measured with Swift 1.5's IQ2_XS (the original's file was not on the machine; the two run at the same
+  speed: 4K 465 vs 467 prompt tokens/s in `2026-09-24`).
+- **IQ3_S at 128K** comes from a first pass during which a large download was running (the clean rerun of that row
+  was stopped when Windows ran short of memory); the other rows are from runs on an otherwise idle machine.
+- **IQ3_XXS and IQ3_S at 262K** are not measured: with 43 / 50 GB of experts, a 260K context exceeds 64 GB of RAM.
+
+## Against 0.1.12
+
+Prompts are 1.2x (1K) to 2.4x (32K-128K) faster: the 0.1.13 prompt path (chunks of up to 8,192 tokens, MMQ experts,
+streamed expert copies; see `2026-09-28-prefill-speed`).
+
+Output speed moves with the text: speculative decoding is faster when more drafts are accepted, and the new prompt
+path's rounding changes the answer, so the same prompt can land a few percent either way (Q2_0 32K: acceptance 0.865
+then, 0.709 now). The engine itself is not slower: on the 4K prompt, run back to back on the same machine state,
+0.1.14 writes 88.5 tokens/s and 0.1.12 85.7 (3 runs each, same free VRAM).
+
+## Found on the way
+
+IQ3_XXS at 64K and 128K ended with `out of memory: cudaFuncSetAttribute` in the MMQ launch: CUDA loads a kernel's code
+when it is first used, and mid-prompt there was no VRAM left for it. Loading every kernel at start
+(`CUDA_MODULE_LOADING=EAGER`, which the engine now sets itself) fixes it for ~30 MB of VRAM; the IQ3_XXS 64K/128K rows
+were measured with that fix.
@@ -0,0 +1,322 @@
+[
+ {
+  "model": "Q2_0",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 41.0,
+  "prefill_tok_s": 494.43,
+  "ttft_s": 3.55,
+  "decode_tok_s": 84.26,
+  "decoded": 256,
+  "spec_accept": 0.709,
+  "tokens_per_round": 2.72,
+  "vram_slots": 4176
+ },
+ {
+  "model": "Q2_0",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 41.2,
+  "prefill_tok_s": 1006.63,
+  "ttft_s": 5.05,
+  "decode_tok_s": 90.26,
+  "decoded": 256,
+  "spec_accept": 0.846,
+  "tokens_per_round": 3.23,
+  "vram_slots": 4123
+ },
+ {
+  "model": "Q2_0",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 63.1,
+  "prefill_tok_s": 1307.6,
+  "ttft_s": 25.63,
+  "decode_tok_s": 73.61,
+  "decoded": 256,
+  "spec_accept": 0.709,
+  "tokens_per_round": 2.68,
+  "vram_slots": 3833
+ },
+ {
+  "model": "Q2_0",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 89.0,
+  "prefill_tok_s": 1293.5,
+  "ttft_s": 51.12,
+  "decode_tok_s": 69.02,
+  "decoded": 256,
+  "spec_accept": 0.679,
+  "tokens_per_round": 2.62,
+  "vram_slots": 3780
+ },
+ {
+  "model": "Q2_0",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 147.4,
+  "prefill_tok_s": 1208.46,
+  "ttft_s": 107.81,
+  "decode_tok_s": 67.21,
+  "decoded": 256,
+  "spec_accept": 0.691,
+  "tokens_per_round": 2.64,
+  "vram_slots": 3697
+ },
+ {
+  "model": "Q2_0",
+  "tier": "262k",
+  "prompt_tokens": 259943,
+  "exit": 0,
+  "wall_s": 309.8,
+  "prefill_tok_s": 966.84,
+  "ttft_s": 270.35,
+  "decode_tok_s": 60.3,
+  "decoded": 256,
+  "spec_accept": 0.623,
+  "tokens_per_round": 2.43,
+  "vram_slots": 3525
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 40.4,
+  "prefill_tok_s": 461.28,
+  "ttft_s": 2.28,
+  "decode_tok_s": 74.38,
+  "decoded": 256,
+  "spec_accept": 0.772,
+  "tokens_per_round": 2.84,
+  "vram_slots": 4386,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 41.8,
+  "prefill_tok_s": 810.71,
+  "ttft_s": 4.48,
+  "decode_tok_s": 73.84,
+  "decoded": 256,
+  "spec_accept": 0.79,
+  "tokens_per_round": 3.24,
+  "vram_slots": 4336,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 65.0,
+  "prefill_tok_s": 1237.98,
+  "ttft_s": 25.59,
+  "decode_tok_s": 71.49,
+  "decoded": 256,
+  "spec_accept": 0.856,
+  "tokens_per_round": 3.28,
+  "vram_slots": 4057,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 97.1,
+  "prefill_tok_s": 1135.63,
+  "ttft_s": 56.59,
+  "decode_tok_s": 64.29,
+  "decoded": 256,
+  "spec_accept": 0.707,
+  "tokens_per_round": 2.76,
+  "vram_slots": 4005,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 158.0,
+  "prefill_tok_s": 1070.61,
+  "ttft_s": 120.1,
+  "decode_tok_s": 59.76,
+  "decoded": 256,
+  "spec_accept": 0.726,
+  "tokens_per_round": 2.69,
+  "vram_slots": 3923,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "262k",
+  "prompt_tokens": 259943,
+  "exit": 0,
+  "wall_s": 332.4,
+  "prefill_tok_s": 886.2,
+  "ttft_s": 293.43,
+  "decode_tok_s": 52.83,
+  "decoded": 256,
+  "spec_accept": 0.628,
+  "tokens_per_round": 2.53,
+  "vram_slots": 3565,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 44.1,
+  "prefill_tok_s": 414.55,
+  "ttft_s": 2.54,
+  "decode_tok_s": 60.26,
+  "decoded": 256,
+  "spec_accept": 0.864,
+  "tokens_per_round": 3.52,
+  "vram_slots": 3304,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 45.8,
+  "prefill_tok_s": 769.57,
+  "ttft_s": 4.73,
+  "decode_tok_s": 62.07,
+  "decoded": 256,
+  "spec_accept": 0.855,
+  "tokens_per_round": 3.24,
+  "vram_slots": 3263,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 69.7,
+  "prefill_tok_s": 1107.75,
+  "ttft_s": 28.61,
+  "decode_tok_s": 51.4,
+  "decoded": 256,
+  "spec_accept": 0.678,
+  "tokens_per_round": 2.6,
+  "vram_slots": 3031,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 104.9,
+  "prefill_tok_s": 1065.31,
+  "ttft_s": 61.77,
+  "decode_tok_s": 49.95,
+  "decoded": 256,
+  "spec_accept": 0.741,
+  "tokens_per_round": 2.94,
+  "vram_slots": 2965,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 172.1,
+  "prefill_tok_s": 1015.11,
+  "ttft_s": 128.12,
+  "decode_tok_s": 45.78,
+  "decoded": 256,
+  "spec_accept": 0.665,
+  "tokens_per_round": 2.35,
+  "vram_slots": 2901,
+  "note": "first pass"
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 56.2,
+  "prefill_tok_s": 396.13,
+  "ttft_s": 4.14,
+  "decode_tok_s": 51.52,
+  "decoded": 256,
+  "spec_accept": 0.826,
+  "tokens_per_round": 3.12,
+  "vram_slots": 2617
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 57.5,
+  "prefill_tok_s": 736.9,
+  "ttft_s": 6.35,
+  "decode_tok_s": 51.6,
+  "decoded": 256,
+  "spec_accept": 0.747,
+  "tokens_per_round": 2.89,
+  "vram_slots": 2580
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 82.5,
+  "prefill_tok_s": 1069.83,
+  "ttft_s": 31.33,
+  "decode_tok_s": 48.19,
+  "decoded": 256,
+  "spec_accept": 0.772,
+  "tokens_per_round": 2.94,
+  "vram_slots": 2383
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 115.0,
+  "prefill_tok_s": 1069.8,
+  "ttft_s": 61.47,
+  "decode_tok_s": 48.75,
+  "decoded": 256,
+  "spec_accept": 0.786,
+  "tokens_per_round": 2.9,
+  "vram_slots": 2347
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 191.8,
+  "prefill_tok_s": 931.09,
+  "ttft_s": 138.11,
+  "decode_tok_s": 40.48,
+  "decoded": 256,
+  "spec_accept": 0.661,
+  "tokens_per_round": 2.71,
+  "vram_slots": 2310,
+  "note": "first pass (a download was running)"
+ }
+]
@@ -0,0 +1,35 @@
+### Prompt processing (tok/s)
+
+| Model | 1K | 4K | 32K | 64K | 128K | 262K |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: |
+| Q2_0 | 494 | 1,007 | 1,308 | 1,294 | 1,208 | 967 |
+| IQ2_XS | 461 | 811 | 1,238 | 1,136 | 1,071 | 886 |
+| IQ3_XXS | 415 | 770 | 1,108 | 1,065 | 1,015 | - |
+| IQ3_S | 396 | 737 | 1,070 | 1,070 | 931 | - |
+
+### Output (tok/s)
+
+| Model | 1K | 4K | 32K | 64K | 128K | 262K |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: |
+| Q2_0 | 84.3 | 90.3 | 73.6 | 69.0 | 67.2 | 60.3 |
+| IQ2_XS | 74.4 | 73.8 | 71.5 | 64.3 | 59.8 | 52.8 |
+| IQ3_XXS | 60.3 | 62.1 | 51.4 | 50.0 | 45.8 | - |
+| IQ3_S | 51.5 | 51.6 | 48.2 | 48.8 | 40.5 | - |
+
+### Draft acceptance
+
+| Model | 1K | 4K | 32K | 64K | 128K | 262K |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: |
+| Q2_0 | 0.709 | 0.846 | 0.709 | 0.679 | 0.691 | 0.623 |
+| IQ2_XS | 0.772 | 0.790 | 0.856 | 0.707 | 0.726 | 0.628 |
+| IQ3_XXS | 0.864 | 0.855 | 0.678 | 0.741 | 0.665 | - |
+| IQ3_S | 0.826 | 0.747 | 0.772 | 0.786 | 0.661 | - |
+
+### Experts in VRAM
+
+| Model | 1K | 4K | 32K | 64K | 128K | 262K |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: |
+| Q2_0 | 4,176 | 4,123 | 3,833 | 3,780 | 3,697 | 3,525 |
+| IQ2_XS | 4,386 | 4,336 | 4,057 | 4,005 | 3,923 | 3,565 |
+| IQ3_XXS | 3,304 | 3,263 | 3,031 | 2,965 | 2,901 | - |
+| IQ3_S | 2,617 | 2,580 | 2,383 | 2,347 | 2,310 | - |
@@ -11,26 +11,36 @@ New here? Start with the [README](../README.md) - it has everything you need to

 ## Speed (measured)

-RTX 5070 **12 GB**, Ryzen 5 7600 (6 cores), 64 GB DDR5-5200, Windows. One code-agent prompt per length, 256 generated
-tokens, MTP speculative decoding on. "262K" is the model's full context window (a 259,943-token prompt).
+RTX 5070 **12 GB**, Ryzen 5 7600 (6 cores), 64 GB DDR5-5200, Windows, engine 0.1.14 with the settings setup writes
+(`--prefill auto`, 8-bit KV above 4K, KV streaming from 64K). One code-agent prompt per length, 256 generated tokens,
+MTP speculative decoding on. "262K" is the model's full context window (a 259,943-token prompt). The IQ2_XS row was
+measured with Swift 1.5's IQ2_XS, which runs at the original's speed.

 ### Prompt processing (tokens/s)

 | Model | 1K | 4K | 32K | 64K | 128K | 262K |
 | --- | ---: | ---: | ---: | ---: | ---: | ---: |
-| **Q2_0** | 389 | 539 | 571 | 561 | 543 | 496 |
-| **IQ2_XS** | 332 | 463 | 495 | 486 | 472 | 437 |
-| **IQ3_XXS** | 285 | 410 | 435 | 427 | 414 | - |
-| **IQ3_S** | 260 | 374 | 397 | - | 378 | - |
+| **Q2_0** | 494 | 1,007 | 1,308 | 1,294 | 1,208 | 967 |
+| **IQ2_XS** | 461 | 811 | 1,238 | 1,136 | 1,071 | 886 |
+| **IQ3_XXS** | 415 | 770 | 1,108 | 1,065 | 1,015 | - |
+| **IQ3_S** | 396 | 737 | 1,070 | 1,070 | 931 | - |
+| **Coder** | 599 | 1,152 | 1,298 | 1,350 | 1,266 | 1,034 |

 ### Output (tokens/s)

 | Model | 1K | 4K | 32K | 64K | 128K | 262K |
 | --- | ---: | ---: | ---: | ---: | ---: | ---: |
-| **Q2_0** | 88.7 | 94.6 | 87.5 | 76.4 | 65.1 | 56.3 |
-| **IQ2_XS** | 82.0 | 78.0 | 65.3 | 63.7 | 52.0 | 48.0 |
-| **IQ3_XXS** | 64.6 | 65.6 | 57.3 | 54.4 | 44.8 | - |
-| **IQ3_S** | 51.2 | 54.4 | 51.1 | - | 42.2 | - |
+| **Q2_0** | 84.3 | 90.3 | 73.6 | 69.0 | 67.2 | 60.3 |
+| **IQ2_XS** | 74.4 | 73.8 | 71.5 | 64.3 | 59.8 | 52.8 |
+| **IQ3_XXS** | 60.3 | 62.1 | 51.4 | 50.0 | 45.8 | - |
+| **IQ3_S** | 51.5 | 51.6 | 48.2 | 48.8 | 40.5 | - |
+| **Coder** | 53.3 | 50.6 | 53.3 | 50.8 | 44.0 | 42.8 |
+
+Output speed depends on the text as well: speculative decoding runs faster when more of the drafted tokens are
+accepted, so a different answer to the same prompt moves it by several percent. Run back to back on the 4K prompt,
+0.1.14 writes 88.5 tokens/s and 0.1.12 85.7. The numbers before 0.1.13 (prompts about half as fast):
+[`bench/results/2026-09-24-final`](../bench/results/2026-09-24-final/matrix.md); these:
+[`bench/results/2026-09-28-speed-0114`](../bench/results/2026-09-28-speed-0114/README.md).

 IQ3_XXS and IQ3_S at 262K are not measured: with their 43 / 50 GB of experts, a 260K-token context brings a 64 GB PC
 to its memory limit. Use up to 128K with them on 64 GB (setup caps it). IQ3_S (engine 0.1.4 or newer) is only published
@@ -47,8 +57,8 @@ halves the KV cache's memory with a Hadamard rotation before 4-bit rounding (PR
 is measurably less precise on long documents (perplexity +8-12%; needle tests still pass). 8-bit stays the default.
 Details: [`bench/results/2026-09-27-kv-q4`](../bench/results/2026-09-27-kv-q4/README.md).

-Time to first token is prompt length / prompt speed: about 7 s at 4K, 55 s at 32K, 4 minutes at 128K and 9 minutes at
-262K (engine 0.1.12; long prompts are about twice as fast since 0.1.13, below).
+Time to first token is prompt length / prompt speed: with Q2_0 about 4 s at 4K, 25 s at 32K, under 2 minutes at 128K
+and 4.5 minutes at 262K (engine 0.1.13 made long prompts about twice as fast, below).

 **Faster prompts (engine 0.1.13):** the prompt is read in chunks of up to 8,192 tokens instead of 2,048 (`--prefill
 auto`: the largest chunk whose buffers fit in the expert-cache slots it borrows, and a request borrows only what its
@@ -66,6 +76,8 @@ switch to `--prefill auto` the next time START-HERE / setup.sh starts them. The

 Not measured - estimated from the runs above (same CPU and 64 GB RAM): the GPU part scaled by memory bandwidth, the CPU
 part by how many more experts the card's VRAM holds. Treat as **±20%**. Numbers are *prompt / output* tokens/s.
+The prompt figures predate engine 0.1.13, which about doubled prompt speed on the measured card; how much of that a
+card gains depends on its PCIe link (the experts stream over it), so they are still the older estimates.

 | GPU | Model | 1K | 4K | 32K | 64K | 128K | 262K |
 | --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: |
@@ -77,7 +89,9 @@ part by how many more experts the card's VRAM holds. Treat as **±20%**. Numbers
 |  | IQ3_XXS | ~260 / ~106 | ~374 / ~103 | ~396 / ~89 | ~390 / ~85 | ~378 / ~71 | - |

 More VRAM matters more than a faster GPU: every extra GB holds ~700 more experts, and every expert on the GPU is one the
-CPU does not have to compute. A 3090's 24 GB takes most of the CPU work away.
+CPU does not have to compute. A 3090's 24 GB takes most of the CPU work away. (Since 0.1.14 the expert profile ranks
+all 24,576 experts; before, the cache stopped at 8,000, about 10-14 GB. `tools/make_profile.py` builds a profile from
+your own prompts: run the engine once with `--dump-routing trace.bin`, see the tool's help.)

 ## Which model?

@@ -92,6 +106,23 @@ of [Qwen3.8-Flash-Next](https://huggingface.co/Qwen/Qwen3.8-Flash-Next).

 With 64 GB of RAM all three fit (close the browser for IQ3_XXS, and keep its context at 128K or less). With 48 GB only Q2_0 / IQ2_XS may fit. 32 GB is not enough.

+### Or: the Coder (half the experts, for code)
+
+**[Qwen3.8-Flash-Next GSQ-RCO Coder](https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF)** is
+ISTA-DASLab's expert-pruned release: 256 of each layer's 512 experts are kept (still 10 active per token), chosen with
+RCO on code, agentic and vision calibration data; its authors report 91.3% of the full model's SWE-bench Verified and
+98.7% of LiveCodeBench v6. One size, named IQ1_M for its 1.89 bits per *original* parameter; the kept experts are
+stored like IQ3_S (IQ2_S-IQ4_XS gate/up, IQ4_NL/Q2_0 down). Shard 1 is 29.6 GB (experts: 23 GB of RAM), so it runs
+on **32 GB of RAM**, and at 262K on 64 GB. Its shard 2 and its vision encoder are the original's files: with the
+original installed, setup downloads only shard 1. Strata ships its expert profile (`data/expert-profile-coder.bin`,
+the shipped ranking mapped onto the kept experts through the release's `rco-allocation.txt`: 72% of the expert
+reads hit the GPU on a 12 GB card). Images work; the experimental speed projection loads and runs on it (it was made
+for the full model).
+
+```
+START-HERE.bat --setup --family coder
+```
+
 ### Or: Swift 1.5 (a fine-tune that thinks shorter)

 The setup's first question also offers **[Swift 1.5](https://huggingface.co/ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF)**,
@@ -122,7 +153,11 @@ You need **only an NVIDIA driver** (version 580 or newer; update it with the NVI
 | Disk | ~70-80 GB free for the model, ~6 GB for the MTP layer (+1 GB with images). **Q2_0 on an AVX-512 CPU** also writes a one-time ~40 GB copy of its experts for the fast CPU kernel. An NVMe SSD is strongly recommended. |
 | OS | Windows 10/11, or Linux (Ubuntu 22.04/24.04 get everything installed automatically). |

-What the first start installs, all inside this folder (`.venv/`, `engine/`, `third_party/`, `models/`, `packs/`, `mtp/`):
+What the first start installs: in this folder `.venv/`, `engine/` and `third_party/`; the model files (`models/`,
+`packs/`, `mtp/`, 70-120 GB) in **`Strata-data` next to this folder**, so a new copy of Strata (an update unzipped
+elsewhere) finds them and sets itself up the same way. The place is remembered per user (`%APPDATA%\Strata\settings.json`,
+`~/.config/strata/settings.json`); `--data-dir` chooses another. Installs from before 0.1.16 are moved there by the next
+start (a rename on the same drive; files on another drive are used where they are).
 Python 3.12 if you have none (for your user account, no admin), a private Python environment, NVIDIA's CUDA libraries
 (from pip, ~0.4 GB), the ready-made Strata engine for RTX 30/40/50, the model and the MTP draft layer. If no
 ready-made engine fits your PC, it offers to install the build tools (Visual Studio Build Tools + CUDA Toolkit on
@@ -160,6 +195,7 @@ downloaded again. Closing the window stops the model.
 START-HERE.bat --setup                          install another model, or change context / images
 START-HERE.bat --model IQ2_XS --context 32768 --vision yes --yes     no questions
 START-HERE.bat --gguf-dir D:\models\IQ2_XS       use GGUF files you already have
+START-HERE.bat --data-dir E:\Strata-data         keep the model files somewhere else
 START-HERE.bat --port 8081                      another port
 ```

@@ -184,6 +220,15 @@ for the build tools), and the same start: `http://127.0.0.1:8080`. Later runs of
 start the model directly. Options as on Windows (`./setup.sh --setup`, `--model Q2_0 --yes`, `--gguf-dir /data/Q2_0`).
 Terminal chat: `.venv/bin/python chat.py`.

+- **Updating:** `git pull`, then `./setup.sh`: it compiles the engine again when its source changed (a minute or
+  two for the changed files). If that compile fails, it says so and starts the engine you had.
+- **Other distributions** (Arch, Fedora, ...): install the C++ compiler and the CUDA Toolkit 13 with your package
+  manager first (Arch: `sudo pacman -S base-devel cuda`); setup finds `nvcc` on PATH, in `/usr/local/cuda*` and in
+  `/opt/cuda*`, and does the rest.
+- **WSL** works (Ubuntu 24.04 tested), with one limit: the NVIDIA driver pins only about 1 GB of RAM there, so KV
+  streaming (`--kv-resident`) is off and the KV cache stays in VRAM, and the experts are copied to the GPU from
+  unpinned RAM (slower prompts than native Linux).
+
 ---

 ## Using it
@@ -384,7 +429,7 @@ the document, +0.4% on the chat. Details: `bench/results/2026-09-27-esp/`.
 | `port 8080 is already in use` | Strata is already running (look for its window), or another program uses the port: `START-HERE.bat --port 8081`. |
 | `cudaHostRegister ... out of memory` in the log | Normal on Windows: the engine pins the experts in per-layer slices instead. Only a problem if the load then fails. |
 | The first start takes minutes | It is reading 34-55 GB into RAM; the second start is faster while the files are in the OS cache. |
-| The PC freezes for a few minutes at the first start | Normal the first time: the engine loads the experts into RAM, pins part of it for the GPU and sizes the expert cache. Wait; don't close the window. Still frozen after 10 minutes: restart the PC, close other programs, try again, or pick a smaller size. |
+| The PC freezes for a few minutes at the start | Normal, most of all the first time (the server window says when it happens): the engine loads the experts into RAM, pins part of it for the GPU and sizes the expert cache. Wait; don't close the window. Still frozen after 10 minutes: restart the PC, close other programs, try again, or pick a smaller size. |
 | `the engine stopped unexpectedly (exit code ...)` | The engine process ended mid-answer - usually out of RAM (Linux ends the biggest program: `sudo dmesg \| grep -i -E 'killed process\|out of memory'`). The next request starts it again by itself. If it repeats: close other programs or pick a smaller size. The server also warns at start when the model's experts leave less than ~6 GB of RAM for everything else. |
 | Slow output, disk light busy | Not enough free RAM: close other programs, or choose Q2_0 / IQ2_XS. |
 | `prompt ... exceeds the context` | The request is longer than the context you chose: run setup again with a bigger `--context`. |
@@ -393,7 +438,8 @@ the document, +0.4% on the chat. Details: `bench/results/2026-09-27-esp/`.
 | A picture is refused or `cannot read the image` | The file is not a picture Pillow can open (JPEG, PNG, WebP, GIF, BMP, TIFF, AVIF work). |
 | Pictures are slow (10-30 s) | The encoder runs on the CPU: run setup again with `--vision gpu` (needs ~1.4 GB of VRAM). |
 | A request never finishes: "reading the prompt", GPU "100%" at low power | The GPU ran out of VRAM (engines before 0.1.9 could end with ~30 MiB free at large contexts). Run `START-HERE.bat` once to get engine 0.1.9 or newer; the log then says `... MiB of VRAM free with everything loaded` (a few hundred) and names the `--vram-reserve-mib` to add if it is low. |
-| Generation stops mid-answer, GPU "100%", one CPU core busy | Fixed in engine 0.1.12 (issue #29, a race in the CPU expert pool on big-VRAM cards). Since then a request that stops moving ends with an error instead of hanging (after 2 minutes; 1 minute from 0.1.13): the log says `no progress for ... s ... (issue #29)` with where it stopped, and the next request starts the engine again. If you see that line, please open an issue with it. Engine 0.1.13 adds a stall report under it (what every expert-pool thread and the GPU handshake were doing, memory and page faults) and, on Windows, a `strata-stall-<pid>.dmp` file with every thread's stack: attach both. (`STRATA_WATCHDOG_S` sets the time in seconds; 0 turns it off.) |
+| Generation stops mid-answer, GPU "100%", one CPU core busy | Fixed in engine 0.1.12 (issue #29, a race in the CPU expert pool on big-VRAM cards). Since then a request that stops moving ends with an error instead of hanging (after 2 minutes; 1 minute from 0.1.13): the log says `no progress for ... s ... (issue #29)` with where it stopped, and the next request starts the engine again. If you see that line, please open an issue with it. Engine 0.1.13 adds a stall report under it (what every expert-pool thread and the GPU handshake were doing, memory and page faults) and, on Windows, a `strata-stall-<pid>.dmp` file with every thread's stack: attach both. (`STRATA_WATCHDOG_S` sets the time in seconds; 0 turns it off.) Engine 0.1.14 fixes the stall those reports found (issue #31: with the IQ packs the host could wait forever inside the NVIDIA driver while copying experts in a verify window; the experts are now copied by a GPU kernel, `--pcie-mode dma` restores the old way). |
+| `out of memory: cudaFuncSetAttribute` in the log (IQ3_XXS, long prompt) | Fixed in engine 0.1.15: CUDA loaded a kernel's code when it was first needed, and mid-prompt there was no VRAM left for it. Run `START-HERE.bat` (Windows) or `./setup.sh` (Linux) once to update. |
 | Anything else | The engine log is `strata-<model>.log` in this folder. |

 ---
@@ -420,8 +466,12 @@ The full story, with measurements, bottlenecks and what comes next: **[docs/pape

 ## Credits and licenses

+Strata itself: [MIT](../LICENSE). The model files are not part of it; their licenses apply to them (below).
+
 - Model: [Qwen/Qwen3.8-Flash-Next](https://huggingface.co/Qwen/Qwen3.8-Flash-Next) by the Qwen team; quantizations:
  [ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF](https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF).
+  The Coder: [ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF](https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF)
+  (Apache-2.0 per its card); its support in Strata came from @pjgmobile's PR #54.
  Swift 1.5: [ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF](https://huggingface.co/ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF)
  by UkisAI. Their licenses apply to the weights.
 - [llama.cpp / ggml](https://github.com/ggml-org/llama.cpp) (MIT): the i-quant formats, the GPU dot products and
@@ -432,3 +482,5 @@ The full story, with measurements, bottlenecks and what comes next: **[docs/pape
  [HyperQwen](https://github.com/syv-ai/HyperQwen); references in the paper.
 - The web app's font: [Outfit](https://github.com/Outfitio/Outfit-Fonts) (SIL Open Font License 1.1, see
  `serve/web/fonts/OFL.txt`). Its Monitor tab started from @code-martin's dashboard idea (PR #22).
+- The experimental speed projection's vector (`data/experimental-speed-projection/`): Qwen Community License 1.0,
+  made from the model's activations (see its README).
@@ -454,8 +454,9 @@ private:
 // ---- architecture guard (P1.S2). The engine is specialised to ONE model; anything else must be
 // refused with a precise error rather than silently mis-run.
 struct Qwen4ExpGuard {
-    uint32_t block_count = 48, hidden = 2560, experts = 512, experts_used = 10, head_count = 24,
-             head_count_kv = 2;
+    uint32_t block_count = 48, hidden = 2560, experts = 0, experts_used = 0, head_count = 24,
+             head_count_kv = 2;   // 0 = presence-only: pruned variants (GSQ-RCO Coder) legitimately ship
+                                  // fewer experts than the canonical 512; the graph reads the true value
 };

 inline std::string check_architecture(const GgufFile& g, const Qwen4ExpGuard& want = {}) {
@@ -477,7 +478,7 @@ inline std::string check_architecture(const GgufFile& g, const Qwen4ExpGuard& wa
    for (const auto& r : reqs) {
        const MetaValue* v = g.get(r.key);
        if (!v) return std::string("missing ") + r.key;
-        if (v->u != r.want)
+        if (r.want && v->u != r.want)
            return std::string(r.key) + " = " + std::to_string(v->u) + ", expected " + std::to_string(r.want);
    }
    return {}; // empty == ok
@@ -0,0 +1,26 @@
+// include/strata/kernels/cpu/iq_avx2.hpp - AVX-2 multi-token dot products for the i-quant expert
+// formats (IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S) against Q8_K activations (ggml's block_q8_K).
+//
+// The AVX-512 kernel's scheme on the CPUs without it (AMD Zen 2/3, Intel Core 12th-14th gen): the weights of a
+// 32-value chunk are decoded once per verify window and every token applies them with five instructions.
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+
+namespace strata::kernels::cpu {
+
+bool iq256_supported(int ggml_type) noexcept;
+/// ff[t][r] = silu(gate_r . a[t]) * (up_r . a[t]), rows [r0, r1); gate rows at blob, up rows at blob + up_off.
+void iq256_gu_rows(int ggml_type, const uint8_t* blob, size_t gu_row, size_t up_off, int n, const void* const* act,
+                   int nt, float* const* ff, int r0, int r1);
+/// out[t][r] = w_r . a[t], rows [r0, r1).
+void iq256_rows(int ggml_type, const uint8_t* w, size_t row_bytes, int n, const void* const* act, int nt,
+                float* const* out, int r0, int r1);
+
+/// out[t][r] = w_r . h[t] for IQ4_NL (type 20) rows against Q8_0 activations (ggml's block_q8_0).
+/// IQ4_NL is a 32-value-block format, so this does not go through iq256_rows (QK_K blocks, Q8_K acts).
+void iq4nl256_down_rows(const uint8_t* w, size_t row_bytes, int n, const void* const* hq, int nt,
+                        float* const* out, int r0, int r1);
+
+}  // namespace strata::kernels::cpu
@@ -37,7 +37,7 @@ void gdn_recurrence(float* state, const float* h, const float* gate, const float

 // ---- MoE
 /// softmax over 512, top-10 (ties to the lower id), weights renormalised over the ten (the native router).
-void route(const float* logits, int32_t* ids, float* weights, int64_t T, void* stream);
+void route(const float* logits, int32_t* ids, float* weights, int64_t T, int64_t n_expert, void* stream);
 /// Expert blob (Strata pack layout, Q2_0) -> BF16 matrices: gate/up interleaved [1280, 2560], down [2560, 640].
 void blob_dequant(const uint8_t* blob, uint16_t* gu16, uint16_t* down16, void* stream);
 /// h16[n, r] = fp16(silu(gu[n, 2r]) * gu[n, 2r + 1])   (the interleaved expert gate/up)
@@ -76,6 +76,57 @@ class EngineDied(RuntimeError):
    """The engine process ended in the middle of a request (issue #27: on Linux, the out-of-memory killer)."""


+def narrate_start(log_path: str, offset: int, args: list, done: threading.Event, heartbeat=20.0) -> None:
+    """While the engine starts, say in the server window what it is doing, from its log: the start reads tens of GB
+    into RAM and locks part of it for the GPU, and on many PCs everything is slow or frozen for a minute or more -
+    people closed the window thinking it had hung.  The warning comes at that step, not after it."""
+    gb = 0.0
+    if "--native" in args:                              # about the size of the experts it will read
+        try:
+            gb = os.path.getsize(args[args.index("--native") + 1]) / 1e9
+        except (OSError, IndexError):
+            pass
+    size = f"about {gb:.0f} GB" if gb >= 1 else "tens of GB"
+    t0 = last = time.time()
+    said = set()
+
+    def say(key, text):
+        nonlocal last
+        if key not in said:
+            said.add(key)
+            last = time.time()
+            print(text, flush=True)
+
+    say("weights", "[strata] starting the engine: reading the model's weights ...")
+    pos = offset
+    while not done.wait(0.5):
+        try:
+            with open(log_path, "rb") as f:
+                f.seek(pos)
+                chunk = f.read()
+        except OSError:
+            chunk = b""
+        if chunk.count(b"\n"):
+            cut = chunk.rfind(b"\n") + 1
+            pos += cut
+            for line in chunk[:cut].decode("utf-8", "replace").splitlines():
+                if "PLE on" in line or "expert arena:" in line:
+                    say("arena", f"[strata] loading the experts into RAM ({size}) and locking part of them for the GPU.\n"
+                                 "         YOUR PC CAN BE SLOW OR STOP RESPONDING FOR 1-3 MINUTES NOW - this is normal.\n"
+                                 "         Please wait and don't close this window; the browser opens when it is ready.")
+                elif " loaded " in line and "GiB at" in line:
+                    say("loaded", "[strata] experts loaded: " + line.split(" loaded ", 1)[1].strip() +
+                        f" ({time.time() - t0:.0f} s so far)")
+                elif "expert cache " in line and " slots, " in line and "auto" not in line:
+                    n = line.split("expert cache ", 1)[1].split(";")[0].replace(" slots,", " experts,").strip()
+                    say("cache", f"[strata] filling the GPU's expert cache ({n}) ...")
+                elif "session is up" in line:
+                    say("up", "[strata] almost ready ...")
+        if time.time() - last > heartbeat:
+            last = time.time()
+            print(f"[strata] still starting ({time.time() - t0:.0f} s) - please wait ...", flush=True)
+
+
 class StrataEngine:
    """The resident engine: `strata --serve` reads `GEN <max_new> <ids>` lines and streams `T <id>` lines, then
    `DONE ...`.  Requests are serialized by the service's FIFO, so one pipe is enough.
@@ -90,6 +141,10 @@ class StrataEngine:
        self.spawn = (exe, list(args), cwd, log, env)   # to start it again after it died (issue #27)
        self.log_path = log
        self.log = open(log, "a", encoding="utf-8") if log else subprocess.DEVNULL
+        loading = threading.Event()                     # set once READY: the narrator below stops
+        if log:
+            threading.Thread(target=narrate_start, args=(log, os.path.getsize(log), args, loading),
+                             daemon=True).start()
        self.proc = subprocess.Popen([exe, "--serve", *args], cwd=cwd, stdin=subprocess.PIPE, stdout=subprocess.PIPE,
                                     stderr=self.log, text=True, encoding="utf-8", bufsize=1, env=env)
        self.max_context = 0
@@ -111,6 +166,7 @@ class StrataEngine:
                self.max_context = int(f[1])
                self.can_stop = "stop" in f[2:]
                break
+        loading.set()
        if self.max_context <= 0:
            raise RuntimeError("the engine exited before it was ready" + (f" (see {log})" if log else ""))
        # the engine's stdout on a thread, so a request can wait with a timeout (heartbeats, cancel checks)
@@ -441,7 +497,10 @@ class Service:
        self.embeddings = threading.local()           # the current request's image embeddings file (GENI)
        self.api_key = ""                              # when set, /v1/* needs it (Bearer or x-api-key)
        self.status = {"busy": False, "queued": 0}      # GET /status: what the model is doing right now
-        self.history = collections.deque(maxlen=30)     # the last finished requests, newest last (GET /metrics)
+        self.history = collections.deque(maxlen=500)    # the last finished requests, newest last (GET /metrics)
+        # since the server started (the Monitor's totals, issue #35)
+        self.totals = {"since": time.time(), "requests": 0, "prompt_tokens": 0, "reused": 0, "output_tokens": 0,
+                       "prompt_ms": 0.0, "decode_ms": 0.0}
        self.status_lock = threading.Lock()
        self.stop_ids = set(tokenizer.encode(IM_END, parse_special=True) +
                            tokenizer.encode("<|endoftext|>", parse_special=True))
@@ -494,12 +553,13 @@ class Service:
            return 0.0
        return s["generated"] / max(1e-6, time.time() - s["first_token"])

-    def metrics(self) -> dict:
+    def metrics(self, all_requests=False) -> dict:
        """GET /metrics: what the Monitor tab shows - the engine's facts, what it is doing, the last requests, and
        the hardware (with a minute of history per series)."""
        with self.status_lock:
            s = dict(self.status)
            hist = list(self.history)
+            totals = dict(self.totals)
        now = time.time()
        progress = getattr(self.engine, "progress", None)
        if s.get("busy") and s.get("first_token") is None:
@@ -519,7 +579,9 @@ class Service:
        engine = {"model": self.model, "max_context": self.engine.max_context, "images": self.vision is not None,
                  **dict(getattr(self.engine, "info", {}) or {})}
        tel = self.telemetry.snapshot() if getattr(self, "telemetry", None) else {"now": {}, "history": {}, "static": {}}
-        return {"engine": engine, "live": live, "requests": hist[::-1], "hardware": tel["now"], "hardware_static":
+        return {"engine": engine, "live": live, "requests": hist[::-1][:None if all_requests else 12],
+                "requests_kept": len(hist), "totals": totals, "hardware": tel["now"],
+                "hardware_static":
                tel["static"], "history": tel["history"], "time": now}

    def prepare(self, messages, tools, kwargs, max_new=None):
@@ -594,7 +656,9 @@ class Service:
            s = dict(self.status)
        el = now - s.get("started", now)
        if s.get("first_token") is None:
-            print(f"[strata] reading the prompt: {s.get('prompt_tokens', 0)} tokens, {el:.0f} s so far", flush=True)
+            pr = getattr(self.engine, "progress", None)   # (position reached, prompt tokens): a reused prefix counts
+            done = f"{pr[0]:,} of {pr[1]:,}" if pr and pr[1] else f"{s.get('prompt_tokens', 0):,}"   # as read (#29)
+            print(f"[strata] reading the prompt: {done} tokens, {el:.0f} s so far", flush=True)
        else:
            rate = s["generated"] / max(1e-6, now - s["first_token"])
            print(f"[strata] {s['phase']}: {s['generated']} of max {s.get('max_tokens')} tokens, {rate:.1f} tok/s, "
@@ -656,6 +720,10 @@ class Service:
                          f"{' Its log: ' + self.engine.log_path if getattr(self.engine, 'log_path', None) else ''}",
                          flush=True)
                    raise
+                except ValueError as e:                 # the engine's ERR line (it may have ended after it)
+                    finish = "error"
+                    print(f"[strata] the engine reported an error: {e}", flush=True)
+                    raise
                finally:
                    gen.close()                         # STOP+drain to THIS request's DONE while still holding the
                    #                                     fifo, so a stop-token break can't leave the shared engine
@@ -679,6 +747,13 @@ class Service:
                        "prompt_ms": last.get("prompt_ms"), "decode_ms": last.get("decode_ms"),
                        "decode_tok_s": round(last["generated"] / (last["decode_ms"] / 1000), 1)
                        if n and last.get("generated") and last.get("decode_ms") else None})
+                    t = self.totals
+                    t["requests"] += 1
+                    t["prompt_tokens"] += len(ids)
+                    t["reused"] += last.get("reused") or 0
+                    t["output_tokens"] += n
+                    t["prompt_ms"] += last.get("prompt_ms") or 0.0
+                    t["decode_ms"] += last.get("decode_ms") or 0.0
                    now = time.time()
                    el = now - self.status.get("started", now)
                    ft = self.status.get("first_token")
@@ -930,7 +1005,8 @@ def make_handler(svc: Service):
                return
            if path == "/metrics":
                if self._authorized():
-                    self._json(200, svc.metrics())
+                    # the last 12 requests; `?requests=all` every one kept (the Monitor's "Show all", issue #35)
+                    self._json(200, svc.metrics(all_requests="requests=all" in self.path))
                return
            if path == "/settings":
                if self._authorized():
@@ -1035,6 +1111,9 @@ def make_handler(svc: Service):
            except EngineDied as e:                          # mid-stream: say so, then end the stream properly
                err = {"error": {"type": "server_error", "message": f"{e}; the next request restarts it"}}
                self.wfile.write(b"data: " + json.dumps(err).encode() + b"\n\ndata: [DONE]\n\n")
+            except ValueError as e:                          # the engine's ERR after the stream started: the
+                err = {"error": {"type": "server_error", "message": str(e)}}   # headers are sent, so no 400 now
+                self.wfile.write(b"data: " + json.dumps(err).encode() + b"\n\ndata: [DONE]\n\n")

        def _anthropic(self, req):
            req = svc.with_shared(req, "anthropic")
@@ -1062,6 +1141,9 @@ def make_handler(svc: Service):
            except EngineDied as e:                          # mid-stream: Anthropic's error event
                err = {"type": "error", "error": {"type": "api_error", "message": f"{e}; the next request restarts it"}}
                self.wfile.write(b"event: error\ndata: " + json.dumps(err).encode() + b"\n\n")
+            except ValueError as e:                          # the engine's ERR after the stream started
+                err = {"type": "error", "error": {"type": "api_error", "message": str(e)}}
+                self.wfile.write(b"event: error\ndata: " + json.dumps(err).encode() + b"\n\n")

    return Handler

@@ -221,6 +221,34 @@ class EngineDeath(unittest.TestCase):
            httpd.shutdown()
            httpd.server_close()

+    def test_engine_err_mid_stream(self):
+        """The engine's ERR line after the stream started reaches the client as an error event (it used to be a
+        400 written into the open stream, which clients read as an empty answer)."""
+        class ErrEngine(MockEngine):
+            def generate(self, ids, max_new, sampling, cancel, embeddings=None):
+                yield None                                  # a prompt-progress heartbeat: the stream has started
+                raise ValueError("verify: layer 31 never rang (an illegal memory access was encountered)")
+
+        tok = ByteTokenizer()
+        svc = Service(ErrEngine(tok, ANSWER, max_context=CTX), tok, ChatTemplate(ROOT / "serve/chat_template.jinja"))
+        httpd = serve(svc, port=0)
+        base = f"http://127.0.0.1:{httpd.server_address[1]}"
+        try:
+            for path, body in [("/v1/chat/completions", {"model": "m", "stream": True, "max_tokens": 20,
+                                                          "messages": [{"role": "user", "content": "hi"}]}),
+                               ("/v1/messages", {"model": "m", "stream": True, "max_tokens": 20,
+                                                 "messages": [{"role": "user", "content": "hi"}]})]:
+                req = urllib.request.Request(base + path, data=json.dumps(body).encode(),
+                                             headers={"Content-Type": "application/json"})
+                with urllib.request.urlopen(req, timeout=30) as r:
+                    text = r.read().decode()
+                self.assertIn("illegal memory access", text, path)
+                self.assertNotIn("HTTP/1", text, path)
+                self.assertEqual(svc.metrics()["requests"][0]["finish"], "error")
+        finally:
+            httpd.shutdown()
+            httpd.server_close()
+

 class SharedSettings(unittest.TestCase):
    """The web app's "Use for other apps too": POST /settings makes its Chat settings every client's defaults."""
@@ -94,6 +94,10 @@ main { flex: 1; min-height: 0; display: flex; }
           margin-top: var(--st-s-2); font-variant-numeric: tabular-nums; }
 .req-card { display: flex; flex-direction: column; gap: var(--st-s-3); min-width: 0; }
 .table-wrap { overflow-x: auto; }
+.req-head { display: flex; align-items: center; justify-content: space-between; gap: var(--st-s-2); }
+.table-wrap.all { max-height: 420px; overflow-y: auto; }
+.req-totals { margin: 0; font-size: var(--st-fs-sm); }
+.req-totals:empty { display: none; }

 /* ---------------------------------------------------------------- about */
 .about { max-width: 860px; }
@@ -143,9 +143,10 @@ function setMetric(key, value, unit, sub) {
 }

 let lastMetrics = null, metricsFailures = 0, keyWarned = false;
+let reqShowAll = false;   // the Monitor's request table: the last 12, or every one the server keeps (issue #35)
 async function poll() {
  try {
-    const r = await fetch("/metrics", {headers: headers()});
+    const r = await fetch(reqShowAll ? "/metrics?requests=all" : "/metrics", {headers: headers()});
    if (r.status === 401) {
      setPill("error", "API key needed");
      if (!keyWarned) { keyWarned = true; toast("warn", "API key needed", "This server needs a key: add it under About > Settings.", 6000); }
@@ -180,11 +181,20 @@ function render(m) {
    setPill("idle", "Idle");
  }
  if (live.queued > 0) setPill("queued", `${live.queued} queued`);
-  if (tab === "monitor") renderMonitor(live, hw, st, eng, h, last, m.requests || []);
+  if (tab === "monitor") renderMonitor(live, hw, st, eng, h, last, m.requests || [], m.totals, m.requests_kept);
  if (tab === "about") renderAbout(eng, hw, st);
 }

-function renderMonitor(live, hw, st, eng, h, last, requests) {
+function renderTotals(t) {
+  if (!t || !t.requests) return "";
+  const since = new Date(t.since * 1000).toLocaleString([], {weekday: "short", hour: "2-digit", minute: "2-digit"});
+  const read = t.prompt_tokens - t.reused;
+  const pSpeed = t.prompt_ms > 0 && read > 0 ? ` at ${fmt(read / (t.prompt_ms / 1000))} tok/s` : "";
+  const oSpeed = t.decode_ms > 0 && t.output_tokens > 0 ? ` at ${fmt(t.output_tokens / (t.decode_ms / 1000), 1)} tok/s` : "";
+  return `Since ${since}: ${fmt(t.requests)} requests · ${fmt(read)} prompt tokens read${pSpeed} (${fmt(t.reused)} reused) · ` +
+         `${fmt(t.output_tokens)} written${oSpeed}`;
+}
+function renderMonitor(live, hw, st, eng, h, last, requests, totals, kept) {
  // model state
  const on = live.queued > 0 ? "queued" : live.state;
  for (const b of document.querySelectorAll("#state-badges .st-badge")) b.classList.toggle("on", b.dataset.s === on || b.dataset.s === live.state);
@@ -268,7 +278,7 @@ function renderMonitor(live, hw, st, eng, h, last, requests) {
  } else {
    const badge = {stop: ["", "Done"], length: ["", "Max tokens"], cancel: ["st-badge--queued", "Stopped"],
                   disconnect: ["st-badge--queued", "Closed"], error: ["st-badge--error", "Error"]};
-    body.innerHTML = requests.slice(0, 12).map((r) => {
+    body.innerHTML = requests.slice(0, reqShowAll ? requests.length : 12).map((r) => {
      const [cls, text] = badge[r.finish] || ["", r.finish || "–"];
      const t = new Date(r.time * 1000).toLocaleTimeString([], {hour: "2-digit", minute: "2-digit", second: "2-digit"});
      const proj = r.projection == null ? "" : ` <span class="st-badge${r.projection ? " st-badge--reading" : ""}" title="experimental speed projection ${r.projection ? "on" : "off"}">${r.projection ? "ESP" : "stock"}</span>`;
@@ -277,6 +287,12 @@ function renderMonitor(live, hw, st, eng, h, last, requests) {
        <td class="num">${fmt(r.duration_s, 1)} s</td></tr>`;
    }).join("");
  }
+  const all = $("req-all");
+  kept = kept == null ? requests.length : kept;
+  all.hidden = kept <= 12;
+  all.textContent = reqShowAll ? "Show fewer" : `Show all (${kept})`;
+  $("req-wrap").classList.toggle("all", reqShowAll);
+  $("req-totals").textContent = renderTotals(totals);
 }

 function facts(el, rows) {
@@ -320,6 +336,7 @@ document.addEventListener("click", (e) => {
  const b = e.target.closest("[data-copy]");
  if (b) copyText(b.dataset.copy, b);
 });
+$("req-all").addEventListener("click", () => { reqShowAll = !reqShowAll; if (lastMetrics) render(lastMetrics); });

 // ------------------------------------------------------------------ Markdown (escaped first, then formatted)
 function inline(s) {
@@ -112,14 +112,16 @@
          <div class="st-progress" data-tone="warn"><div class="st-progress__bar" id="temp-bar" style="width:0%"></div></div>
        </div>
        <div class="st-card req-card">
-          <span class="card-title">Recent requests</span>
-          <div class="table-wrap">
+          <div class="req-head"><span class="card-title">Recent requests</span>
+            <button class="st-btn st-btn--secondary" id="req-all" hidden>Show all</button></div>
+          <div class="table-wrap" id="req-wrap">
            <table class="st-table">
              <thead><tr><th>Time</th><th>Status</th><th class="num">Prompt</th><th class="num">Reused</th><th class="num">Output</th>
                <th class="num">Tok/s</th><th class="num">Duration</th></tr></thead>
              <tbody id="req-body"><tr><td colspan="7" class="muted">No requests yet</td></tr></tbody>
            </table>
          </div>
+          <p class="muted req-totals" id="req-totals"></p>
        </div>
      </div>
    </div>
@@ -57,7 +57,7 @@ PREBUILT_ASSET = "strata-windows-x64.zip" if WIN else "strata-linux-x64.zip"
 # the CUDA libraries the ready-made engine loads (the same CUDA 13.0 it is built with), from NVIDIA's pip packages
 CUDA_WHEELS = ["nvidia-cublas==13.0.2.14", "nvidia-cuda-runtime==13.0.96"]
 MIN_DRIVER = 580                       # CUDA 13.0
-MIN_ENGINE = (0, 1, 13)                # the faster prompt path and `--prefill auto`, v0.1.13
+MIN_ENGINE = (0, 1, 16)                # v0.1.16: the Coder (256 experts, PR #54); v0.1.15: EAGER kernels; v0.1.14: #31
 PY_PACKAGES = ["numpy", "jinja2", "regex", "pyyaml", "tqdm", "requests", "cmake", "ninja", "pillow", "psutil"]

 MODELS = {
@@ -70,6 +70,9 @@ MODELS = {
    "IQ3_S": {"about": "3.5-bit i-quant, the best quality (matches the full model), the slowest; needs a 64 GB PC "
                       "with little else running", "download_gb": 83.6, "ram_gb": 62, "arena_gb": 50.3,
              "families": ("qwen",)},
+    # the Coder release: 256 of the 512 experts kept (the ones code, tools and vision use), IQ2_S-IQ4_XS like IQ3_S
+    "IQ1_M": {"about": "the Coder's only size: half the experts, stored like IQ3_S (3.5 bits)", "download_gb": 58.4,
+              "ram_gb": 32, "arena_gb": 23.4, "families": ("coder",)},
 }
 CONTEXTS = [8192, 32768, 65536, 131072, 262144]
 # The model families: the same architecture, weights in the same three GSQ-RCO sizes, different files.
@@ -87,6 +90,15 @@ FAMILIES = {
              "mmproj_hf": "https://huggingface.co/ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF/resolve/main/",
              "mmproj": "mmproj-Swift-Qwen3.8-Flash-Next-BF16.gguf", "name": "swift-1.5",
              "license": "Swift Open License 1.0: https://huggingface.co/ukisai/Swift-1.5-Qwen3.8-Flash-Next-GSQ-RCO-GGUF"},
+    # ISTA-DASLab's expert-pruned release: half of each layer's experts removed, chosen for code, agentic tool use and
+    # vision; its shard 2 (the n-gram table) and vision encoder are the original's files, shared with it
+    "coder": {"title": "Qwen3.8-Flash-Next Coder", "by": "ISTA-DASLab's coding version",
+              "about": "half the experts (code, tools, images kept): needs ~32 GB of RAM, faster; weaker outside coding",
+              "hf": "https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF/resolve/main/{q}/",
+              "file": "Qwen3.8-Flash-Next-GSQ-RCO-{q}-0000{i}-of-00002.gguf", "tag": "coder-",
+              "mmproj_hf": "https://huggingface.co/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-Coder-GGUF/resolve/main/",
+              "mmproj": "mmproj-Qwen3.8-Flash-Next-BF16.gguf", "name": "qwen3.8-flash-next-coder",
+              "profile": "expert-profile-coder.bin"},
 }
 MMPROJ = "mmproj-Qwen3.8-Flash-Next-BF16.gguf"
 # EXPERIMENTAL, off by default (setup asks): a control vector shipped with the repository, see its README
@@ -249,6 +261,7 @@ def find_nvcc():
            cands += [str(p / "bin" / "nvcc.exe") for p in sorted(base.iterdir(), reverse=True)]
    else:
        cands += [str(p / "bin" / "nvcc") for p in sorted(Path("/usr/local").glob("cuda*"), reverse=True)]
+        cands += [str(p / "bin" / "nvcc") for p in sorted(Path("/opt").glob("cuda*"), reverse=True)]   # Arch (#46)
    best = (None, None)
    for c in dict.fromkeys(cands):                     # every toolkit found; the newest wins
        if c and Path(c).exists():
@@ -550,8 +563,9 @@ def install_build_tools(gpu, yes):
    else:
        apt = shutil.which("apt-get")
        if apt is None:
-            fail("automatic install is only done on Ubuntu/Debian",
-                 "install g++ and the CUDA Toolkit 13 (https://developer.nvidia.com/cuda-downloads), then run it again")
+            fail("missing: " + " and ".join(missing) + " (the automatic install is only done on Ubuntu/Debian)",
+                 "install them with your distribution's packages (Arch: pacman -S base-devel cuda; nvcc is found on "
+                 "PATH, in /usr/local/cuda* and in /opt/cuda*), then run it again")
        if not have_cc:
            run(["sudo", "apt-get", "install", "-y", "build-essential"])
        if nvcc is None or cuda_v < need_cuda:
@@ -582,15 +596,20 @@ def cmake_build(src, bdir, target, defs, vcvars, bat_name):
    conf = [cmake, "-G", "Ninja", f"-DCMAKE_MAKE_PROGRAM={ninja}", "-S", str(src), "-B", str(bdir),
            "-DCMAKE_BUILD_TYPE=Release", *defs]
    build = [cmake, "--build", str(bdir), "--target", target, "-j", str(max(2, (os.cpu_count() or 4) // 2))]
+    # A failed build is tried once more: CUDA 13.0's ptxas now and then fails to parse a PTX file it just wrote, and
+    # the same command then gets past it (issue #45); a second attempt only compiles what is still missing.
    if WIN:
        bat = ROOT / bat_name
        q = lambda c: " ".join(f'"{x}"' if " " in str(x) else str(x) for x in c)  # noqa: E731
-        bat.write_text(f'@echo off\r\ncall "{vcvars}" >nul\r\n{q(conf)} || exit /b 1\r\n{q(build)} || exit /b 1\r\n',
+        bat.write_text(f'@echo off\r\ncall "{vcvars}" >nul\r\n{q(conf)} || exit /b 1\r\n{q(build)} && exit /b 0\r\n'
+                       f'echo   (the build stopped - trying it once more)\r\n{q(build)} || exit /b 1\r\n',
                       encoding="utf-8")
        run(["cmd", "/c", str(bat)])
    else:
        run(conf)
-        run(build)
+        if run(build, check=False).returncode != 0:
+            say("  (the build stopped - trying it once more)")
+            run(build)


 ENGINE_SOURCES = ("CMakeLists.txt", "src", "include", "third_party/ggml")
@@ -640,36 +659,265 @@ def build_engine(gpu, vision, yes, llama) -> Path:
        shutil.copy2(ROOT / "build-vision" / "bin" / VEXE, eng / VEXE)
    bindir = Path(nvcc).parent                            # the toolkit's own libraries (bin, bin/x64, lib64)
    dirs = [str(d) for d in (bindir, bindir / "x64", bindir.parent / "lib64") if d.is_dir()]
-    stamp.write_text(json.dumps({"source": "local", "archs": [int(gpu["arch"])], "vision": vision,
+    stamp.write_text(json.dumps({"source": "local", "version": source_version(), "archs": [int(gpu["arch"])],
+                                 "vision": vision,
                                 "cuda_dirs": dirs, "src": src, "vision_src": vsrc if want_vision else None}, indent=1))
    ok(f"engine compiled: {eng / EXE}")
    return eng


+# ------------------------------------------------------------------------------------------------ the data folder
+# The model files - the GGUFs, the prepared packs and the MTP layer, 70-120 GB - live in a data folder NEXT TO the
+# Strata folder (`Strata-data`), not inside it: updating Strata by unzipping a new copy used to give a new, empty
+# folder and a full download again.  Where it is, and which Strata folders this user ran, is kept in a small
+# per-user file, so every Strata folder on the PC finds the same files.
+DATA_ITEMS = ("models", "packs", "mtp")
+
+
+def settings_path() -> Path:
+    if WIN:
+        return Path(os.environ.get("APPDATA") or Path.home() / "AppData" / "Roaming") / "Strata" / "settings.json"
+    return Path(os.environ.get("XDG_CONFIG_HOME") or Path.home() / ".config") / "strata" / "settings.json"
+
+
+def load_settings() -> dict:
+    try:
+        return json.loads(settings_path().read_text(encoding="utf-8"))
+    except (OSError, ValueError):
+        return {}
+
+
+def save_settings(s: dict) -> None:
+    try:
+        settings_path().parent.mkdir(parents=True, exist_ok=True)
+        settings_path().write_text(json.dumps(s, indent=1), encoding="utf-8")
+    except OSError as e:
+        warn(f"could not save {settings_path()} ({e})")
+
+
+def has_data(folder: Path) -> bool:
+    for d in DATA_ITEMS:
+        try:
+            if (folder / d).is_dir() and any((folder / d).iterdir()):
+                return True
+        except OSError:
+            pass
+    return False
+
+
+def other_installs(settings: dict) -> list:
+    """Strata folders besides this one that may hold model files: the ones this user ran before, and Strata* folders
+    next to this one (a zip unpacked again lands in e.g. `Strata-main (1)\\Strata-main`)."""
+    cands = [Path(p) for p in settings.get("installs", [])]
+    for base in dict.fromkeys((ROOT.parent, ROOT.parent.parent)):
+        try:
+            for d in base.iterdir():
+                if d.is_dir() and d.name.lower().startswith("strata"):
+                    cands.append(d)
+                    cands += [c for c in d.iterdir() if c.is_dir() and c.name.lower().startswith("strata")]
+        except OSError:
+            pass
+    found = []
+    for d in cands:
+        try:
+            d = d.resolve()
+            if d != ROOT and d not in found and (d / "setup.py").is_file():
+                found.append(d)
+        except OSError:
+            pass
+    return found
+
+
+def same_drive(a: Path, b: Path) -> bool:
+    try:
+        return os.stat(a).st_dev == os.stat(b).st_dev
+    except OSError:
+        return False
+
+
+def move_into(src: Path, dst: Path) -> None:
+    """A rename into the data folder (same drive: instant); a folder merges into one already there, keeping what the
+    destination has.  Whatever cannot be moved (a file in use) stays where it is."""
+    if not dst.exists():
+        try:
+            dst.parent.mkdir(parents=True, exist_ok=True)
+            os.replace(src, dst)
+            return
+        except OSError:
+            if not src.is_dir():
+                return
+            dst.mkdir(parents=True, exist_ok=True)
+    if src.is_dir() and dst.is_dir():
+        for c in list(src.iterdir()):
+            move_into(c, dst / c.name)
+        try:
+            src.rmdir()
+        except OSError:
+            pass
+
+
+def repoint_config(cfg_file: Path, old: Path, new: Path) -> None:
+    """A config whose model files moved from `old` to `new` points at them there (each path only if its file is
+    now there and no longer at the old place)."""
+    try:
+        cfg = json.loads(cfg_file.read_text(encoding="utf-8-sig"))
+    except (OSError, ValueError):
+        return
+
+    def fix(v):
+        if isinstance(v, list):
+            return [fix(x) for x in v]
+        if isinstance(v, dict):
+            return {k: fix(x) for k, x in v.items()}
+        if isinstance(v, str):
+            for d in DATA_ITEMS:
+                o = str(old / d)
+                nv, no = os.path.normcase(v), os.path.normcase(o)   # Windows: C:\ and c:\ are the same place
+                if nv == no or nv.startswith(no + os.sep):
+                    n = str(new / d) + v[len(o):]
+                    if Path(n).exists() and not Path(v).exists():
+                        return n
+        return v
+
+    new_cfg = fix(cfg)
+    if new_cfg != cfg:
+        cfg_file.write_text(json.dumps(new_cfg, indent=1), encoding="utf-8")
+
+
+def data_folder(requested: str | None) -> tuple:
+    """(the data folder, folders on other drives that still hold model files).  Moves the model files of this folder
+    and of earlier Strata folders on the same drive into the data folder, and points their configs there."""
+    settings = load_settings()
+    dest = Path(requested).expanduser().resolve() if requested else \
+        Path(settings["data_dir"]) if settings.get("data_dir") else ROOT.parent / "Strata-data"
+    try:
+        dest.mkdir(parents=True, exist_ok=True)
+    except OSError as e:                                # e.g. no write access next to the Strata folder
+        warn(f"cannot use {dest} for the model files ({e}): keeping them in {ROOT}")
+        dest = ROOT
+    elsewhere = []
+    for folder in [ROOT, *other_installs(settings)]:
+        if folder == dest or not has_data(folder):
+            continue
+        if not same_drive(folder, dest):
+            elsewhere.append(folder)                    # another drive: used where it is (no 70 GB copy)
+            continue
+        # the downloads merge file by file (the same file wherever it came from); a prepared pack or MTP layer moves
+        # whole or not at all, so two copies are never mixed
+        if (folder / "models").is_dir():
+            move_into(folder / "models", dest / "models")
+        for item in [*((folder / "packs").glob("*") if (folder / "packs").is_dir() else []), folder / "mtp"]:
+            rel = item.relative_to(folder)
+            if item.exists() and not (dest / rel).exists():
+                move_into(item, dest / rel)
+        for d in ("packs",):
+            try:
+                (folder / d).rmdir()                    # empty now
+            except OSError:
+                pass
+        for c in folder.glob("strata-*.json"):
+            repoint_config(c, folder, dest)
+        if has_data(folder):
+            elsewhere.append(folder)                    # in use, or a copy the data folder already has
+            warn(f"some model files are still in {folder} (in use, or already in {dest})")
+        else:
+            ok(f"model files from {folder} moved to {dest} (a new copy of Strata finds them there)")
+    installs = [str(ROOT)] + [p for p in settings.get("installs", []) if p != str(ROOT) and Path(p).is_dir()]
+    save_settings({**settings, "data_dir": str(dest), "installs": installs[:20]})
+    return dest, elsewhere
+
+
+def previous_config(elsewhere_first: list, settings: dict):
+    """The most recently used model config of another Strata folder on this PC, for a folder that has none yet."""
+    cands = []
+    for folder in [*elsewhere_first, *other_installs(settings)]:
+        cands += list(folder.glob("strata-*.json"))
+    cands = [c for c in dict.fromkeys(cands) if c.is_file()]
+    return max(cands, key=lambda p: p.stat().st_mtime) if cands else None
+
+
+def choices_from_config(cfg_path: Path) -> dict:
+    """The setup answers a config was written with (family, size, context, KV, images, projection, network)."""
+    cfg = json.loads(cfg_path.read_text(encoding="utf-8-sig"))
+    tag = cfg_path.stem[len("strata-"):]
+    family = next((f for f, d in FAMILIES.items() if d["tag"] and tag.startswith(d["tag"])), "qwen")
+    model = tag.split("-")[-1].upper()
+    a = cfg.get("args", [])
+    val = lambda k: a[a.index(k) + 1] if k in a and a.index(k) + 1 < len(a) else None   # noqa: E731
+    vis = cfg.get("vision")
+    esp = val("--control-vector-scaled")
+    esp_path = esp.rsplit(":", 1)[0] if esp else None
+    return {"family": family, "model": model if model in MODELS else None,
+            "context": int(val("--max-context")) if val("--max-context") else None,
+            "kv": val("--kv") if val("--kv") in ("int8", "q4_0") else None,
+            "vision": ("gpu" if vis.get("gpu") else "cpu") if isinstance(vis, dict) else "none",
+            "esp": ("on" if Path(esp_path).name == ESP_VECTOR.name else esp_path) if esp_path else "off",
+            "host": cfg.get("host"), "api_key": cfg.get("api_key"), "port": cfg.get("port")}
+
+
+def find_in(roots: list, rel: str):
+    """The first of roots/rel that exists."""
+    for r in roots:
+        if (r / rel).exists():
+            return r / rel
+    return None
+
+
 # ------------------------------------------------------------------------------------------------ start
 def installed_configs():
    return sorted(ROOT.glob("strata-*.json"), key=lambda p: p.stat().st_mtime, reverse=True)


+def source_version() -> str:
+    """The engine version the source tree builds (CMakeLists.txt's project version)."""
+    m = re.search(r"project\(strata VERSION ([\d.]+)", (ROOT / "CMakeLists.txt").read_text(encoding="utf-8"))
+    return m.group(1) if m else "0"
+
+
 def engine_version(exe: Path) -> tuple:
-    """The version in the engine folder's BUILD.json; a locally compiled engine is the current source's."""
+    """The version in the engine folder's BUILD.json, or else the one compiled into the binary.  A locally compiled
+    engine is not necessarily the source's version: when compiling a `git pull` fails, the previous engine is kept
+    (issue #49)."""
    try:
        meta = json.loads((Path(exe).parent / "BUILD.json").read_text())
    except (OSError, ValueError):
-        return (0, 0, 0)
-    if meta.get("source") == "local":
-        return MIN_ENGINE
-    return tuple(int(x) for x in str(meta.get("version", "0")).split(".")[:3] if x.isdigit())
+        meta = {}
+    v = str(meta.get("version") or "")
+    if not v:                                          # the version compiled into the binary: 0.1.13 and newer
+        try:                                           # carry it, so a binary without it is older
+            m = re.search(rb"engine=(\d+\.\d+\.\d+)\n", Path(exe).read_bytes())
+            v = m.group(1).decode() if m else "0.1.12"
+        except OSError:
+            v = "0"
+    return tuple(int(x) for x in v.split(".")[:3] if x.isdigit())
+
+
+def is_wsl() -> bool:
+    return sys.platform.startswith("linux") and "microsoft" in platform.uname().release.lower()


 def upgrade_config(cfg_path: Path, cfg: dict) -> dict:
    """Configs written before v0.1.13 read prompts in fixed 2048-token chunks; the engine now picks the chunk
-    itself (`--prefill auto`: up to 8192, as the free VRAM allows - about 2x faster on long prompts)."""
+    itself (`--prefill auto`: up to 8192, as the free VRAM allows - about 2x faster on long prompts).  Under WSL,
+    KV streaming is dropped: its RAM copy must be pinned, and the driver pins only about 1 GB there."""
    a = cfg.get("args", [])
-    if "--prefill" in a and a[a.index("--prefill") + 1] == "2048" and engine_version(cfg["exe"]) >= (0, 1, 13):
+    changed = False
+    ver = engine_version(cfg["exe"]) if "--prefill" in a else (0, 0, 0)
+    if "--prefill" in a and a[a.index("--prefill") + 1] == "2048" and ver >= (0, 1, 13):
        a[a.index("--prefill") + 1] = "auto"
-        cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
+        changed = True
        ok("prompt reading: the engine now picks its chunk size (--prefill auto)")
+    elif "--prefill" in a and a[a.index("--prefill") + 1] == "auto" and (0, 0, 0) < ver < (0, 1, 13):
+        a[a.index("--prefill") + 1] = "2048"           # an older engine kept after a failed update (issue #49)
+        changed = True
+        warn(f"the installed engine is {'.'.join(map(str, ver))}: prompts are read in 2048-token chunks until it is updated")
+    if is_wsl() and "--kv-resident" in a:
+        i = a.index("--kv-resident")
+        del a[i:i + 2]
+        changed = True
+        ok("WSL: KV streaming off (the driver pins only about 1 GB of RAM); the KV cache stays in VRAM")
+    if changed:
+        cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
    return cfg


@@ -683,8 +931,20 @@ def start(cfg_path: Path, port: int | None, open_browser=True) -> int:
           "--port", str(port or cfg.get("port", 8080))]
    if open_browser:
        cmd.append("--open")
+    gb = 0.0
+    if "--native" in cfg["args"]:
+        try:
+            gb = Path(cfg["args"][cfg["args"].index("--native") + 1]).stat().st_size / 1e9
+        except (OSError, IndexError):
+            pass
    say()
-    say(f"Starting {cfg.get('model_name', 'the model')} (loads 34-43 GB into RAM: 30-90 s). Close this window to stop it.")
+    say("  " + "-" * 100)
+    say(f"  Starting {cfg.get('model_name', 'the model')}: it loads {f'about {gb:.0f} GB' if gb >= 1 else '34-55 GB'} "
+        "into RAM and locks part of it for the GPU.")
+    say("  While it does, YOUR PC CAN BE SLOW OR STOP RESPONDING FOR 1-3 MINUTES (longer the first time after a")
+    say("  restart). That is normal: please wait and don't close this window - the browser opens when it is ready.")
+    say("  Later, closing this window stops the model.")
+    say("  " + "-" * 100)
    return subprocess.call(cmd)


@@ -717,11 +977,13 @@ def main() -> int:
    ap.add_argument("--experimental-speed-projection", metavar="on|off|GGUF",
                    help="EXPERIMENTAL, off by default: the control vector in data/experimental-speed-projection "
                         "(or another GGUF) as a projection on layers 4-44; see docs/DETAILS.md")
-    ap.add_argument("--port", type=int, default=8080)
+    ap.add_argument("--port", type=int, help="the server's port (default: the one the install was set up with, 8080 for a new one)")
    ap.add_argument("--host", help="where the server listens: 127.0.0.1 = this PC only (default), 0.0.0.0 = also other "
                                   "devices on your network (issue #26; set --api-key too)")
    ap.add_argument("--api-key", help="require this key from clients (recommended with --host 0.0.0.0)")
-    ap.add_argument("--models-dir", default=str(ROOT / "models"), help="where the GGUF files go (~70 GB)")
+    ap.add_argument("--data-dir", help="where the model files go (~70-120 GB): default Strata-data next to this folder, "
+                                       "remembered for every Strata folder on this PC")
+    ap.add_argument("--models-dir", help="where the GGUF files go (default: <data folder>/models)")
    ap.add_argument("--gguf-dir", help="use GGUF files you already have (a folder with the two shards)")
    ap.add_argument("--yes", action="store_true", help="accept the recommended answers")
    ap.add_argument("--setup", action="store_true", help="install another model or change settings")
@@ -732,23 +994,42 @@ def main() -> int:
    ap.add_argument("--check", action="store_true", help="only check this PC and exit")
    ap.add_argument("--skip-build", action="store_true", help=argparse.SUPPRESS)
    a = ap.parse_args()
-
    say("Strata - Qwen3.8-Flash-Next on a normal PC (NVIDIA GPU + system RAM + CPU)")
+    data, elsewhere = data_folder(a.data_dir)          # the model files: in the data folder, found from any copy
+    roots = [data, *elsewhere]
+    if a.models_dir is None:
+        a.models_dir = str(data / "models")

    # ---- 0. already installed: just start it
    have = installed_configs()
+    explicit = a.setup or a.model or a.family or a.check or a.no_start
+    if not have and not explicit:                      # a new copy of Strata (an update unzipped elsewhere): set it
+        prev = previous_config(elsewhere, load_settings())   # up like the last one, from the files already here
+        if prev is not None:
+            ch = choices_from_config(prev)
+            if ch["model"]:
+                say(f"  Found your earlier install in {prev.parent} ({prev.stem[len('strata-'):]}): setting up this "
+                    "copy the same way - the model files are reused, nothing big is downloaded.")
+                a.family, a.model, a.context = ch["family"], ch["model"], a.context or ch["context"]
+                a.kv = a.kv or ch["kv"]
+                a.vision = a.vision or ch["vision"]
+                a.experimental_speed_projection = a.experimental_speed_projection or ch["esp"]
+                a.host, a.api_key = a.host or ch["host"], a.api_key or ch["api_key"]
+                a.port = a.port or ch["port"]
+                a.yes = True
+    port = a.port or 8080                              # a new install's port (issue #32: --port for an existing one)
    if have and not (a.setup or a.model or a.family or a.check or a.no_start):
        if not a.build:
            update_installed_engine(a.prebuilt)
        if len(have) == 1:
-            return start(have[0], None)
+            return start(have[0], a.port)
        say()
        for i, c in enumerate(have, 1):
            say(f"  {i}) {json.loads(c.read_text(encoding='utf-8-sig')).get('model_name', c.stem)}")
        say(f"  {len(have) + 1}) install another model / change settings")
        pick = int(ask("Which one?", [str(i) for i in range(1, len(have) + 2)], "1", a.yes))
        if pick <= len(have):
-            return start(have[pick - 1], None)
+            return start(have[pick - 1], a.port)

    # ---- 1. the PC
    step(1, "checking your PC")
@@ -769,11 +1050,11 @@ def main() -> int:
    cpu, avx2, avx512 = cpu_info()
    need = min(d["ram_gb"] for d in MODELS.values())
    if ram < need - 4 and not a.check:
-        # every model keeps ALL its experts in RAM (34+ GB); VRAM only holds a copy of the most-used ones, so a
+        # every model keeps ALL its experts in RAM (23+ GB); VRAM only holds a copy of the most-used ones, so a
        # bigger GPU does not lower this
-        fail(f"RAM: {ram:.0f} GB - the smallest model (Q2_0 / IQ2_XS) needs about {need} GB",
-             "Strata keeps all of the model's experts in RAM (34-50 GB, whatever the GPU) and the GPU holds a copy "
-             "of the most-used ones: it needs 48 GB of RAM or more")
+        fail(f"RAM: {ram:.0f} GB - the smallest model (the Coder) needs about {need} GB",
+             "Strata keeps all of the model's experts in RAM (23-50 GB, whatever the GPU) and the GPU holds a copy "
+             "of the most-used ones: it needs 32 GB of RAM or more (48 GB for the full model)")
    ok(f"RAM: {ram:.0f} GB" if ram >= need - 4 else f"RAM: {ram:.0f} GB (less than the {need} GB the smallest model needs)")
    ok(f"CPU: {cpu} ({'AVX-512' if avx512 else 'AVX2' if avx2 else 'no AVX2'})")
    if not avx2:
@@ -801,14 +1082,14 @@ def main() -> int:
    if fam.get("license"):
        say(f"  Its license: {fam['license']}")
    say()
-    names = [m for m in MODELS if family in MODELS[m].get("families", FAMILIES)]
+    names = [m for m in MODELS if family in MODELS[m].get("families", ("qwen", "swift"))]
    if a.model and a.model not in names:
        fail(f"{fam['title']} has no {a.model} model file", "choose one of: " + ", ".join(names))
    for i, m in enumerate(names, 1):
        d = MODELS[m]
        fit = "" if ram >= d["ram_gb"] else f"   <- needs {d['ram_gb']} GB RAM, you have {ram:.0f}"
        say(f"  {i}) {m:8s} {d['about']}; download {d['download_gb']:.0f} GB, uses ~{d['arena_gb']:.0f} GB of RAM{fit}")
-    rec = str(names.index("IQ3_XXS") + 1) if ram >= 60 else "1"
+    rec = str(names.index("IQ3_XXS") + 1) if ram >= 60 and "IQ3_XXS" in names else "1"
    model = a.model or names[int(ask("Which size?", [str(i) for i in range(1, len(names) + 1)], rec, a.yes)) - 1]
    if ram < MODELS[model]["ram_gb"] - 4:
        fail(f"{model} needs about {MODELS[model]['ram_gb']} GB of RAM; this PC has {ram:.0f} GB",
@@ -853,7 +1134,7 @@ def main() -> int:
    # chosen here; with it loaded, the web app and the API switch it off per request
    esp = None
    esp_choice = (a.experimental_speed_projection or "").strip()
-    if family == "qwen":
+    if family in ("qwen", "coder"):                   # the Coder: the same model's residual stream
        if not esp_choice:
            say()
            say("  EXPERIMENTAL - speed projection: a small control vector applied while the model runs (layers 4-44).")
@@ -870,6 +1151,13 @@ def main() -> int:
        warn("the experimental speed projection is made for the original Qwen3.8-Flash-Next, not Swift 1.5: left off")
    models_dir = Path(a.gguf_dir) if a.gguf_dir else Path(a.models_dir) / tag
    shards = [models_dir / fam["file"].format(q=model, i=i) for i in (1, 2)]
+    if not a.gguf_dir and not all(sh.exists() and done(sh) for sh in shards):
+        for r in elsewhere:                            # already downloaded in a Strata folder on another drive
+            cand = [r / "models" / tag / sh.name for sh in shards]
+            if all(c.exists() and done(c) for c in cand):
+                models_dir, shards = cand[0].parent, cand
+                ok(f"model files found in {models_dir}")
+                break
    have_model = all(s.exists() and (done(s) or a.gguf_dir) for s in shards)
    need = (0 if a.gguf_dir or have_model else MODELS[model]["download_gb"]) + 8 + \
        (40 if model == "Q2_0" and avx512 and family == "qwen" else 0) + (1 if vision != "none" else 0)
@@ -903,9 +1191,9 @@ def main() -> int:
            if s.exists() and done(s):
                ok(f"{s.name} already downloaded")
                continue
-            # the original's shard 2 is the same file for all three sizes: reuse one that is already here
+            # the original's shard 2 is the same file for all its sizes and the Coder: reuse one that is already here
            other = [p for p in Path(a.models_dir).glob("*/Qwen3.8-Flash-Next-GSQ-RCO-*-00002-of-00002.gguf") if done(p)]
-            if family == "qwen" and s.name.endswith("00002-of-00002.gguf") and other and not s.exists():
+            if family in ("qwen", "coder") and s.name.endswith("00002-of-00002.gguf") and other and not s.exists():
                try:
                    os.link(other[0], s)
                    mark(s)
@@ -919,6 +1207,8 @@ def main() -> int:
            fail(f"missing {s}")
    ok("model files present")
    mmproj = Path(a.models_dir) / fam["mmproj"]
+    if not mmproj.exists():
+        mmproj = find_in(roots, f"models/{fam['mmproj']}") or mmproj
    if vision != "none":
        if not mmproj.exists() and a.gguf_dir and (Path(a.gguf_dir) / fam["mmproj"]).exists():
            mmproj = Path(a.gguf_dir) / fam["mmproj"]
@@ -928,7 +1218,7 @@ def main() -> int:

    # ---- 6. the pack and the MTP draft layer
    step(6, "preparing the model for Strata")
-    pack = ROOT / "packs" / tag.lower()
+    pack = find_in(roots, f"packs/{tag.lower()}") or data / "packs" / tag.lower()
    env = dict(os.environ, STRATA_GGUF_PY=str(llama / "gguf-py"))
    if model == "Q2_0" and avx512 and family == "qwen":
        # the Q2_0 experts repacked for the AVX-512 kernel (the measured speed): a one-time ~40 GB conversion
@@ -944,7 +1234,7 @@ def main() -> int:
        # every tensor as the GGUF stores it; the experts are read from the GGUF at start (seconds to build)
        run([sys.executable, str(ROOT / "tools" / "iq_pack.py"), "--gguf", str(shards[0]), "--out", str(pack)], env=env)
    ok(f"model prepared: {pack}")
-    mtp = ROOT / "mtp"
+    mtp = (find_in(roots, "mtp/rt/experts.bin") or data / "mtp/rt/experts.bin").parent.parent
    rt = mtp / "rt"
    if not (rt / "experts.bin").exists():
        say("  The MTP draft layer (speculative decoding, ~2x faster output) comes from the original Qwen checkpoint:")
@@ -966,7 +1256,7 @@ def main() -> int:
    if ple is None:
        fail("the model has no per_layer_token_embd tensor (is this a Qwen3.8-Flash-Next GGUF?)")
    args = ["--pack", str(pack), "--native", str(shards[0]), "--ple-gguf", str(ple),
-            "--expert-profile", str(ROOT / "data" / "expert-profile.bin"), "--expert-cache", "auto",
+            "--expert-profile", str(ROOT / "data" / fam.get("profile", "expert-profile.bin")), "--expert-cache", "auto",
            "--prefill", "auto", "--spec", "4", "--spec-min-p", "0.5", "--mtp", str(rt),
            "--max-context", str(ctx)]
    if ctx > 8192:
@@ -975,7 +1265,10 @@ def main() -> int:
    # per layer) stays in VRAM; the VRAM it frees holds more experts (+6% at 128K, +23% at 262K with Q2_0). It
    # costs ~13.7 KB of RAM per context token with 8-bit KV (1.7 GB at 128K), 7.5 KB with 4-bit, so only when it fits.
    kv_ram_gb = ctx * (13 * (576 if kv == "q4_0" else 1056)) / 1e9   # 12 QSA layers + the draft layer
-    if ctx >= 65536 and ram >= MODELS[model]["ram_gb"] + kv_ram_gb + 1:
+    # the RAM copy must be pinned, and under WSL the NVIDIA driver pins only about 1 GB in all
+    if is_wsl() and ctx >= 65536:
+        ok("WSL: KV streaming off (the driver pins only about 1 GB of RAM); the KV cache stays in VRAM")
+    elif ctx >= 65536 and ram >= MODELS[model]["ram_gb"] + kv_ram_gb + 1:
        args += ["--kv-resident", "32768"]
        ok(f"KV streaming on: the context's KV cache lives in RAM ({kv_ram_gb:.1f} GB), more experts fit in VRAM")
    if vision != "none":
@@ -986,7 +1279,7 @@ def main() -> int:
                 "--cvec-mode", "project", "--cvec-dir", "per-layer"]
    cfg = {"exe": str(eng / EXE), "args": args, "cwd": str(ROOT), "tokenizer": str(pack / "tokenizer"),
           "model_name": f"{fam['name']}-{model.lower()}", "log": str(ROOT / f"strata-{tag.lower()}.log"),
-           "lib_dirs": lib_dirs, "port": a.port}
+           "lib_dirs": lib_dirs, "port": port}
    if a.host:
        cfg["host"] = a.host
    if a.api_key:
@@ -998,22 +1291,22 @@ def main() -> int:
            cfg["vision"]["threads"] = max(1, (os.cpu_count() or 8) // 2)
    cfg_path = ROOT / f"strata-{tag.lower()}.json"
    cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
-    script = write_run_script(tag, cfg_path, a.port)
+    script = write_run_script(tag, cfg_path, port)
    ok(f"start script: {script.name}")

    say()
    say("All set.")
-    say(f"  API (OpenAI):     http://127.0.0.1:{a.port}/v1   (any API key; model name: anything)")
-    say(f"  API (Anthropic):  http://127.0.0.1:{a.port}/v1/messages")
+    say(f"  API (OpenAI):     http://127.0.0.1:{port}/v1   (any API key; model name: anything)")
+    say(f"  API (Anthropic):  http://127.0.0.1:{port}/v1/messages")
    if a.host and a.host not in ("127.0.0.1", "localhost"):
-        say(f"  Other devices:    the server window prints this PC's address (http://<IP>:{a.port}/)"
+        say(f"  Other devices:    the server window prints this PC's address (http://<IP>:{port}/)"
            + ("" if a.api_key else " - no API key set: anyone on your network can use it"))
    say(f"  Next time:        just run {'START-HERE.bat' if WIN else './setup.sh'} (or {script.name}) - it starts right away")
    if vision != "none":
        say("  Images:           send them in the chat page, in chat.py (/image <path>) or over the API")
    if a.no_start:
        return 0
-    return start(cfg_path, a.port)
+    return start(cfg_path, port)


 if __name__ == "__main__":
@@ -2,10 +2,17 @@
 # Strata for Linux: the first run installs everything and starts the model; later runs just start it.
 # Needs only an NVIDIA driver. Python (with venv) is installed through apt/dnf if it is missing (asks for sudo).
 cd "$(dirname "$0")" || exit 1
+# Python 3.10+ that can make a venv WITH pip: Debian/Ubuntu ship `venv` without `ensurepip` (that is the separate
+# python3-venv package), and a venv made without it has no pip
+ok_py() { "$1" -c 'import sys, venv, ensurepip; sys.exit(0 if sys.version_info >= (3, 10) else 1)' 2>/dev/null; }
+# a .venv from an earlier run that failed half-way has a python but no pip: start it again
+if [ -x .venv/bin/python ] && ! .venv/bin/python -m pip --version >/dev/null 2>&1; then
+  rm -rf .venv
+fi
 if [ ! -x .venv/bin/python ]; then
  PY=""
  for c in python3 python; do
-    if command -v $c >/dev/null 2>&1 && $c -c 'import sys, venv; sys.exit(0 if sys.version_info >= (3, 10) else 1)' 2>/dev/null; then
+    if command -v $c >/dev/null 2>&1 && ok_py $c; then
      PY=$c; break
    fi
  done
@@ -19,12 +26,13 @@ if [ ! -x .venv/bin/python ]; then
      sudo pacman -S --noconfirm python python-pip
    fi
    PY=python3
-    if ! $PY -c 'import sys, venv; sys.exit(0 if sys.version_info >= (3, 10) else 1)' 2>/dev/null; then
-      echo "Please install Python 3.10 or newer (with venv), then run ./setup.sh again."
+    if ! ok_py $PY; then
+      echo "Please install Python 3.10 or newer with venv (Ubuntu/Debian: sudo apt install python3-venv), then run"
+      echo "./setup.sh again."
      exit 1
    fi
  fi
  # a private environment inside this folder (system Python stays untouched; newer distros refuse global pip)
-  $PY -m venv .venv || { echo "could not create .venv: sudo apt install python3-venv"; exit 1; }
+  $PY -m venv .venv || { rm -rf .venv; echo "could not create .venv: sudo apt install python3-venv"; exit 1; }
 fi
 exec .venv/bin/python setup.py "$@"
@@ -362,11 +362,8 @@ if (!native_bf16_projections) f32_to_bf16_bulk(x, b.x_bf16, g.n_embd, stream);
 //      LEDGER L41 -> L42.
 project_bf16(x, b.x_bf16, (const uint16_t*) w_router->data, b.logits, g.n_embd, g.n_expert, true, stream);
 // ---- routing: softmax over ALL experts, stable descending argsort with ties by index, gather, renormalise
-if (native_router_enabled()) {
-    if (g.n_expert != 512 || k != 10) {
-        err = v.name("router") + ": native router requires 512 experts and k=10";
-        return false;
-    }
+// the native fused router is canonical-512x10 only; anything else takes the generic top-k kernel
+if (native_router_enabled() && g.n_expert == 512 && k == 10) {
    try { native_router_top10(b.logits, b.ids, b.weights, stream); }
    catch (const std::exception& error) { err = v.name("router") + ": " + error.what(); return false; }
 } else router_top10(b.logits, 1, (int) g.n_expert, (int) k, b.ids, b.weights, stream);
@@ -610,8 +607,11 @@ uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, Q
        uint8_t* d = nullptr;
        if (cudaHostAlloc((void**) &h, bytes, cudaHostAllocMapped | cudaHostAllocPortable) != cudaSuccess ||
            cudaHostGetDevicePointer((void**) &d, h, 0) != cudaSuccess) {
+            // under WSL the NVIDIA driver pins only ~1 GiB in all, which is less than 128K of 8-bit KV needs
            if (p.mode == 1) std::fprintf(stderr, "strata: KV streaming: cannot pin %.2f GiB of RAM for a layer's KV copy "
-                                 "(lower the context, or run without --kv-resident)\n", (double) bytes / 1073741824.0);
+                                 "(%.2f GiB pinned so far) - lower the context, or run without --kv-resident (under "
+                                 "WSL the driver pins only about 1 GiB in all)\n", (double) bytes / 1073741824.0,
+                                 (double) g_kv_host_bytes / 1073741824.0);
            return 0;
        }
        g_kv_host_bytes += bytes;
@@ -29,20 +29,46 @@ namespace {
 // A 2 MB-aligned reservation.  Large pages first, then the largest alignment the OS will give us for free.
 void* reserve(uint64_t bytes, PageBacking& got, std::string& note) {
 #ifdef _WIN32
+    // MEM_LARGE_PAGES needs SeLockMemoryPrivilege.  Having it assigned to the account is not enough: the
+    // PROCESS must enable it in its own token (AdjustTokenPrivileges) before VirtualAlloc, or the call fails.
+    // An account without the assignment, or a failure to enable, leaves the process as it was: VirtualAlloc
+    // then refuses and the 4 KB fallback below runs - that is the EXPECTED outcome on a desktop.
+    {
+        HANDLE tok = nullptr;
+        if (OpenProcessToken(GetCurrentProcess(), TOKEN_ADJUST_PRIVILEGES | TOKEN_QUERY, &tok)) {
+            TOKEN_PRIVILEGES tp{};
+            tp.PrivilegeCount = 1;
+            if (LookupPrivilegeValueW(nullptr, L"SeLockMemoryPrivilege", &tp.Privileges[0].Luid)) {
+                tp.Privileges[0].Attributes = SE_PRIVILEGE_ENABLED;
+                if (!AdjustTokenPrivileges(tok, FALSE, &tp, 0, nullptr, nullptr) && GetLastError() != ERROR_NOT_ALL_ASSIGNED)
+                    (void) 0;   // nothing actionable: the large-page attempt below reports the outcome
+            }
+            CloseHandle(tok);
+        }
+    }
    // MEM_LARGE_PAGES needs SeLockMemoryPrivilege; a normal account does not have it and VirtualAlloc then
    // fails with ERROR_PRIVILEGE_NOT_HELD.  That is the EXPECTED outcome on a desktop, not an error.
    SIZE_T large = GetLargePageMinimum();
-    if (large > 0) {
-        void* p = VirtualAlloc(nullptr, (SIZE_T) bytes, MEM_RESERVE | MEM_COMMIT | MEM_LARGE_PAGES,
+    // A/B switch: STRATA_NO_LARGEPAGES=1 skips the large-page attempt, same run, same boot.
+    if (large > 0 && std::getenv("STRATA_NO_LARGEPAGES") == nullptr) {
+        // MEM_LARGE_PAGES requires the allocation size to be an exact multiple of the large page size -
+        // anything else is ERROR_INVALID_PARAMETER (87), which reads like a privilege problem but is not.
+        // Round up: the slack is under 2 MB and the tail stays unused.
+        const SIZE_T lbytes = (SIZE_T) (((SIZE_T) bytes + large - 1) / large * large);
+        void* p = VirtualAlloc(nullptr, lbytes, MEM_RESERVE | MEM_COMMIT | MEM_LARGE_PAGES,
                               PAGE_READWRITE);
        if (p) {
            got = PageBacking::LargePages;
            note = "large pages (" + std::to_string((unsigned long long) large) + " B)";
            return p;
        }
-        note = "large pages refused (GetLargePageMinimum=" + std::to_string((unsigned long long) large) +
-               ", VirtualAlloc error " + std::to_string((unsigned long long) GetLastError()) +
-               " - needs SeLockMemoryPrivilege); using 4 KB pages";
+        // 1450 (ERROR_NO_SYSTEM_RESOURCES) is the large-page pool saying no, 87 is a size that is not a
+        // multiple of the minimum, 1314 is the privilege: without the byte count the three read as one bug.
+        note = "large pages refused for " + std::to_string((unsigned long long) lbytes) + " B (GetLargePageMinimum=" +
+               std::to_string((unsigned long long) large) + ", VirtualAlloc error " +
+               std::to_string((unsigned long long) GetLastError()) + "); using 4 KB pages";
+    } else if (std::getenv("STRATA_NO_LARGEPAGES") != nullptr) {
+        note = "large pages skipped (STRATA_NO_LARGEPAGES); using 4 KB pages";
    } else {
        note = "this system has no large-page minimum; using 4 KB pages";
    }
@@ -88,9 +88,12 @@ uint64_t session_init(const ModelGeometry& g, int64_t max_cells, int64_t k, void
    s.qsa_buf_arena = take(qsa_buffers_bytes(g, max_cells));

    uint8_t* qp = (uint8_t*) s.qsa_state_arena;
+    // a layer whose pinned RAM could not be had (KV streaming's host copy) is half-built: going on would have the
+    // attention read null host pointers at the first request ("illegal memory access"), so the session fails here
    for (int64_t i = 0; i < g.n_qsa_layers(); ++i)
-        qsa_state_init(g, max_cells, qp + (i == 0 ? 0 : first + (uint64_t) (i - 1) * rest), s.qsa_states[i],
-                       i == 0 ? nullptr : &s.qsa_states[0]);
+        if (qsa_state_init(g, max_cells, qp + (i == 0 ? 0 : first + (uint64_t) (i - 1) * rest), s.qsa_states[i],
+                           i == 0 ? nullptr : &s.qsa_states[0]) == 0)
+            return 0;
    qsa_buffers_init(g, max_cells, s.qsa_buf_arena, s.qsa_bufs);

    s.moe_arena = take(moe_buffers_bytes(g, k));
@@ -34,6 +34,7 @@

 #include <atomic>
 #include <chrono>
+#include <cmath>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
@@ -93,7 +93,16 @@ bool expert_layout_load(const std::string& pack_dir, int64_t n_layers, int64_t n
    L.max_blob = 0;
    std::string line;
    while (std::getline(in, line)) {
-        if (line.empty() || line[0] == '#') continue;
+        if (line.empty() || line[0] == '#') {
+            if (!line.empty() && line[0] == '#') {
+                // v3 packs record their expert count in the header; a pruned model (GSQ-RCO Coder) ships
+                // fewer experts than the canonical geometry the caller passes, which is a compile-time
+                // default, so the header wins.
+                const size_t at = line.find("(n_expert ");
+                if (at != std::string::npos) L.n_expert = std::atoll(line.c_str() + at + 10);
+            }
+            continue;
+        }
        std::istringstream ss(line);
        long long l = -1, gt = -1, dt = -1;
        unsigned long long off = 0, blob = 0, go = 0, uo = 0, dox = 0;
@@ -130,7 +139,7 @@ bool expert_layout_load(const std::string& pack_dir, int64_t n_layers, int64_t n
            err = "native_experts.txt: layer " + std::to_string(l) + " is missing or not contiguous";
            return false;
        }
-        at += L.bytes[(size_t) l] * (uint64_t) n_expert;
+        at += L.bytes[(size_t) l] * (uint64_t) L.n_expert;
    }
    L.total = at;
    g_layout = L;
@@ -0,0 +1,419 @@
+// src/kernels/cpu/iq_avx2.cpp - the i-quant expert rows in 256-bit lanes, several tokens at once.
+//
+// The AVX-512 kernels' (iq_avx512.cpp) multi-token scheme on the CPUs without AVX-512 (AMD Zen 2/3, Intel
+// Core 12th-14th gen and Core Ultra): each 32-value chunk is decoded once - grid lookups, one sign vector,
+// one scale vector - and every token of the verify window applies them with a load, a `sign`, a `maddubs`,
+// a `madd` and an `add`.  ggml-cpu's own AVX2 dot products for these formats are single-token: every token
+// re-does the codebook lookups and the sign expansion.  The sign vector is ggml's bit_selector pattern
+// (ggml-cpu/arch/x86/quants.c): pshufb-broadcast of each sign byte, AND with the bit selector, CMPEQ, OR
+// with one, `vpsignb`.  The arithmetic is ggml's (ggml-cpu/quants.c, the `_generic` references) - only the
+// order of the float additions differs.
+//
+// Formats: IQ2_XXS (16), IQ2_XS (17), IQ3_XXS (18), IQ3_S (21), IQ2_S (22).  IQ1_M stays on ggml-cpu.
+#include "strata/kernels/cpu/iq_avx2.hpp"
+
+#define GGML_COMMON_DECL_CPP
+#define GGML_COMMON_IMPL_CPP
+#include "ggml-common.h"
+
+#include <immintrin.h>
+
+#include <cmath>
+#include <cstring>
+
+namespace strata::kernels::cpu {
+namespace {
+
+inline float h2f(uint16_t h) { return _mm_cvtss_f32(_mm_cvtph_ps(_mm_cvtsi32_si128((int) h))); }
+inline uint32_t u32(const uint8_t* p) { uint32_t v; std::memcpy(&v, p, 4); return v; }
+inline uint16_t u16(const uint8_t* p) { uint16_t v; std::memcpy(&v, p, 2); return v; }
+inline uint64_t u64(const uint8_t* p) { uint64_t v; std::memcpy(&v, p, 8); return v; }
+
+// 32 sign bits -> 32 bytes of -1 (bit set) / +1: ggml's bit_selector pattern, shared by all tokens.
+inline __m256i sgn_vec(uint32_t m) {
+    const __m128i bm = _mm_set1_epi32((int) m);
+    const __m128i lo = _mm_shuffle_epi8(bm, _mm_setr_epi8(0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1));
+    const __m128i hi = _mm_shuffle_epi8(bm, _mm_setr_epi8(2, 2, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3, 3, 3, 3, 3));
+    const __m256i bits = _mm256_inserti128_si256(_mm256_castsi128_si256(lo), hi, 1);
+    const __m256i sel = _mm256_setr_epi8(1, 2, 4, 8, 16, 32, 64, (char) 0x80,
+                                         1, 2, 4, 8, 16, 32, 64, (char) 0x80,
+                                         1, 2, 4, 8, 16, 32, 64, (char) 0x80,
+                                         1, 2, 4, 8, 16, 32, 64, (char) 0x80);
+    const __m256i nz = _mm256_cmpeq_epi8(_mm256_and_si256(bits, sel), sel);
+    return _mm256_or_si256(nz, _mm256_set1_epi8(1));
+}
+
+// The two 16-value scales of one 32-value half (IQ2_XS, IQ2_S): int16 lanes 0-7 = a (values 0-15),
+// 8-15 = b (values 16-31).  NOT a broadcast dword [a,b] - that would alternate the scales lane by lane.
+inline __m256i sc16(int a, int b) {
+    return _mm256_inserti128_si256(_mm256_castsi128_si256(_mm_set1_epi16((short) a)),
+                                   _mm_set1_epi16((short) b), 1);
+}
+// One scale for the whole 32-value half (IQ2_XXS, IQ3_XXS, IQ3_S).
+inline __m256i sc32(int s) { return _mm256_set1_epi16(s); }
+
+// keven_signs_q2xs (ggml keeps it static in arch/x86/quants.c): one u64 per 7-bit sign index, byte k
+// = 0xFF when bit k of ksigns_iq2xs[i] is set, 0x01 otherwise.  With this a whole 32-value sign vector
+// is four scalar loads and a set_epi64x - no ksigns byte packing chain and no bit_selector expansion.
+// Built once at load time from the shared ggml-common table so it can never drift from it.
+struct EvenSigns {
+    uint64_t v[128];
+    EvenSigns() {
+        for (int i = 0; i < 128; ++i) {
+            uint64_t r = 0;
+            for (int k = 0; k < 8; ++k)
+                r |= (uint64_t) (((ksigns_iq2xs[i] >> k) & 1) ? 0xFF : 0x01) << (8 * k);
+            v[i] = r;
+        }
+    }
+};
+static const EvenSigns even_signs;
+
+inline float hsum8(__m256 v) {
+    const __m128 lo = _mm256_castps256_ps128(v), hi = _mm256_extractf128_ps(v, 1);
+    __m128 s = _mm_add_ps(lo, hi);
+    s = _mm_hadd_ps(s, s);
+    s = _mm_hadd_ps(s, s);
+    return _mm_cvtss_f32(s);
+}
+
+// ---- per format: one 32-value half (values 64*j + 32*half .. +31) -> grid magnitudes, sign vector, scales
+template <int TY> struct Fmt32;
+
+template <> struct Fmt32<16> {   // IQ2_XXS: d, qs[32] u16
+    static constexpr int bytes = 66;
+    static constexpr float K = 0.125f;
+    static inline void decode(const uint8_t* b, int j, int half, __m256i& g, __m256i& sgn, __m256i& sc) {
+        const uint8_t* q = b + 2 + 16 * j + 8 * half;
+        const uint32_t w0 = u32(q), w1 = u32(q + 4);
+        g = _mm256_set_epi64x((long long) iq2xxs_grid[w0 >> 24], (long long) iq2xxs_grid[(w0 >> 16) & 255],
+                              (long long) iq2xxs_grid[(w0 >> 8) & 255], (long long) iq2xxs_grid[w0 & 255]);
+        sgn = _mm256_set_epi64x((long long) even_signs.v[(w1 >> 21) & 127], (long long) even_signs.v[(w1 >> 14) & 127],
+                                (long long) even_signs.v[(w1 >> 7) & 127], (long long) even_signs.v[w1 & 127]);
+        sc = sc32(2 * (int) (w1 >> 28) + 1);
+    }
+};
+
+template <> struct Fmt32<17> {   // IQ2_XS: d, qs[32] u16 (9-bit grid index + 7-bit sign index), scales[8]
+    static constexpr int bytes = 74;
+    static constexpr float K = 0.125f;
+    static inline void decode(const uint8_t* b, int j, int half, __m256i& g, __m256i& sgn, __m256i& sc) {
+        uint16_t v[8];
+        std::memcpy(v, b + 2 + 16 * j, 16);
+        const int o = 4 * half;
+        g = _mm256_set_epi64x((long long) iq2xs_grid[v[o + 3] & 511], (long long) iq2xs_grid[v[o + 2] & 511],
+                              (long long) iq2xs_grid[v[o + 1] & 511], (long long) iq2xs_grid[v[o] & 511]);
+        uint32_t s = 0;
+        for (int l = 0; l < 4; ++l) s |= (uint32_t) ksigns_iq2xs[v[o + l] >> 9] << (8 * l);
+        sgn = sgn_vec(s);
+        const uint8_t sb = b[66 + 2 * j + half];
+        sc = sc16(2 * (sb & 15) + 1, 2 * (sb >> 4) + 1);
+    }
+};
+
+template <> struct Fmt32<22> {   // IQ2_S: d, qs[64] (32 grid bytes, 32 sign bytes), qh[8], scales[8]
+    static constexpr int bytes = 82;
+    static constexpr float K = 0.125f;
+    static inline void decode(const uint8_t* b, int j, int half, __m256i& g, __m256i& sgn, __m256i& sc) {
+        const uint8_t* qs = b + 2 + 8 * j;
+        const uint8_t h = b[66 + 2 * j + half];
+        const int o = 4 * half;
+        g = _mm256_set_epi64x((long long) iq2s_grid[qs[o + 3] | ((h << 2) & 0x300)],
+                              (long long) iq2s_grid[qs[o + 2] | ((h << 4) & 0x300)],
+                              (long long) iq2s_grid[qs[o + 1] | ((h << 6) & 0x300)],
+                              (long long) iq2s_grid[qs[o] | ((h << 8) & 0x300)]);
+        const uint64_t m = u64(b + 2 + 32 + 8 * j);
+        sgn = sgn_vec(half ? (uint32_t) (m >> 32) : (uint32_t) m);
+        const uint8_t sb = b[74 + 2 * j + half];
+        sc = sc16(2 * (sb & 15) + 1, 2 * (sb >> 4) + 1);
+    }
+};
+
+template <> struct Fmt32<18> {   // IQ3_XXS: d, qs[64] grid bytes, 8 x u32 (4 x 7-bit sign index + 4-bit scale)
+    static constexpr int bytes = 98;
+    static constexpr float K = 0.25f;
+    static inline void decode(const uint8_t* b, int j, int half, __m256i& g, __m256i& sgn, __m256i& sc) {
+        const uint8_t* q = b + 2 + 16 * j + 8 * half;
+        g = _mm256_set_epi32((int) iq3xxs_grid[q[7]], (int) iq3xxs_grid[q[6]], (int) iq3xxs_grid[q[5]], (int) iq3xxs_grid[q[4]],
+                             (int) iq3xxs_grid[q[3]], (int) iq3xxs_grid[q[2]], (int) iq3xxs_grid[q[1]], (int) iq3xxs_grid[q[0]]);
+        const uint32_t w = u32(b + 2 + 64 + 8 * j + 4 * half);
+        sgn = _mm256_set_epi64x((long long) even_signs.v[(w >> 21) & 127], (long long) even_signs.v[(w >> 14) & 127],
+                                (long long) even_signs.v[(w >> 7) & 127], (long long) even_signs.v[w & 127]);
+        sc = sc32(2 * (int) (w >> 28) + 1);
+    }
+};
+
+template <> struct Fmt32<21> {   // IQ3_S: d, qs[64], qh[8], signs[32], scales[4]
+    static constexpr int bytes = 110;
+    static constexpr float K = 1.0f;
+    static inline void decode(const uint8_t* b, int j, int half, __m256i& g, __m256i& sgn, __m256i& sc) {
+        const uint8_t* q = b + 2 + 16 * j + 8 * half;
+        const uint32_t h = b[66 + 2 * j + half];
+#define G3(k) (int) iq3s_grid[q[k] | (((h >> k) & 1u) << 8)]
+        g = _mm256_set_epi32(G3(7), G3(6), G3(5), G3(4), G3(3), G3(2), G3(1), G3(0));
+#undef G3
+        const uint64_t m = u64(b + 74 + 8 * j);
+        sgn = sgn_vec(half ? (uint32_t) (m >> 32) : (uint32_t) m);
+        const uint8_t s = b[106 + j];
+        sc = sc32(half ? 2 * (s >> 4) + 1 : 2 * (s & 15) + 1);
+    }
+};
+
+template <int TY, int NT>
+inline void row_dot(const uint8_t* row, int nblocks, const block_q8_K* const* y, float* res) {
+    __m256 accf[NT];
+    for (int t = 0; t < NT; ++t) accf[t] = _mm256_setzero_ps();
+    for (int i = 0; i < nblocks; ++i) {
+        const uint8_t* blk = row + (size_t) i * Fmt32<TY>::bytes;
+        __m256i acci[NT];
+        for (int t = 0; t < NT; ++t) acci[t] = _mm256_setzero_si256();
+        for (int j = 0; j < 4; ++j) {
+            for (int half = 0; half < 2; ++half) {
+                __m256i g, sgn, sc;
+                Fmt32<TY>::decode(blk, j, half, g, sgn, sc);
+                const int off = 64 * j + 32 * half;
+                for (int t = 0; t < NT; ++t) {
+                    const __m256i yv = _mm256_loadu_si256((const __m256i*) (y[t][i].qs + off));
+                    const __m256i ys = _mm256_sign_epi8(yv, sgn);
+                    acci[t] = _mm256_add_epi32(acci[t], _mm256_madd_epi16(_mm256_maddubs_epi16(g, ys), sc));
+                }
+            }
+        }
+        const float dx = h2f(u16(blk)) * Fmt32<TY>::K;
+        for (int t = 0; t < NT; ++t)
+            accf[t] = _mm256_fmadd_ps(_mm256_set1_ps(dx * y[t][i].d), _mm256_cvtepi32_ps(acci[t]), accf[t]);
+    }
+    for (int t = 0; t < NT; ++t) res[t] = hsum8(accf[t]);
+}
+
+// ---- IQ2_XS (17) with ggml's vectorized sign decode (arch/x86/quants.c, ggml_vec_dot_iq2_xs_q8_K):
+// the 7-bit sign indices never touch ksigns_iq2xs - the 8th sign bit is reconstructed by parity (two
+// shifts, a xor and ggml's bit_helper pshufb) and the per-half sign vectors come out of shuffles of one
+// 32-byte load.  The 8 scale bytes become all 16 half-scales with ggml's unpack trick.  The grid lookups
+// stay scalar loads into set_epi64x (ggml does the same).  Every token still only pays a load, a sign,
+// a maddubs, a madd and an add per half, into alternating accumulators.
+template <int NT>
+inline void row_dot_iq2xs(const uint8_t* row, int nblocks, const block_q8_K* const* y, float* res) {
+    static const uint8_t bit_sel[32] = {
+        1, 2, 4, 8, 16, 32, 64, (uint8_t) 0x80, 1, 2, 4, 8, 16, 32, 64, (uint8_t) 0x80,
+        1, 2, 4, 8, 16, 32, 64, (uint8_t) 0x80, 1, 2, 4, 8, 16, 32, 64, (uint8_t) 0x80 };
+    static const char shuf_even[32] = {   // sign bits of u16 0,1,2,3 (even bytes of the sign register)
+        0, 0, 0, 0, 0, 0, 0, 0, 2, 2, 2, 2, 2, 2, 2, 2,
+        4, 4, 4, 4, 4, 4, 4, 4, 6, 6, 6, 6, 6, 6, 6, 6 };
+    static const char shuf_odd[32] = {    // sign bits of u16 4,5,6,7
+        8, 8, 8, 8, 8, 8, 8, 8, 10, 10, 10, 10, 10, 10, 10, 10,
+        12, 12, 12, 12, 12, 12, 12, 12, 14, 14, 14, 14, 14, 14, 14, 14 };
+    static const uint8_t bit_help[32] = {  // 0x80 when the 4-bit index has odd popcount (parity bit 7)
+        0x00, (uint8_t) 0x80, (uint8_t) 0x80, 0x00, (uint8_t) 0x80, 0x00, 0x00, (uint8_t) 0x80,
+        (uint8_t) 0x80, 0x00, 0x00, (uint8_t) 0x80, 0x00, (uint8_t) 0x80, (uint8_t) 0x80, 0x00,
+        0x00, (uint8_t) 0x80, (uint8_t) 0x80, 0x00, (uint8_t) 0x80, 0x00, 0x00, (uint8_t) 0x80,
+        (uint8_t) 0x80, 0x00, 0x00, (uint8_t) 0x80, 0x00, (uint8_t) 0x80, (uint8_t) 0x80, 0x00 };
+    static const uint8_t k_sc_shuffle[128] = {  // half h -> scale bytes 2h (lanes 0-7), 2h+1 (lanes 8-15)
+        0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1,
+        2, 2, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3, 3, 3, 3, 3,
+        4, 4, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5, 5,
+        6, 6, 6, 6, 6, 6, 6, 6, 7, 7, 7, 7, 7, 7, 7, 7,
+        8, 8, 8, 8, 8, 8, 8, 8, 9, 9, 9, 9, 9, 9, 9, 9,
+        10, 10, 10, 10, 10, 10, 10, 10, 11, 11, 11, 11, 11, 11, 11, 11,
+        12, 12, 12, 12, 12, 12, 12, 12, 13, 13, 13, 13, 13, 13, 13, 13,
+        14, 14, 14, 14, 14, 14, 14, 14, 15, 15, 15, 15, 15, 15, 15, 15 };
+    const __m256i bsel = _mm256_loadu_si256((const __m256i*) bit_sel);
+    const __m256i shev = _mm256_loadu_si256((const __m256i*) shuf_even);
+    const __m256i shod = _mm256_loadu_si256((const __m256i*) shuf_odd);
+    const __m256i bhelp = _mm256_loadu_si256((const __m256i*) bit_help);
+    const __m256i m511 = _mm256_set1_epi16(511);
+    const __m128i m4 = _mm_set1_epi8(0xf), m1 = _mm_set1_epi8(1);
+    const __m256i one8 = _mm256_set1_epi8(1);
+
+    __m256 accf[NT];
+    for (int t = 0; t < NT; ++t) accf[t] = _mm256_setzero_ps();
+    for (int i = 0; i < nblocks; ++i) {
+        const uint8_t* blk = row + (size_t) i * 74;
+        // the 8 scale bytes -> 16 half-scales of 2*s+1, interleaved [a0, b0, a1, b1, ...] (ggml's unpack)
+        __m128i st = _mm_set1_epi64x((long long) u64(blk + 66));
+        st = _mm_unpacklo_epi8(_mm_and_si128(st, m4), _mm_and_si128(_mm_srli_epi16(st, 4), m4));
+        const __m128i scales = _mm_add_epi8(_mm_slli_epi16(st, 1), m1);
+        __m256i acc[NT][2];
+        for (int t = 0; t < NT; ++t) acc[t][0] = acc[t][1] = _mm256_setzero_si256();
+        for (int jj = 0; jj < 2; ++jj) {   // 128 values (16 u16) per step
+            const __m256i q2 = _mm256_loadu_si256((const __m256i*) (blk + 2 + 32 * jj));
+            const __m256i p7 = _mm256_srli_epi16(q2, 9);              // 7 sign bits per u16, even bytes
+            const __m256i p3 = _mm256_srli_epi16(q2, 13);             // sign index bits 4-6
+            const __m256i fsb = _mm256_or_si256(p7,
+                _mm256_shuffle_epi8(bhelp, _mm256_xor_si256(p7, p3)));  // + the parity sign bit 7
+            const __m256i s01 = _mm256_broadcastsi128_si256(_mm256_castsi256_si128(fsb)); // u16 0-7
+            const __m256i s23 = _mm256_broadcastsi128_si256(_mm256_extracti128_si256(fsb, 1)); // u16 8-15
+            alignas(32) uint16_t gi[16];
+            _mm256_store_si256((__m256i*) gi, _mm256_and_si256(q2, m511));
+            for (int h = 0; h < 4; ++h) {   // the four 32-value halves of this group
+                const int H = 4 * jj + h;   // global half index -> scales byte pair
+                const __m256i g = _mm256_set_epi64x((long long) iq2xs_grid[gi[4 * h + 3]],
+                                                   (long long) iq2xs_grid[gi[4 * h + 2]],
+                                                   (long long) iq2xs_grid[gi[4 * h + 1]],
+                                                   (long long) iq2xs_grid[gi[4 * h]]);
+                const __m256i sb = _mm256_shuffle_epi8(h < 2 ? s01 : s23, (h & 1) ? shod : shev);
+                const __m256i sgn = _mm256_or_si256(
+                    _mm256_cmpeq_epi8(_mm256_and_si256(sb, bsel), bsel), one8);
+                const __m256i sc = _mm256_cvtepi8_epi16(
+                    _mm_shuffle_epi8(scales, _mm_loadu_si128((const __m128i*) k_sc_shuffle + H)));
+                const int off = 128 * jj + 32 * h;
+                for (int t = 0; t < NT; ++t) {
+                    const __m256i yv = _mm256_loadu_si256((const __m256i*) (y[t][i].qs + off));
+                    const __m256i ys = _mm256_sign_epi8(yv, sgn);
+                    acc[t][h & 1] = _mm256_add_epi32(acc[t][h & 1],
+                        _mm256_madd_epi16(_mm256_maddubs_epi16(g, ys), sc));
+                }
+            }
+        }
+        const float dx = h2f(u16(blk)) * 0.125f;
+        for (int t = 0; t < NT; ++t)
+            accf[t] = _mm256_fmadd_ps(_mm256_set1_ps(dx * y[t][i].d),
+                _mm256_cvtepi32_ps(_mm256_add_epi32(acc[t][0], acc[t][1])), accf[t]);
+    }
+    for (int t = 0; t < NT; ++t) res[t] = hsum8(accf[t]);
+}
+
+template <int TY, int NT>
+inline void row_dot_any(const uint8_t* row, int nblocks, const block_q8_K* const* y, float* res) {
+    if constexpr (TY == 17) row_dot_iq2xs<NT>(row, nblocks, y, res);
+    else                    row_dot<TY, NT>(row, nblocks, y, res);
+}
+
+template <int TY, int NT>
+void gu_rows(const uint8_t* blob, size_t gu_row, size_t up_off, int n, const void* const* act, float* const* ff,
+             int r0, int r1) {
+    const block_q8_K* y[NT];
+    for (int t = 0; t < NT; ++t) y[t] = (const block_q8_K*) act[t];
+    const int nb = n / QK_K;
+    float g[NT], u[NT];
+    for (int r = r0; r < r1; ++r) {
+        row_dot_any<TY, NT>(blob + (size_t) r * gu_row, nb, y, g);
+        row_dot_any<TY, NT>(blob + up_off + (size_t) r * gu_row, nb, y, u);
+        for (int t = 0; t < NT; ++t) ff[t][r] = (g[t] / (1.f + std::exp(-g[t]))) * u[t];
+    }
+}
+
+template <int TY, int NT>
+void dot_rows(const uint8_t* w, size_t row_bytes, int n, const void* const* act, float* const* out, int r0, int r1) {
+    const block_q8_K* y[NT];
+    for (int t = 0; t < NT; ++t) y[t] = (const block_q8_K*) act[t];
+    float res[NT];
+    for (int r = r0; r < r1; ++r) {
+        row_dot_any<TY, NT>(w + (size_t) r * row_bytes, n / QK_K, y, res);
+        for (int t = 0; t < NT; ++t) out[t][r] = res[t];
+    }
+}
+
+template <int TY>
+void gu_rows_nt(int nt, const uint8_t* blob, size_t gu_row, size_t up_off, int n, const void* const* act,
+                float* const* ff, int r0, int r1) {
+    switch (nt) {
+        case 1: gu_rows<TY, 1>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 2: gu_rows<TY, 2>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 3: gu_rows<TY, 3>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 4: gu_rows<TY, 4>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 5: gu_rows<TY, 5>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 6: gu_rows<TY, 6>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 7: gu_rows<TY, 7>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        default: gu_rows<TY, 8>(blob, gu_row, up_off, n, act, ff, r0, r1); break;
+    }
+}
+
+template <int TY>
+void dot_rows_nt(int nt, const uint8_t* w, size_t row_bytes, int n, const void* const* act, float* const* out, int r0,
+                 int r1) {
+    switch (nt) {
+        case 1: dot_rows<TY, 1>(w, row_bytes, n, act, out, r0, r1); break;
+        case 2: dot_rows<TY, 2>(w, row_bytes, n, act, out, r0, r1); break;
+        case 3: dot_rows<TY, 3>(w, row_bytes, n, act, out, r0, r1); break;
+        case 4: dot_rows<TY, 4>(w, row_bytes, n, act, out, r0, r1); break;
+        default: for (int t0 = 0; t0 < nt; t0 += 4) {
+            const int k = nt - t0 < 4 ? nt - t0 : 4;
+            dot_rows_nt<TY>(k, w, row_bytes, n, act + t0, out + t0, r0, r1);
+        }
+    }
+}
+
+}  // namespace
+
+bool iq256_supported(int type) noexcept {
+    return type == 16 || type == 17 || type == 18 || type == 21 || type == 22;
+}
+
+void iq256_gu_rows(int type, const uint8_t* blob, size_t gu_row, size_t up_off, int n, const void* const* act, int nt,
+                   float* const* ff, int r0, int r1) {
+    switch (type) {
+        case 16: gu_rows_nt<16>(nt, blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 17: gu_rows_nt<17>(nt, blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 18: gu_rows_nt<18>(nt, blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 21: gu_rows_nt<21>(nt, blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        case 22: gu_rows_nt<22>(nt, blob, gu_row, up_off, n, act, ff, r0, r1); break;
+        default: break;
+    }
+}
+
+void iq256_rows(int type, const uint8_t* w, size_t row_bytes, int n, const void* const* act, int nt, float* const* out,
+                int r0, int r1) {
+    switch (type) {
+        case 16: dot_rows_nt<16>(nt, w, row_bytes, n, act, out, r0, r1); break;
+        case 17: dot_rows_nt<17>(nt, w, row_bytes, n, act, out, r0, r1); break;
+        case 18: dot_rows_nt<18>(nt, w, row_bytes, n, act, out, r0, r1); break;
+        case 21: dot_rows_nt<21>(nt, w, row_bytes, n, act, out, r0, r1); break;
+        case 22: dot_rows_nt<22>(nt, w, row_bytes, n, act, out, r0, r1); break;
+        default: break;
+    }
+}
+
+// ---- IQ4_NL (type 20): 32-value blocks of 18 bytes (f16 d + 16 nibble bytes) against Q8_0 activations.
+// ggml-cpu's own AVX-2 dot (ggml_vec_dot_iq4_nl_q8_0) is single-token: every token redoes the nibble ->
+// kvalues_iq4nl pshufb decode and the |w| half of the signed x signed product.  Here both are computed once
+// per block; every token then costs a sign, a maddubs, a madd and an fmadd.  The arithmetic is ggml's - only
+// the order of the float additions differs.
+template <int NT>
+void iq4nl_rows(const uint8_t* w, size_t row_bytes, int n, const block_q8_0* const* y, float* const* out,
+                int r0, int r1) {
+    const __m128i values = _mm_loadu_si128((const __m128i*) kvalues_iq4nl);
+    const __m128i m4b = _mm_set1_epi8(0x0f);
+    const __m256i ones = _mm256_set1_epi16(1);
+    const int nb = n / QK4_NL;
+    for (int r = r0; r < r1; ++r) {
+        const uint8_t* row = w + (size_t) r * row_bytes;
+        __m256 accf[NT];
+        for (int t = 0; t < NT; ++t) accf[t] = _mm256_setzero_ps();
+        for (int ib = 0; ib < nb; ++ib) {
+            const uint8_t* blk = row + (size_t) ib * sizeof(block_iq4_nl);
+            const __m128i bits = _mm_loadu_si128((const __m128i*) (blk + 2));
+            const __m128i lo = _mm_and_si128(bits, m4b);                      // values 0..15
+            const __m128i hi = _mm_and_si128(_mm_srli_epi16(bits, 4), m4b);    // values 16..31
+            const __m256i q4 = _mm256_inserti128_si256(_mm256_castsi128_si256(_mm_shuffle_epi8(values, lo)),
+                                                       _mm_shuffle_epi8(values, hi), 1);
+            const __m256i aq = _mm256_sign_epi8(q4, q4);   // |w|: the unsigned operand of maddubs
+            const float dx = h2f(u16(blk));
+            for (int t = 0; t < NT; ++t) {
+                const block_q8_0& b = y[t][ib];
+                const __m256i q8 = _mm256_loadu_si256((const __m256i*) b.qs);
+                const __m256i p = _mm256_madd_epi16(_mm256_maddubs_epi16(aq, _mm256_sign_epi8(q8, q4)), ones);
+                accf[t] = _mm256_fmadd_ps(_mm256_set1_ps(dx * h2f(b.d)), _mm256_cvtepi32_ps(p), accf[t]);
+            }
+        }
+        for (int t = 0; t < NT; ++t) out[t][r] = hsum8(accf[t]);
+    }
+}
+
+void iq4nl256_down_rows(const uint8_t* w, size_t row_bytes, int n, const void* const* hq, int nt, float* const* out,
+                        int r0, int r1) {
+    const block_q8_0* y[8];
+    for (int t = 0; t < nt; ++t) y[t] = (const block_q8_0*) hq[t];
+    switch (nt) {
+        case 1: iq4nl_rows<1>(w, row_bytes, n, y, out, r0, r1); break;
+        case 2: iq4nl_rows<2>(w, row_bytes, n, y, out, r0, r1); break;
+        case 3: iq4nl_rows<3>(w, row_bytes, n, y, out, r0, r1); break;
+        case 4: iq4nl_rows<4>(w, row_bytes, n, y, out, r0, r1); break;
+        default: for (int t0 = 0; t0 < nt; t0 += 4) {
+            const int k = nt - t0 < 4 ? nt - t0 : 4;
+            iq4nl256_down_rows(w, row_bytes, n, hq + t0, k, out + t0, r0, r1);
+        }
+    }
+}
+
+}  // namespace strata::kernels::cpu
@@ -4,6 +4,7 @@
 #include "strata/kernels/cpu/native_expert.hpp"
 #include "strata/kernels/cpu/expert.hpp"
 #include "strata/kernels/cpu/iq_avx512.hpp"
+#include "strata/kernels/cpu/iq_avx2.hpp"
 #include "strata/kernels/cpu/expert_layout.hpp"

 #include "ggml.h"
@@ -76,12 +77,21 @@ void native_quant_h(const NativeFmt& f, const float* h, void* dst) {

 void native_gu_rows(const NativeFmt& f, const uint8_t* blob, const void* const* act, int nt, float* const* ff,
                    int r0, int r1) {
-    // the AVX-512 kernels decode the weights once for all tokens: 2.0-2.4x ggml-cpu at three tokens, no faster at
-    // one (both are bound by the codebook lookups, ~5 GB/s per core), measured by native_expert_parity
+    // the multi-token kernels decode the weights once for all tokens: 2.0-2.4x ggml-cpu at three tokens, no faster
+    // at one (all are bound by the codebook lookups, ~5 GB/s per core), measured by native_expert_parity.  AVX-512
+    // first, then the AVX-2 one (Zen 2/3, Intel 12th-14th gen); STRATA_NO_IQ512 / STRATA_NO_IQ256 fall back to
+    // ggml-cpu's single-token vec_dot.
    static const bool avx512 = cpu_avx512_ok() && std::getenv("STRATA_NO_IQ512") == nullptr;
-    if (avx512 && nt >= 2 && iq512_supported(f.gu_type)) {   // one token: ggml-cpu is as fast or faster
-        iq512_gu_rows(f.gu_type, blob, f.gu_row, f.up_off, (int) f.n_embd, act, nt, ff, r0, r1);
-        return;
+    static const bool avx2 = std::getenv("STRATA_NO_IQ256") == nullptr;
+    if (nt >= 2 && iq512_supported(f.gu_type)) {   // one token: ggml-cpu is as fast or faster
+        if (avx512) {
+            iq512_gu_rows(f.gu_type, blob, f.gu_row, f.up_off, (int) f.n_embd, act, nt, ff, r0, r1);
+            return;
+        }
+        if (avx2) {
+            iq256_gu_rows(f.gu_type, blob, f.gu_row, f.up_off, (int) f.n_embd, act, nt, ff, r0, r1);
+            return;
+        }
    }
    const ggml_vec_dot_t dot = traits(f.gu_type)->vec_dot;
    const int n = (int) f.n_embd;
@@ -99,6 +109,13 @@ void native_gu_rows(const NativeFmt& f, const uint8_t* blob, const void* const*

 void native_down_rows(const NativeFmt& f, const uint8_t* blob, const void* const* hq, int nt, float* const* out,
                      int r0, int r1) {
+    // IQ4_NL down rows: the AVX-2 multi-token kernel decodes the nibbles and absolutises the weights once per
+    // block instead of once per token; ggml-cpu's dot is single-token.  STRATA_NO_IQ4NL falls back to it.
+    static const bool iq4nl_mt = std::getenv("STRATA_NO_IQ4NL") == nullptr;
+    if (nt >= 2 && f.d_type == 20 && iq4nl_mt) {
+        iq4nl256_down_rows(blob + f.down_off, f.d_row, (int) f.n_ff, hq, nt, out, r0, r1);
+        return;
+    }
    const ggml_vec_dot_t dot = traits(f.d_type)->vec_dot;
    const int n = (int) f.n_ff;
    for (int r = r0; r < r1; ++r) {
@@ -55,13 +55,38 @@ std::vector<int> physical_cores(bool skip_first) {
        }
    }
 #else
+    // The logical CPUs this process may run on, ONE PER PHYSICAL CORE (issue #40): SMT siblings share a core's
+    // load/store bandwidth, so a worker on each would put two workers on one core, as the Windows branch above
+    // explains.  sysfs names each CPU's (package, core); the first allowed CPU of each pair is kept, so a taskset
+    // that leaves out the first sibling still gets its core.  Without sysfs every allowed CPU counts, as before.
+    auto topo = [](int cpu, const char* what) -> long {
+        char path[96];
+        std::snprintf(path, sizeof path, "/sys/devices/system/cpu/cpu%d/topology/%s", cpu, what);
+        long v = -1;
+        if (std::FILE* f = std::fopen(path, "r")) {
+            if (std::fscanf(f, "%ld", &v) != 1) v = -1;
+            std::fclose(f);
+        }
+        return v;
+    };
+    std::vector<int> allowed;
    cpu_set_t set;
    CPU_ZERO(&set);
    if (sched_getaffinity(0, sizeof set, &set) == 0) {
        for (int i = 0; i < CPU_SETSIZE; ++i)
-            if (CPU_ISSET(i, &set)) cores.push_back(i);
+            if (CPU_ISSET(i, &set)) allowed.push_back(i);
    } else {
-        for (unsigned i = 0; i < std::thread::hardware_concurrency(); ++i) cores.push_back((int) i);
+        for (unsigned i = 0; i < std::thread::hardware_concurrency(); ++i) allowed.push_back((int) i);
+    }
+    std::vector<std::pair<long, long>> seen;
+    for (int cpu : allowed) {
+        const long pkg = topo(cpu, "physical_package_id"), core = topo(cpu, "core_id");
+        if (pkg >= 0 && core >= 0) {
+            const std::pair<long, long> key{pkg, core};
+            if (std::find(seen.begin(), seen.end(), key) != seen.end()) continue;   // an SMT sibling
+            seen.push_back(key);
+        }
+        cores.push_back(cpu);
    }
 #endif
    if (skip_first && !cores.empty()) cores.erase(cores.begin());
@@ -9,6 +9,8 @@
 #include "strata/kernels/cpu/native_expert.hpp"
 #include "strata/kernels/cpu/expert.hpp"
 #include "strata/kernels/cpu/iq_avx512.hpp"
+#include "strata/kernels/cpu/iq_avx2.hpp"
+#include "strata/kernels/cpu/expert_layout.hpp"
 #include "ggml-cpu.h"
 #include "strata/kernels/iq_kernels.hpp"

@@ -33,6 +35,7 @@ static double rel(const std::vector<float>& a, const std::vector<float>& b) {
 }

 int main(int argc, char** argv) {
+    setvbuf(stdout, nullptr, _IONBF, 0);   // keep the trail on a crash
    if (argc < 2) { std::fprintf(stderr, "usage: native_expert_parity <shard1.gguf> [layer ...]\n"); return 2; }
    strata::GgufFile gguf(argv[1]);
    std::vector<int> layers;
@@ -103,40 +106,50 @@ int main(int argc, char** argv) {
            }
            cpu::native_gu_rows(f, blob.data(), a, NT, ffp, 0, (int) FF);
            if (cpu::iq512_supported(f.gu_type)) {
-                // the AVX-512 rows against ggml's own vec_dot, same Q8_K activations: float-order differences only
-                std::vector<float> g512((size_t) NT * FF), gref((size_t) NT * FF);
-                float* gp[NT];
-                for (int k = 0; k < NT; ++k) gp[k] = g512.data() + k * FF;
-                cpu::iq512_rows(f.gu_type, blob.data(), f.gu_row, (int) H, a, NT, gp, 0, (int) FF);
+                // ggml's own vec_dot, same Q8_K activations: the reference for both multi-token kernels
+                // (float-order differences only)
                const auto* tc = ggml_get_type_traits_cpu((ggml_type) f.gu_type);
+                std::vector<float> gref((size_t) NT * FF);
                for (int k = 0; k < NT; ++k)
                    for (int64_t r = 0; r < FF; ++r)
                        tc->vec_dot((int) H, &gref[k * FF + r], 0, blob.data() + r * f.gu_row, 0, a[k], 0, 1);
-                {   // single-thread throughput, the weights cache-resident (compute bound)
-                    float* f1[1] = {ffp[0]};
+                double usg = 0.0;
+                {   // ggml single-token throughput, the weights cache-resident (compute bound)
                    const int it = 50;
                    auto t0 = std::chrono::steady_clock::now();
-                    for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, 1, f1, 0, (int) FF);
-                    auto t1 = std::chrono::steady_clock::now();
                    for (int i = 0; i < it; ++i)
                        for (int64_t r = 0; r < FF; ++r) {
                            float gg, uu;
                            tc->vec_dot((int) H, &gg, 0, blob.data() + r * f.gu_row, 0, a[0], 0, 1);
                            tc->vec_dot((int) H, &uu, 0, blob.data() + f.up_off + r * f.gu_row, 0, a[0], 0, 1);
                        }
-                    auto t2 = std::chrono::steady_clock::now();
-                    const double us512 = std::chrono::duration<double, std::micro>(t1 - t0).count() / it;
-                    const double usg = std::chrono::duration<double, std::micro>(t2 - t1).count() / it;
-                    std::printf("          gate+up one token one thread: avx512 %.0f us (%.2f GB/s), ggml %.0f us (%.2f GB/s)\n",
-                                us512, 2.0 * f.up_off / us512 / 1e3, usg, 2.0 * f.up_off / usg / 1e3);
-                    auto t3 = std::chrono::steady_clock::now();
-                    for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, NT, ffp, 0, (int) FF);
-                    auto t4 = std::chrono::steady_clock::now();
-                    std::printf("          gate+up %d tokens one thread: avx512 %.0f us vs ggml %d x %.0f us\n", NT,
-                                std::chrono::duration<double, std::micro>(t4 - t3).count() / it, NT, usg);
+                    auto t1 = std::chrono::steady_clock::now();
+                    usg = std::chrono::duration<double, std::micro>(t1 - t0).count() / it;
                }
-                std::printf("          %s AVX-512 gate rows vs ggml vec_dot: rel %.2e\n",
-                            ggml_type_name((ggml_type) f.gu_type), rel(g512, gref));
+                // one multi-token kernel: parity against ggml's vec_dot, then one-token and NT-token throughput
+                auto check = [&](const char* tag, bool is512) {
+                    std::vector<float> g((size_t) NT * FF);
+                    float* gp[NT];
+                    for (int k = 0; k < NT; ++k) gp[k] = g.data() + k * FF;
+                    (is512 ? cpu::iq512_rows : cpu::iq256_rows)(f.gu_type, blob.data(), f.gu_row, (int) H, a, NT, gp, 0, (int) FF);
+                    std::printf("          %s %s gate rows vs ggml vec_dot: rel %.2e\n",
+                                ggml_type_name((ggml_type) f.gu_type), tag, rel(g, gref));
+                    float* f1[1] = {ffp[0]};
+                    const int it = 50;
+                    const auto gu = is512 ? cpu::iq512_gu_rows : cpu::iq256_gu_rows;
+                    auto t0 = std::chrono::steady_clock::now();
+                    for (int i = 0; i < it; ++i) gu(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, 1, f1, 0, (int) FF);
+                    auto t1 = std::chrono::steady_clock::now();
+                    for (int i = 0; i < it; ++i) gu(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, NT, ffp, 0, (int) FF);
+                    auto t2 = std::chrono::steady_clock::now();
+                    const double us1 = std::chrono::duration<double, std::micro>(t1 - t0).count() / it;
+                    const double usn = std::chrono::duration<double, std::micro>(t2 - t1).count() / it;
+                    std::printf("          gate+up one token one thread: %s %.0f us (%.2f GB/s), ggml %.0f us (%.2f GB/s)\n",
+                                tag, us1, 2.0 * f.up_off / us1 / 1e3, usg, 2.0 * f.up_off / usg / 1e3);
+                    std::printf("          gate+up %d tokens one thread: %s %.0f us vs ggml %d x %.0f us\n", NT, tag, usn, NT, usg);
+                };
+                if (cpu::cpu_avx512_ok()) check("avx512", true);   // guarded: the binary runs on AVX-2 CPUs too
+                check("avx2", false);
            }
            for (int k = 0; k < NT; ++k) {
                cpu::native_quant_h(f, ff[k].data(), hq[k].data());
@@ -145,18 +158,51 @@ int main(int argc, char** argv) {
            }
            cpu::native_down_rows(f, blob.data(), hp, NT, op, 0, (int) H);
            if (f.d_type == 42) {
-                // (b2) the AVX-512 GGUF-layout Q2_0 kernel the pool uses for Q2_0 down projections
+                // (b2) the GGUF-layout Q2_0 kernel the pool uses for Q2_0 down projections - the AVX-512 one
+                // where the CPU has it, the AVX-2 one (q2_avx2.cpp) where it does not.  Calling the AVX-512
+                // kernel unconditionally faults on a Zen 2/3 CPU.
                std::vector<cpu::ActQ> a2(NT);
                const cpu::ActQ* ap[NT];
                std::vector<float> alt((size_t) NT * H);
                float* altp[NT];
                for (int k = 0; k < NT; ++k) {
-                    cpu::act_quant_q8_1(ff[k].data(), (int) FF, a2[k]);
+                    cpu::act_quant_any(ff[k].data(), (int) FF, a2[k]);   // gated: AVX-512 or AVX-2 per CPU
                    ap[k] = &a2[k];
                    altp[k] = alt.data() + k * H;
                }
-                cpu::q2_0_gguf_rows_multi(blob.data() + f.down_off, f.d_row, (int) (FF / 64), ap, NT, altp, 0, (int) H);
-                std::printf("          q2_0 AVX-512 down vs ggml down: rel %.2e\n", rel(alt, got_c));
+                cpu::q2_rows_any(blob.data() + f.down_off, f.d_row, (int) (FF / 64), ap, NT, altp, 0, (int) H);
+                std::printf("          q2_0 %s down vs ggml down: rel %.2e\n",
+                            cpu::cpu_avx512_ok() ? "AVX-512" : "AVX-2", rel(alt, got_c));
+            }
+            if (f.d_type == 20) {
+                // (b3) the IQ4_NL multi-token AVX-2 kernel the pool now uses for IQ4_NL down projections,
+                // against ggml-cpu's single-token vec_dot on the SAME Q8_0 activations (h), plus timing.
+                std::vector<float> alt((size_t) NT * H), ref((size_t) NT * H);
+                float* altp[NT];
+                for (int k = 0; k < NT; ++k) altp[k] = alt.data() + k * H;
+                cpu::iq4nl256_down_rows(blob.data() + f.down_off, f.d_row, (int) FF, hp, NT, altp, 0, (int) H);
+                const auto* tdc = ggml_get_type_traits_cpu((ggml_type) f.d_type);
+                for (int k = 0; k < NT; ++k)
+                    for (int64_t r = 0; r < H; ++r)
+                        tdc->vec_dot((int) FF, ref.data() + k * H + r, 0,
+                                     blob.data() + f.down_off + (size_t) r * f.d_row, 0, hp[k], 0, 1);
+                std::printf("          iq4_nl AVX-2 multi-token down vs ggml down: rel %.2e\n", rel(alt, ref));
+                const int it = 50;
+                const auto t0 = std::chrono::steady_clock::now();
+                for (int i = 0; i < it; ++i)
+                    cpu::iq4nl256_down_rows(blob.data() + f.down_off, f.d_row, (int) FF, hp, NT, altp, 0, (int) H);
+                const auto t1 = std::chrono::steady_clock::now();
+                for (int i = 0; i < it; ++i)
+                    for (int k = 0; k < NT; ++k)
+                        for (int64_t r = 0; r < H; ++r)
+                            tdc->vec_dot((int) FF, ref.data() + k * H + r, 0,
+                                        blob.data() + f.down_off + (size_t) r * f.d_row, 0, hp[k], 0, 1);
+                const auto t2 = std::chrono::steady_clock::now();
+                const double usn = std::chrono::duration<double, std::micro>(t1 - t0).count() / it;
+                const double usg = std::chrono::duration<double, std::micro>(t2 - t1).count() / it;
+                std::printf("          down %d tokens one thread: avx2 %.0f us (%.2f GB/s) vs ggml %.0f us (%.2f GB/s)\n",
+                            NT, usn, (double) (f.bytes - f.down_off) / usn / 1e3, usg,
+                            (double) (f.bytes - f.down_off) / usg / 1e3);
            }
        }
        // (c) the GPU: one group holding the NT entries
@@ -1,6 +1,7 @@
 // src/prefill/kernels.cu - see include/strata/prefill/kernels.hpp.
 #include "strata/prefill/kernels.hpp"
 #include "strata/kernels/mrope.hpp"
+#include "strata/kernels/router_top10.hpp"

 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
@@ -201,31 +202,32 @@ __global__ void __launch_bounds__(S * RG) gdn_rec_kernel(float* __restrict__ sta
 }

 // ---------------------------------------------------------------- MoE
+template <int REG>
 __global__ void route_kernel(const float* __restrict__ logits, int32_t* __restrict__ ids, float* __restrict__ wout,
                             int64_t T) {
    const int64_t t = (int64_t) blockIdx.x * (blockDim.x >> 5) + (threadIdx.x >> 5);
    if (t >= T) return;
    const int lane = threadIdx.x & 31;
-    const float* lg = logits + t * 512;
-    float v[16];
+    const float* lg = logits + t * (REG * 32);
+    float v[REG];
 #pragma unroll
-    for (int i = 0; i < 16; ++i) v[i] = lg[lane + i * 32];
+    for (int i = 0; i < REG; ++i) v[i] = lg[lane + i * 32];
    float mx = -INFINITY;
 #pragma unroll
-    for (int i = 0; i < 16; ++i) mx = fmaxf(mx, v[i]);
+    for (int i = 0; i < REG; ++i) mx = fmaxf(mx, v[i]);
    mx = warp_max(mx);
    float sum = 0.0f;
 #pragma unroll
-    for (int i = 0; i < 16; ++i) { v[i] = expf(v[i] - mx); sum += v[i]; }
+    for (int i = 0; i < REG; ++i) { v[i] = expf(v[i] - mx); sum += v[i]; }
    const float rcp = 1.0f / warp_sum(sum);
 #pragma unroll
-    for (int i = 0; i < 16; ++i) { v[i] *= rcp; if (isnan(v[i])) v[i] = -FLT_MAX; }
+    for (int i = 0; i < REG; ++i) { v[i] *= rcp; if (isnan(v[i])) v[i] = -FLT_MAX; }
    float selected = 0.0f, selected_sum = 0.0f;
    for (int rank = 0; rank < 10; ++rank) {
        float best = v[0];
        int ex = lane;
 #pragma unroll
-        for (int i = 1; i < 16; ++i) if (v[i] > best) { best = v[i]; ex = lane + i * 32; }
+        for (int i = 1; i < REG; ++i) if (v[i] > best) { best = v[i]; ex = lane + i * 32; }
 #pragma unroll
        for (int m = 16; m; m >>= 1) {
            const float ob = __shfl_xor_sync(0xffffffffu, best, m);
@@ -455,8 +457,13 @@ void gdn_recurrence(float* state, const float* h, const float* gate, const float
    gdn_rec_kernel<<<HV, dim3(S, RG), 0, (cudaStream_t) stream>>>(state, h, gate, beta, z, gamma, eps, y, y16, T);
    check("gdn_recurrence");
 }
-void route(const float* logits, int32_t* ids, float* weights, int64_t T, void* stream) {
-    route_kernel<<<(unsigned) ((T + 7) / 8), 256, 0, (cudaStream_t) stream>>>(logits, ids, weights, T);
+void route(const float* logits, int32_t* ids, float* weights, int64_t T, int64_t n_expert, void* stream) {
+    if (n_expert == 512)
+        route_kernel<16><<<(unsigned) ((T + 7) / 8), 256, 0, (cudaStream_t) stream>>>(logits, ids, weights, T);
+    else if (n_expert == 256)
+        route_kernel<8><<<(unsigned) ((T + 7) / 8), 256, 0, (cudaStream_t) stream>>>(logits, ids, weights, T);
+    else
+        strata::kernels::router_top10(logits, (int) T, (int) n_expert, 10, ids, weights, stream);
    check("route");
 }
 void blob_dequant(const uint8_t* blob, uint16_t* gu16, uint16_t* down16, void* stream) {
@@ -407,10 +407,10 @@ const MmqPlan& mmq_plan() {
    }();
    return plan;
 }
-uint64_t moe_set_bytes(size_t T) {
+uint64_t moe_set_bytes(size_t T, int64_t n_expert) {
    const MmqPlan& mp = mmq_plan();
    Alloc a; a.count_only = true; bool ok = true;
-    a.take<float>(T * NE, ok); a.take<float>(T * K, ok); a.take<int32_t>(T * K, ok); a.take<int32_t>(T * K, ok);
+    a.take<float>(T * n_expert, ok); a.take<float>(T * K, ok); a.take<int32_t>(T * K, ok); a.take<int32_t>(T * K, ok);
    a.take<int32_t>(T * K, ok);
    if (mp.fallback) a.take<uint16_t>(T * K * N, ok);
    a.take<float>(T * K * 1280, ok);
@@ -432,7 +432,7 @@ bool Prefill::init(const core::WeightTable& wt, const core::ModelGeometry& g, co
    Impl& m = *impl_;
    m.wt = &wt; m.g = &g; m.ss = &ss; m.src = src; m.cache = cache; m.host_res = host_res;
    m.T = chunk; m.cs = (cudaStream_t) stream; m.stats = &stats_;
-    if (g.n_embd != N || g.hc != HC || g.hc_lr != LR || g.n_expert != NE || ss.k != K) {
+    if (g.n_embd != N || g.hc != HC || g.hc_lr != LR || g.n_expert < 1 || ss.k != K) {
        err = "prefill: geometry differs from the artifact's"; return false;
    }
    if (cudaStreamCreateWithFlags(&m.copy, cudaStreamNonBlocking) != cudaSuccess) { err = "prefill: copy stream"; return false; }
@@ -451,7 +451,7 @@ bool Prefill::init(const core::WeightTable& wt, const core::ModelGeometry& g, co
        if (!m.stager->init((size_t) MAXBLOB(), std::max(2, std::min(4, hw / 4)))) ok = false;
    }
    m.steps_host.resize(T * strata::kernels::kStepCount);
-    m.ids_host.resize(T * K); m.slot_host.resize(T * K); m.src_host.resize(T * K); m.cnt.resize(NE); m.off.resize(NE + 1);
+    m.ids_host.resize(T * K); m.slot_host.resize(T * K); m.src_host.resize(T * K); m.cnt.resize(m.g->n_expert); m.off.resize(m.g->n_expert + 1);
    for (int b = 0; b < 2; ++b) {
        if (!m.ple_emb_host[b] &&
            cudaHostAlloc((void**) &m.ple_emb_host[b], (size_t) T * N * 4, cudaHostAllocDefault) != cudaSuccess) {
@@ -513,7 +513,7 @@ bool Prefill::carve(size_t T, void* alloc) {
    {
        // one region for the attention half's and the MoE half's scratch (see gdn_set_bytes)
        const uint64_t region = std::max({gdn_set_bytes(T), qsa_set_bytes(T, m.cap, m.max_blocks, m.sel_batch,
-                                                                           m.attn_batch, s), moe_set_bytes(T)});
+                                                                           m.attn_batch, s), moe_set_bytes(T, m.g->n_expert)});
        uint8_t* base = o.take<uint8_t>((size_t) region, ok);
        m.region = base;
        m.region_bytes = region;
@@ -532,7 +532,7 @@ bool Prefill::carve(size_t T, void* alloc) {
        m.attn_scratch = b.take<float>((size_t) m.attn_batch * strata::kernels::qsa_decode_attn_scratch_floats(m.cap, s), ok);
        Alloc c;
        c.base = base; c.cap = region; c.owned = &m.owned;
-        m.logits = c.take<float>(T * NE, ok); m.w = c.take<float>(T * K, ok); m.ids = c.take<int32_t>(T * K, ok);
+        m.logits = c.take<float>(T * m.g->n_expert, ok); m.w = c.take<float>(T * K, ok); m.ids = c.take<int32_t>(T * K, ok);
        m.slot_dev = c.take<int32_t>(T * K, ok); m.src_dev = c.take<int32_t>(T * K, ok);
        const MmqPlan& mp = mmq_plan();
        m.Xs = mp.fallback ? c.take<uint16_t>(T * K * N, ok) : nullptr;
@@ -552,7 +552,7 @@ bool Prefill::carve(size_t T, void* alloc) {
    if (mmq_plan().any) {
        const MmqPlan& mp = mmq_plan();
        m.ids_identity = o.take<int32_t>(T * K, ok);
-        m.bounds_dev = o.take<int32_t>((size_t) (2 * (NE + NE / MMQ_GROUP + 2)), ok);
+        m.bounds_dev = o.take<int32_t>((size_t) (2 * (m.g->n_expert + m.g->n_expert / MMQ_GROUP + 2)), ok);
        m.grp_gu = o.take<uint8_t>(MMQ_GROUP * mp.gu_max + MMQ_TAIL, ok);
        m.grp_d = o.take<uint8_t>(MMQ_GROUP * mp.d_max + MMQ_TAIL, ok);
        // (written at every run's start, not here: when serving, these are live expert-cache slots until a request
@@ -618,12 +618,12 @@ uint64_t Prefill::bytes_needed(const core::ModelGeometry& g, const core::Session
    const int64_t cap = strata::kernels::qsa_selection_width(strata::kernels::kTopkMaxCells, s);
    const int64_t max_blocks = ss.qsa_states[0].max_cells / s.idx_block + 2;
    o.take<uint8_t>((size_t) std::max({gdn_set_bytes(T), qsa_set_bytes(T, cap, max_blocks, 256, 32, s),
-                                       moe_set_bytes(T)}), ok);
+                                       moe_set_bytes(T, g.n_expert)}), ok);
    for (int i = 0; i < DQ; ++i) { o.take<uint16_t>(1280 * 2560, ok); o.take<uint16_t>(2560 * 640, ok); }
    if (mmq_plan().any) {
        const MmqPlan& mp = mmq_plan();
        o.take<int32_t>(T * K, ok);
-        o.take<int32_t>((size_t) (2 * (NE + NE / MMQ_GROUP + 2)), ok);
+        o.take<int32_t>((size_t) (2 * (g.n_expert + g.n_expert / MMQ_GROUP + 2)), ok);
        o.take<uint8_t>(MMQ_GROUP * mp.gu_max + MMQ_TAIL, ok);
        o.take<uint8_t>(MMQ_GROUP * mp.d_max + MMQ_TAIL, ok);
    }
@@ -809,8 +809,8 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
            std::vector<Stager::Job> js;
            for (int64_t l = 0; l < g.n_layers; ++l) {
                seq_start[(size_t) l] = seq.size();
-                for (int32_t e = 0; e < NE; ++e) {
-                    if (m.host_res && m.cache && m.host_res[(size_t) l * NE + e] >= 0) continue;
+                for (int32_t e = 0; e < m.g->n_expert; ++e) {
+                    if (m.host_res && m.cache && m.host_res[(size_t) l * m.g->n_expert + e] >= 0) continue;
                    const uint8_t* b = m.src->blob(l, e);
                    if (!b) { err = "prefill: expert source has no blob"; return false; }
                    int job = -1;
@@ -1095,7 +1095,7 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                    if (!wr || !wgi || !wsg || !wsu || !wsd) return false;
                    pt.mark(kPfRouter, cs);
                    if (!bf16_proj(m.gemm, wr, m.mixed_bf, m.logits, T, v.name("ffn_gate_inp.weight"), err)) return false;
-                    route(m.logits, m.ids, m.w, T, m.cs);
+                    route(m.logits, m.ids, m.w, T, m.g->n_expert, m.cs);
                    // the shared expert and its scalar gate
                    if (!native_proj(m.gemm, wsg, m.mixed_h, m.sgate, T, v.name("ffn_gate_shexp.weight"), err)) return false;
                    if (!native_proj(m.gemm, wsu, m.mixed_h, m.sup, T, v.name("ffn_up_shexp.weight"), err)) return false;
@@ -1111,11 +1111,11 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                    std::fill(m.cnt.begin(), m.cnt.end(), 0);
                    for (int64_t i = 0; i < T * K; ++i) {
                        const int32_t e = m.ids_host[(size_t) i];
-                        if (e < 0 || e >= NE) { err = "prefill: routed id out of range"; return false; }
+                        if (e < 0 || e >= m.g->n_expert) { err = "prefill: routed id out of range"; return false; }
                        ++m.cnt[(size_t) e];
                    }
                    m.off[0] = 0;
-                    for (int64_t e = 0; e < NE; ++e) m.off[(size_t) e + 1] = m.off[(size_t) e] + m.cnt[(size_t) e];
+                    for (int64_t e = 0; e < m.g->n_expert; ++e) m.off[(size_t) e + 1] = m.off[(size_t) e] + m.cnt[(size_t) e];
                    std::vector<int32_t> fill(m.off.begin(), m.off.end() - 1);
                    for (int64_t i = 0; i < T * K; ++i) {
                        const int32_t e = m.ids_host[(size_t) i];
@@ -1127,7 +1127,7 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                    cudaMemcpyAsync(m.src_dev, m.src_host.data(), (size_t) T * K * 4, cudaMemcpyHostToDevice, m.cs);
                    // the experts, in id order: resident ones from VRAM, the others through the staging ring
                    std::vector<int32_t> order;
-                    for (int32_t e = 0; e < NE; ++e) if (m.cnt[(size_t) e] > 0) order.push_back(e);
+                    for (int32_t e = 0; e < m.g->n_expert; ++e) if (m.cnt[(size_t) e] > 0) order.push_back(e);
                    const strata::kernels::cpu::ExpertLayout& lay = strata::kernels::cpu::expert_layout();
                    const bool use_mmq = mmq_plan().any && mmq_plan().layer[(size_t) l];
                    const int mmq_gt = lay.native ? lay.fmt[(size_t) l].gu_type : 42;
@@ -1162,7 +1162,7 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                        std::vector<Stager::Job> js;
                        for (size_t j = 0; j < order.size(); ++j) {
                            const int32_t e = order[j];
-                            if (m.host_res && m.cache && m.host_res[(size_t) l * NE + e] >= 0) continue;
+                            if (m.host_res && m.cache && m.host_res[(size_t) l * m.g->n_expert + e] >= 0) continue;
                            if (m.src->pinned(l, e)) continue;
                            const uint8_t* b = m.src->blob(l, e);
                            if (!b) { err = "prefill: expert source has no blob"; return false; }
@@ -1174,7 +1174,7 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                    StagerDone stager_done{stream_all ? nullptr : m.stager.get()};
                    auto stage_one = [&](size_t j) -> bool {
                        const int32_t e = order[j];
-                        const bool resident = m.host_res && m.cache && m.host_res[(size_t) l * NE + e] >= 0;
+                        const bool resident = m.host_res && m.cache && m.host_res[(size_t) l * m.g->n_expert + e] >= 0;
                        if (resident) return true;
                        const int sl = stage_next;
                        stage_next = (stage_next + 1) % STAGE;
@@ -1273,7 +1273,7 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                            const int32_t e = order[j];
                            if (stage_of[j] < 0) {
                                ++stats_.experts_resident;
-                                if (!compute(j, m.cache->device_slot(m.host_res[(size_t) l * NE + e]), -1)) return false;
+                                if (!compute(j, m.cache->device_slot(m.host_res[(size_t) l * m.g->n_expert + e]), -1)) return false;
                            } else {
                                pt.mark(kPfWaitCopy, cs);
                                cudaStreamWaitEvent(m.cs, m.copied[stage_of[j]], 0);
@@ -1304,10 +1304,10 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                                issue_until(consumed + (size_t) m.ring);
                            } else {
                                ++stats_.experts_resident;
-                                if (!compute(j, m.cache->device_slot(m.host_res[(size_t) l * NE + e]), -1)) return false;
+                                if (!compute(j, m.cache->device_slot(m.host_res[(size_t) l * m.g->n_expert + e]), -1)) return false;
                            }
                        }
-                        release_to(NE);
+                        release_to(m.g->n_expert);
                    }
                    pt.mark(kPfCombine, cs);
                    moe_combine(m.Dm, m.slot_dev, m.w, m.shared, m.sg, m.bo, T, m.cs);
@@ -53,6 +53,10 @@
 #if defined(_WIN32)
 #include <windows.h>
 #include <psapi.h>
+#include <io.h>
+#else
+#include <unistd.h>
+#include <cerrno>
 #endif

 #include <cuda_runtime.h>
@@ -755,6 +759,16 @@ int main(int argc, char** argv) {
    // pipe or a file is block-buffered, so a program that dies loses every line it had already printed - which
    // turns "it crashed at step 7" into "it crashed somewhere", and the difference is a debugging session.
    std::setvbuf(stdout, nullptr, _IONBF, 0);
+    // Load every CUDA kernel when the context is created, before the expert cache takes the free VRAM.  With the
+    // default lazy loading, a kernel first used mid-prompt (MMQ for IQ3_XXS at 64K+ on a 12 GB card) found no VRAM
+    // left for its code and the engine ended ("out of memory: cudaFuncSetAttribute").  Costs ~30 MB of VRAM.
+    if (std::getenv("CUDA_MODULE_LOADING") == nullptr) {
+#if defined(_WIN32)
+        _putenv_s("CUDA_MODULE_LOADING", "EAGER");
+#else
+        setenv("CUDA_MODULE_LOADING", "EAGER", 0);
+#endif
+    }
    Options o;
    bool have_tokens = false;
    bool have_logits_stride = false;
@@ -1074,7 +1088,9 @@ int main(int argc, char** argv) {
    // the canonical Q2_0 pack's CPU kernels are AVX-512 only; a native pack runs on AVX2 CPUs as well
    if (!native_pack) strata::kernels::cpu::cpu_require_expert_support();
    else if (!strata::kernels::cpu::cpu_avx512_ok())
-        std::fprintf(stderr, "strata generate: this CPU has no AVX-512: the expert kernels run on AVX2\n");
+        std::fprintf(stderr, "strata generate: this CPU has no AVX-512: the expert kernels run on %s "
+                             "(multi-token for the i-quant gate/up rows)\n",
+                     std::getenv("STRATA_NO_IQ256") == nullptr ? "AVX-2" : "ggml-cpu vec_dot (STRATA_NO_IQ256 set)");
    strata::core::NativeEmbed native_embed;
    if (native_pack) {
        if (o.native_preset.empty() || o.spec < 2 || o.keep_canonical ||
@@ -1176,8 +1192,21 @@ int main(int argc, char** argv) {
        strata::kernels::mrope_table_set(d_mrope);
    }
    strata::kernels::ple_set_native_postops(o.native_ple_postops);
-    const strata::core::ModelGeometry g;
-    const int64_t K = 10;
+    strata::core::ModelGeometry g;   // canonical defaults; the model file overrides the MoE shape below
+    int64_t K = 10;
+    if (!o.native_preset.empty()) {
+        // a pruned variant (GSQ-RCO Coder) ships fewer experts than the canonical 512x10; the model file
+        // is the authority on its own MoE shape - everything else in the geometry is unchanged
+        try {
+            strata::GgufFile model_gguf(o.native_preset);
+            if (const strata::MetaValue* v = model_gguf.get("qwen4exp.expert_count")) g.n_expert = (int64_t) v->u;
+            if (const strata::MetaValue* v = model_gguf.get("qwen4exp.expert_used_count")) K = (int64_t) v->u;
+        } catch (const std::exception& e) {
+            std::fprintf(stderr, "strata generate: reading the model's expert shape from %s: %s\n",
+                         o.native_preset.c_str(), e.what());
+            return 1;
+        }
+    }
    // before session_init: every graph captured from here on has the vector's kernels where it applies
    std::string cvec_summary = "0";
    if (!o.cvec_files.empty()) {
@@ -1389,7 +1418,10 @@ int main(int argc, char** argv) {
            o.mtp.clear();
        }
        if (!o.mtp.empty()) mtp.set_prompt_len((int64_t) o.tokens.size());
-        if (!o.mtp.empty() && !mtp.load(o.mtp, g, ss, o.spec, err, o.mtp_window)) { std::fprintf(stderr, "strata generate: %s\n", err.c_str()); return 1; }
+        // the draft layer is the canonical model's MTP head (512 experts) even when the target is pruned,
+        // so it always sees the canonical geometry; `static` because MtpDrafter keeps a reference
+        static const strata::core::ModelGeometry draft_geometry{};
+        if (!o.mtp.empty() && !mtp.load(o.mtp, draft_geometry, ss, o.spec, err, o.mtp_window)) { std::fprintf(stderr, "strata generate: %s\n", err.c_str()); return 1; }
    }
    strata::kernels::cpu::ExpertPool pool(o.pool_workers, /*pin=*/true, /*host_works=*/!o.no_host_worker);
    if (o.no_ple_prefetch) strata::kernels::ple_prefetch_enable(false);
@@ -1405,10 +1437,10 @@ int main(int argc, char** argv) {
            std::fprintf(stderr, "strata generate: %s\n", err.c_str());
            return 1;
        }
-        // The profile knows how many slots it was built for.  `--expert-cache 0` means "take the profile's";
-        // an explicit smaller number is allowed and simply truncates the ranked list, which is the right
-        // behaviour for asking "what would 2,000 slots give" without rebuilding the file.
-        if (o.expert_cache == 0) o.expert_cache = (int) pslots;
+        // An explicit number truncates the ranked list ("what would 2,000 slots give" without rebuilding the
+        // file).  `--expert-cache 0` used to take the count the profile was built for; the profile now ranks
+        // every pair (issue #46: a card that holds more than the old 8,000 used to stop there), so it means auto.
+        if (o.expert_cache == 0) o.expert_cache = -1;
        std::fprintf(stderr, "strata generate: profile %s: %zu ranked pairs, built for %lld slots\n",
                     o.expert_profile.c_str(), profile.size(), (long long) pslots);
    }
@@ -2262,9 +2294,10 @@ int main(int argc, char** argv) {
        return 0;
    };
    if (o.serve) {
-        if (o.spec < 2 || o.mtp.empty() || o.prefill_chunk <= 0 || thits.d_res == nullptr || host_res.empty()) {
-            std::fprintf(stderr, "strata serve: needs --spec T, --mtp DIR, --prefill CHUNK, --expert-profile P and "
-                                 "--expert-cache\n");
+        if (o.spec < 2 || o.mtp.empty() || o.prefill_chunk <= 0 ||
+            (graph_hits && (thits.d_res == nullptr || host_res.empty()))) {
+            std::fprintf(stderr, "strata serve: needs --spec T, --mtp DIR and --prefill CHUNK (and a fillable "
+                                 "--expert-cache; the graphed hit path additionally needs --expert-profile P)\n");
            return 2;
        }
        strata::prefill::Prefill sp;
@@ -2326,8 +2359,11 @@ int main(int argc, char** argv) {
        }
        mem_mark("the verifier and the drafter's binding");
        ver.set_split(o.spec_split);
-        ver.set_pcie_mode(o.pcie_mode == "dma" ? 0 : o.pcie_mode == "direct" ? 1 : o.pcie_mode == "kernel" ? 2
-                          : native_pack ? 0 : 2);   // auto: DMA for the native packs, the copy kernel for Q2_0
+        // auto: the copy kernel for every pack.  DMA (the native packs' default until 0.1.13) has the host call
+        // cudaMemcpyAsync + cudaLaunchHostFunc inside a verify window while the GPU spins on the flag they raise;
+        // issue #31's thread dumps show the host stuck in that cudaMemcpyAsync on a driver lock for good.  The copy
+        // kernel needs no host CUDA call there, and costs ~1-3% decode on IQ3_S (45.3 -> 44.8 tok/s, 8 requests).
+        ver.set_pcie_mode(o.pcie_mode == "dma" ? 0 : o.pcie_mode == "direct" ? 1 : 2);
        std::vector<int64_t> cur;
        // ---- the conversation cache (see ConvCheckpoint).  `live` is what the session holds right now: the tokens
        // it has consumed, so a request that starts with exactly them continues without any copy.  `checks` are the
@@ -2445,8 +2481,35 @@ int main(int argc, char** argv) {
        std::deque<std::string> in_lines;
        bool in_eof = false;
        std::thread([&] {
-            std::string l;
-            while (std::getline(std::cin, l)) {
+            // read(2) on the descriptor, not std::cin: glibc's exit() flushes every stdio stream and waits for
+            // stdin's lock, which getline holds while it waits for input - an engine ending on an error (every
+            // std::exit) would hang in exit() on Linux, and the server would wait for it forever
+            std::string l, buf;
+            char chunk[4096];
+            auto getline_fd = [&](std::string& out) -> bool {
+                for (;;) {
+                    const size_t nlpos = buf.find('\n');
+                    if (nlpos != std::string::npos) {
+                        out.assign(buf, 0, nlpos);
+                        buf.erase(0, nlpos + 1);
+                        return true;
+                    }
+#if defined(_WIN32)
+                    const int n = _read(0, chunk, (unsigned) sizeof chunk);
+#else
+                    const ssize_t n = ::read(0, chunk, sizeof chunk);
+                    if (n < 0 && errno == EINTR) continue;
+#endif
+                    if (n <= 0) {
+                        if (buf.empty()) return false;
+                        out.swap(buf);
+                        buf.clear();
+                        return true;
+                    }
+                    buf.append(chunk, (size_t) n);
+                }
+            };
+            while (getline_fd(l)) {
                if (!l.empty() && l.back() == '\r') l.pop_back();
                if (l == "STOP") { stop_req.store(true); continue; }
                std::lock_guard<std::mutex> lk(in_mu);
@@ -3510,8 +3573,11 @@ int main(int argc, char** argv) {
        }
        mem_mark("the verifier and the drafter's binding");
        ver.set_split(o.spec_split);
-        ver.set_pcie_mode(o.pcie_mode == "dma" ? 0 : o.pcie_mode == "direct" ? 1 : o.pcie_mode == "kernel" ? 2
-                          : native_pack ? 0 : 2);   // auto: DMA for the native packs, the copy kernel for Q2_0
+        // auto: the copy kernel for every pack.  DMA (the native packs' default until 0.1.13) has the host call
+        // cudaMemcpyAsync + cudaLaunchHostFunc inside a verify window while the GPU spins on the flag they raise;
+        // issue #31's thread dumps show the host stuck in that cudaMemcpyAsync on a driver lock for good.  The copy
+        // kernel needs no host CUDA call there, and costs ~1-3% decode on IQ3_S (45.3 -> 44.8 tok/s, 8 requests).
+        ver.set_pcie_mode(o.pcie_mode == "dma" ? 0 : o.pcie_mode == "direct" ? 1 : 2);
        drive.d.plan = ver.plan_sink();
        drive.d.pcie_num = (int) (o.pcie_frac * 256.0 + 0.5);
        if (drive.d.pcie_num < 0) drive.d.pcie_num = 0;
@@ -227,20 +227,24 @@ def main() -> int:

    # ---- the experts
    n_layers = 1 + max(int(n.split(".")[1]) for n in T if n.startswith("blk.") and n.endswith("_exps.weight"))
+    n_expert = int(T["blk.0.ffn_gate_inp.weight"].shape[1])   # router rows = experts kept (pruned models ship < 512)
+    if any(int(T["blk.%d.ffn_gate_inp.weight" % l].shape[1]) != n_expert for l in range(n_layers)):
+        print("the routers disagree on the expert count; a per-layer pruned model cannot be packed")
+        return 1
    layout, offset = [], 0
    for l in range(n_layers):
        ts = [T["blk.%d.ffn_%s_exps.weight" % (l, r)] for r in ROLES]
-        per = [t.expected_bytes() // N_EXPERT for t in ts]
+        per = [t.expected_bytes() // n_expert for t in ts]
        if per[0] != per[1] or ts[0].type_name != ts[1].type_name:
            print("layer %d: gate and up differ in type" % l)
            return 1
        blob = per[0] + per[1] + per[2]
        layout.append((l, ts[0].type_id, ts[2].type_id, offset, blob, ts))
-        offset += blob * N_EXPERT
+        offset += blob * n_expert
    with open(out / "native_experts.txt", "w", encoding="utf-8", newline="\n") as fo:
        fo.write("# strata native experts v3: layer gu_type d_type offset blob_bytes gate_off up_off down_off [shard] "
                 "(n_expert %d, total %d; absolute offsets in %s, or in the named shard beside it)\n"
-                 % (N_EXPERT, offset, src.name))
+                 % (n_expert, offset, src.name))
        for l, gt, dt, off, blob, ts in layout:
            ws = [model.where[t.name] for t in ts]
            if len({w[3] for w in ws}) != 1:
@@ -259,9 +263,9 @@ def main() -> int:
        return 0
    with open(path, "wb") as fo:
        for l, gt, dt, off, blob, ts in layout:
-            parts = [model.bytes(t.name).reshape(N_EXPERT, -1) for t in ts]
-            chunk = np.concatenate(parts, axis=1)          # (512, blob): gate | up | down per expert
-            assert chunk.shape == (N_EXPERT, blob)
+            parts = [model.bytes(t.name).reshape(n_expert, -1) for t in ts]
+            chunk = np.concatenate(parts, axis=1)          # (n_expert, blob): gate | up | down per expert
+            assert chunk.shape == (n_expert, blob)
            fo.write(chunk.tobytes())
            if l % 8 == 0:
                print("  layer %2d  %-8s/%-7s blob %8d  at %.2f GiB" % (l, ts[0].type_name, ts[2].type_name, blob,
@@ -0,0 +1,99 @@
+"""tools/make_profile.py - the expert-cache profile (`data/expert-profile.bin`): every (layer, expert) pair, ranked.
+
+`--expert-cache auto` fills the VRAM slots it can afford with the profile's pairs in order, so the ranking decides
+which experts start resident (the adaptive tier then swaps in what a conversation routes most).  A profile that
+ranks fewer pairs than a card can hold caps the cache (issue #46: a 32 GB card stopped at 8,000 slots); this tool
+writes one that ranks all 24,576.
+
+The order: the base profile's ranking (default: the shipped data/expert-profile.bin), then the pairs your routing
+traces used, most frequent first, then every pair still missing, interleaved across the layers.
+
+    python tools/make_profile.py [TRACE ...] [--base data/expert-profile.bin | --no-base] [--out PATH]
+                                 [--n-expert 256]      (a pruned model: GSQ-RCO Coder keeps 256 of 512)
+
+A routing trace comes from a one-shot engine run with `--dump-routing FILE` (a prompt typical of your use; the
+routed experts of every layer and position are written).  Point the model config's `--expert-profile` at the result.
+"""
+import argparse
+import struct
+from collections import defaultdict
+from pathlib import Path
+
+ROOT = Path(__file__).resolve().parent.parent
+N_LAYER, N_EXPERT = 48, 512
+MAGIC, VERSION = b"STRP", 1
+
+
+def read_profile(path, n_expert=N_EXPERT):
+    blob = Path(path).read_bytes()
+    if blob[:4] != MAGIC:
+        raise SystemExit(f"{path}: not a Strata profile")
+    ver, nl, ne, slots, n = struct.unpack_from("<5I", blob, 4)
+    if (nl, ne) != (N_LAYER, n_expert):
+        raise SystemExit(f"{path}: {nl}x{ne}, not {N_LAYER}x{n_expert}")
+    return [struct.unpack_from("<HH", blob, 24 + 4 * i) for i in range(n)]
+
+
+def read_trace(path, n_expert=N_EXPERT):
+    """(layer, k, k expert ids, k weights) records, as `--dump-routing` writes them."""
+    blob = Path(path).read_bytes()
+    off, freq = 0, defaultdict(int)
+    while off + 8 <= len(blob):
+        layer, k = struct.unpack_from("<ii", blob, off)
+        off += 8
+        for e in struct.unpack_from("<%di" % k, blob, off):
+            if 0 <= layer < N_LAYER and 0 <= e < n_expert:
+                freq[(layer, e)] += 1
+        off += 8 * k                                    # the ids and the weights
+    return freq
+
+
+def write_profile(path, ranked, n_expert=N_EXPERT):
+    table = [[-1] * n_expert for _ in range(N_LAYER)]
+    for slot, (layer, e) in enumerate(ranked):
+        table[layer][e] = slot
+    with open(path, "wb") as f:
+        f.write(MAGIC + struct.pack("<5I", VERSION, N_LAYER, n_expert, len(ranked), len(ranked)))
+        for layer, e in ranked:
+            f.write(struct.pack("<HH", layer, e))
+        for layer in range(N_LAYER):
+            f.write(struct.pack("<%di" % n_expert, *table[layer]))
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__.split("\n\n")[0])
+    ap.add_argument("traces", nargs="*", help="routing traces from --dump-routing")
+    ap.add_argument("--base", default=str(ROOT / "data" / "expert-profile.bin"), help="ranking to keep first")
+    ap.add_argument("--no-base", action="store_true", help="rank by the traces only")
+    ap.add_argument("--out", default=str(ROOT / "data" / "expert-profile.bin"))
+    ap.add_argument("--n-expert", type=int, default=N_EXPERT, help="experts per layer (default 512)")
+    a = ap.parse_args()
+
+    ranked, seen = [], set()
+
+    def take(pairs):
+        for p in pairs:
+            p = (int(p[0]), int(p[1]))
+            if p not in seen:
+                seen.add(p)
+                ranked.append(p)
+
+    ne = a.n_expert
+    if not a.no_base:
+        take(read_profile(a.base, ne))
+    n_base = len(ranked)
+    freq = defaultdict(int)
+    for t in a.traces:
+        for p, c in read_trace(t, ne).items():
+            freq[p] += c
+    take(p for p, _ in sorted(freq.items(), key=lambda kv: (-kv[1], kv[0])))
+    n_trace = len(ranked) - n_base
+    take((layer, e) for e in range(ne) for layer in range(N_LAYER))   # the rest, across the layers
+    write_profile(a.out, ranked, ne)
+    assert read_profile(a.out, ne) == ranked, "the profile did not survive the round trip"
+    print(f"wrote {a.out}: {len(ranked)} ranked pairs ({n_base} from the base, {n_trace} from the traces, "
+          f"{len(ranked) - n_base - n_trace} filled in)")
+
+
+if __name__ == "__main__":
+    main()
Author	SHA1	Message	Date
Niko1221	b500a81e4e	License: MIT (the ggml code, the font and the projection vector keep their own licenses)	2026-09-28 14:29:00 +02:00
Niko1221	26cdd4cde7	Engine 0.1.16; setup requires it (the Coder needs the model's own expert count, PR #54 )	2026-09-28 14:27:45 +02:00
Niko1221	47a83696e9	The GSQ-RCO Coder (256 of 512 experts): setup, profile, speeds (PR #54 ) - setup: a third family, "coder" (IQ1_M), with its own expert profile; its shard 2 and vision encoder are the original's files and are reused; 32 GB of RAM is now enough to install (the Coder's experts are 23 GB). - data/expert-profile-coder.bin: the shipped ranking mapped onto the kept experts through the release's rco-allocation.txt (72% GPU hits, 12 GB card); tools/make_profile.py takes --n-expert. - README/DETAILS: the Coder's speeds (prompts 599-1,350 tok/s, output 43-53 tok/s, 1K-262K) and what it is; bench/results/2026-09-28-coder. - The experimental speed projection can be turned on for the Coder too.	2026-09-28 14:20:58 +02:00
Niko1221	ef206bb671	Merge PR #54 (pjgmobile): pruned expert variants (GSQ-RCO Coder: 256 of 512 experts)	2026-09-28 13:45:47 +02:00
Niko1221	71c60dafce	setup: model files in a data folder next to Strata (no re-download after an update); freeze warnings - The model files (models/, packs/, mtp/: 70-120 GB) live in Strata-data next to the Strata folder, remembered per user (%APPDATA%\Strata\settings.json, ~/.config/strata/settings.json). Updating by unzipping a new copy used to give an empty folder and a full download again. Model files of this folder and of earlier Strata folders on the same drive are moved there (a rename) and their configs repointed; files on another drive are used where they are. A new copy with no config sets itself up like the last install, reusing the files. --data-dir chooses another place. - The start says, before and while it happens, that the PC can be slow or stop responding for 1-3 minutes while the experts are loaded and locked: setup's start box, and the server narrates the engine's start from its log (weights, experts + the warning, expert cache, ready, a "still starting" line every 20 s).	2026-09-28 13:45:39 +02:00
pete	97ae3b2418	Merge remote-tracking branch 'origin/main' into pruned-geometry-support # Conflicts: # src/prefill/prefill.cpp	2026-09-28 12:34:36 +01:00
pete	8bef4161cf	Support pruned expert variants (GSQ-RCO Coder: 256 of 512 experts) - iq_pack.py: derive n_expert from the router tensor shape (was hardcoded 512) - expert_layout.cpp: parse the pack header's recorded n_expert instead of the caller's compile-time canonical geometry; use it for the contiguity stride - gguf_reader.hpp: Qwen4ExpGuard expert-count checks become presence-only - generate.cpp: read qwen4exp.expert_count/expert_used_count from the model file into the runtime geometry; MTP drafter keeps canonical geometry (the draft is the original 512-expert head); serve gate allows profile-less caches; K from metadata - layer.cpp: native fused router stays 512x10-only, generic top-k fallback - prefill: expert count taken from the geometry everywhere; route_kernel templatized (512->16, 256->8 registers/lane, else generic router_top10)	2026-09-28 12:29:32 +01:00
Niko1221	256485621e	docs: troubleshooting row for the IQ3_XXS long-prompt stop (fixed in 0.1.15)	2026-09-28 13:18:18 +02:00
Niko1221	9ac549fd32	Engine 0.1.15; setup requires it (kernels loaded at start: IQ3_XXS at 64K+ no longer runs out of VRAM mid-prompt)	2026-09-28 13:17:58 +02:00
Niko1221	06f41f49a8	docs: speed tables measured with engine 0.1.14 (prompts 1.2-2.4x faster than the old tables)	2026-09-28 13:07:35 +02:00
Niko1221	bef32ed6c3	engine: load every CUDA kernel at start (CUDA_MODULE_LOADING=EAGER) With lazy loading, a kernel first used mid-prompt found no VRAM left for its code once the expert cache and the prompt path had taken it: IQ3_XXS at 64K and 128K on a 12 GB card ended with 'out of memory: cudaFuncSetAttribute' in the MMQ launch. Loaded at start the kernels cost ~30 MB, accounted before the cache is sized (~23 fewer slots). A user's own CUDA_MODULE_LOADING is kept.	2026-09-28 12:49:25 +02:00
Niko1221	dcb0c24422	docs: Linux updates, other distributions and WSL; the 0.1.14 stall fix; the full expert profile	2026-09-28 11:46:37 +02:00
Niko1221	235cf4bd09	serve: the prompt progress line says how far the engine is (N of M tokens), not only the prompt's size (#29 )	2026-09-28 11:22:55 +02:00
Niko1221	b38d11761a	Engine 0.1.14; setup requires it (the PCIe share without host CUDA calls in a verify window, #31 )	2026-09-28 11:21:01 +02:00
Niko1221	fbbfb9ce60	Merge PR #43 (pipeob0): AVX2 multi-token i-quant row kernels	2026-09-28 11:04:17 +02:00
Niko1221	5c52c0fd00	Merge PR #42 (pipeob0): MEM_LARGE_PAGES on Windows - enable the privilege, round the size	2026-09-28 11:02:44 +02:00
Niko1221	9253c92ad7	The PCIe share by the copy kernel for every pack (#31 , #29 ); setup and profile fixes (#46 , #49 , #45 ) - --pcie-mode auto is the copy kernel for native packs too. Issue #31's thread dumps show the host stuck for good in the DMA path's cudaMemcpyAsync (a driver lock) inside a verify window while the GPU spins on the flag those copies raise. The copy kernel needs no host CUDA call there; IQ3_S decode 45.3 -> 44.8 tok/s. --pcie-mode dma keeps the old path. - setup: a locally compiled engine records its version; one kept after a failed compile is recognised by the version in its binary, and a config it cannot read (--prefill auto on 0.1.12) goes back to 2048 (#49). - setup: nvcc in /opt/cuda (Arch), a clearer message off Ubuntu/Debian (#46); a failed build is tried once more (#45: ptxas's intermittent PTX parse error). - data/expert-profile.bin ranks all 24,576 experts (the first 8,000 unchanged): the cache stopped at 8,000 slots on 24-32 GB cards (#46). tools/make_profile.py builds or extends a profile from --dump-routing traces; --expert-cache 0 with a profile now means auto.	2026-09-28 11:02:44 +02:00
pipeob0	d7a66bd33e	generate: say which expert kernels an AVX2-only CPU gets The notice printed 'the expert kernels run on AVX2' whether or not the AVX2 i-quant row kernels were used, and with STRATA_NO_IQ256 set they are not: the run is on ggml's vec_dot. Name the kernel, and name the switch that changed it.	2026-09-28 09:51:47 +02:00
pipeob0	d1b30db332	pinned: say how many bytes were refused, and stop blaming the privilege The refusal note blamed SeLockMemoryPrivilege whatever VirtualAlloc answered, and did not say what was asked for: 1450 (ERROR_NO_SYSTEM_RESOURCES, the large-page pool has that many 2 MB pages free right now), 87 (a size that is not a multiple of the minimum) and 1314 (the privilege is genuinely off) all read as the same one line. Observed for real on a 5700X3D box: a run minutes after an engine was killed printed 1450 while the next one succeeded - the killed process had not released its pinned large pages yet, which is only visible if the byte count and the error are both in the log.	2026-09-28 09:51:47 +02:00
pipeob0	0015035642	parity tool: guard the AVX-512 calls, and check the AVX2 kernels too The tool called the AVX-512 row kernels and the AVX-512 quantizer unconditionally, so on an AVX2-only CPU it died with SIGILL before printing anything - the one tool you want on that machine. Both calls are now behind cpu_avx512_ok(). Adds a block that runs the new AVX2 kernels directly against ggml's vec_dot on Q8_0 activations (the IQ4_NL down rows' real input) and times them, so a kernel change can be validated and measured without the engine.	2026-09-28 09:51:47 +02:00
pipeob0	89c62e5361	kernels: AVX2 multi-token i-quant kernels for an AVX2-only CPU The native pack's i-quant gate/up rows fell back to ggml's single-token vec_dot on a CPU without AVX-512 (Ryzen 5000 series, Ice Lake, ...): the weights were decoded once per token, and the 5700X3D's pool was decode bound. iq_avx2.cpp adds multi-token row kernels for IQ2_XXS/IQ2_XS/IQ3_XXS/IQ3_S/IQ2_S and IQ4_NL down rows, and for IQ2_XS the vectorized sign decode of ggml's AVX2 vec_dot (the 8th sign bit by parity through the bit_helper pshufb, the sign vectors out of shuffles of one 32-byte load, the 8 scale bytes through the unpack trick), with an even_signs table equivalent to keven_signs_q2xs. A row is decoded once and every token in the verify window only pays a load, a sign, a maddubs, a madd and an add. Measured on a 5700X3D + RTX 5060 Ti, Swift native pack, greedy, 300 tokens, medians of interleaved A/B pairs: gate/up 15.16 -> 13.84 ms/round, rows 22.7 -> 23.8 GB/s, 48.5 ms/round vs 50.9 (-4.7%). Single thread, one expert's gate+up: IQ2_XS 3 tokens 536 -> 416 us, IQ2_XXS 390 -> 337 us. IQ4_NL down rows 248 -> 174 us for 3 tokens. native_expert_parity: rel 3.1e-08 vs ggml's vec_dot, 0 failures. STRATA_NO_IQ256 and STRATA_NO_IQ4NL fall back to ggml per format.	2026-09-28 09:51:47 +02:00
pipeob0	97cf327c1e	pinned: enable SeLockMemoryPrivilege and round the arena to GetLargePageMinimum reserve() asked for MEM_LARGE_PAGES without ever enabling SeLockMemoryPrivilege in the process token, and passed the arena size unrounded. VirtualAlloc requires an exact multiple of GetLargePageMinimum, so the call failed with 87 (ERROR_INVALID_PARAMETER) and the log blamed the privilege. Both fixed; STRATA_NO_LARGEPAGES skips the attempt and the note now says what actually happened.	2026-09-28 09:51:46 +02:00
Niko1221	6c06571710	Linux/WSL: engine errors end the engine and reach the client - The --serve stdin reader reads the descriptor instead of std::cin: glibc's exit() waits for stdin's lock, which getline held while waiting for input, so an engine ending on any error hung in exit() and the server waited forever. - session_init fails when a QSA layer could not pin its KV streaming copy; it used to continue half-built and hit an illegal memory access at the first request (seen under WSL, where the driver pins only about 1 GiB). The message says how much was pinned and what to change. - setup: no KV streaming under WSL (new installs and existing configs). - server: an engine ERR after an SSE stream started is sent as an error event (it was a 400 response written into the open stream: an empty answer), and the request is recorded as "error".	2026-09-28 08:51:22 +02:00
Niko1221	fb9b924e5c	verify.cpp: include <cmath> (engine 0.1.13 did not compile on Linux: std::isfinite) MSVC brought <cmath> in through another header; GCC does not. Found by building in WSL (Ubuntu 24.04, CUDA 13.0) through setup.sh the way a Linux user does.	2026-09-28 08:03:22 +02:00
Niko1221	ffcdb437b8	Linux setup fixes, one pool worker per physical core on Linux, --port for installs, Monitor totals - setup.sh is executable in the repository (a fresh clone could not run `./setup.sh`: Permission denied). - setup.sh: Debian/Ubuntu ship Python's venv without ensurepip (the separate python3-venv package); the check now needs both, so it installs python3-venv instead of failing to create .venv - and a .venv left without pip by a failed run is made again. - Linux: physical_cores() keeps one allowed CPU per (package, core) from sysfs, like the Windows branch; SMT siblings no longer get their own pool worker (issue #40: 23 -> 15 workers on a 13700K, +16% decode). - START-HERE / setup.sh `--port N` now applies to an existing install; without it the install's own port is kept (issue #32). - Monitor: the server keeps the last 500 requests and totals since it started; the Monitor shows the totals and a "Show all" toggle (`/metrics?requests=all`), 12 rows otherwise (issue #35).	2026-09-28 07:55:39 +02:00
pete	898d6a4c3e	Merge remote-tracking branch 'origin/main' # Conflicts: # serve/server.py	2026-09-27 17:10:41 +01:00
pete	ba8c55e1cc	serve: optional --fit-max-tokens clamps the output cap to remaining context instead of rejecting Agentic clients legitimately request the model's full output cap with a long prompt (e.g. prompt 131,888 + max_tokens 130,298 on a 262,144 window: rejected over a 42-token overshoot). Default behaviour is unchanged: without the flag the server still rejects with 400, like llama.cpp. With --fit-max-tokens the server clamps max_tokens to the remaining room (max_context - CTX_SLACK - prompt, floor 1 token) and answers. Measured on a 262,144-token host: prompt 251,455 + max_tokens 131,072 -> HTTP 200, completion_tokens 10,681 (= 262,144 - 251,455 - 8), finish length. Flag off: byte-identical 400. Small requests unaffected in both modes.	2026-09-27 13:24:23 +01:00
pete	b78d9da751	Merge remote-tracking branch 'origin/main'	2026-09-27 13:24:23 +01:00
pete	ad29df2396	Merge branch 'pr-14'	2026-09-26 19:20:59 +01:00
pete	8936605cc5	Merge branch 'pr-8'	2026-09-26 17:39:42 +01:00
pete	3503b5ddd8	Merge branch 'pr-7'	2026-09-26 17:39:42 +01:00