Engine 0.1.28: the expert cache reserves the draft head's VRAM (0.1.27 left 16 GB cards below the stall line), a cancelled request no longer fails the next one, the API key protects /status; setup requires it

docs: the draft layer's token subset and --draft-vocab
setup: --host/--api-key/--draft-vocab on a start are kept for the model (#179 ); --draft-vocab en keeps the English/code subset (~110 MiB less VRAM); CUDA 13.0 for RTX 50 builds (#220 ); llama.cpp's web UI is not unpacked (Windows path limit, #206 ); Swift Q2_0 hidden (#171 ); the earlier data folder and nested Strata-data move into --data-dir (#198 ); whole shards without a finish mark count (#173 )
2026-11-07 14:21:06 -10:00 · 2026-09-30 11:18:54 +02:00 · 2026-09-30 11:18:54 +02:00 · 2026-09-30 10:54:45 +02:00 · 2026-09-30 10:54:45 +02:00 · 2026-09-30 10:54:45 +02:00
23 changed files with 1251 additions and 119 deletions
@@ -8,7 +8,7 @@
 # below refuses anything older: silently building for a pre-Ampere 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.25 LANGUAGES CXX)
+project(strata VERSION 0.1.28 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
@@ -62,14 +62,15 @@ if(STRATA_ENABLE_CUDA)
  endif()
  foreach(_a IN LISTS CMAKE_CUDA_ARCHITECTURES)
    string(REGEX REPLACE "-.*$" "" _base "${_a}")     # 120-real / 120-virtual are still 120
-    # Plan v0.3 P6: developed and measured on sm_120 (RTX 50); the kernels need sm_80 or newer (tf32 mma in
-    # the attention scorer, bf16 math), so RTX 30 (86) and RTX 40 (89) build too.  Older cards are refused.
+    # Turing port: developed and measured on sm_120 (RTX 50); the tf32 mma in the QSA scorer has a
+    # portable fp32-FMA fallback, so RTX 20 (75), 30 (86), RTX 40 (89) and RTX 50 (120) build too.
+    # Pre-sm_75 is refused.
    if(_base STREQUAL "native" OR _base STREQUAL "all" OR _base STREQUAL "all-major")
      continue()
    endif()
-    if(_base LESS 80 AND NOT STRATA_EXPERIMENTAL_SM75)   # the A/B of a Turing card as a layer-split stage
+    if(_base LESS 75)
      message(FATAL_ERROR
-        "Strata needs an NVIDIA GPU of compute capability 8.0 or newer (RTX 30 / 40 / 50 series), but "
+        "Strata needs an NVIDIA GPU of compute capability 7.5 or newer (RTX 20 / 30 / 40 / 50 series), but "
        "CMAKE_CUDA_ARCHITECTURES is '${CMAKE_CUDA_ARCHITECTURES}'.")
    endif()
  endforeach()
@@ -664,7 +665,7 @@ if(STRATA_NATIVE_EXPERTS)
    endif()
    set(_strata_mmq ${_strata_gg}/src/ggml-cuda)
    add_library(strata_mmq STATIC src/prefill/moe_mmq.cu src/prefill/ggml_cuda_host.cu ${_strata_mmq}/quantize.cu)
-    foreach(_t q2_0 iq2_xxs iq2_xs iq2_s iq3_xxs iq3_s iq4_nl iq4_xs)
+    foreach(_t q2_0 iq2_xxs iq2_xs iq2_s iq3_xxs iq3_s iq4_nl iq4_xs q8_0)   # q8_0: the drafter's matrices (E-9)
      target_sources(strata_mmq PRIVATE ${_strata_mmq}/template-instances/mmq-instance-${_t}.cu)
    endforeach()
    target_include_directories(strata_mmq PRIVATE ${_strata_gg}/include ${_strata_gg}/src ${_strata_mmq}
@@ -692,7 +693,7 @@ if(STRATA_NATIVE_EXPERTS)
      src/prefill/moe_mmq.cu
      src/prefill/ggml_cuda_host.cu
      "${_strata_mmq}/quantize.cu")
-    foreach(_t q2_0 iq2_xxs iq2_xs iq2_s iq3_xxs iq3_s iq4_nl iq4_xs)
+    foreach(_t q2_0 iq2_xxs iq2_xs iq2_s iq3_xxs iq3_s iq4_nl iq4_xs q8_0)
      list(APPEND _strata_mmq_sources "${_strata_mmq}/template-instances/mmq-instance-${_t}.cu")
    endforeach()
    set_source_files_properties(${_strata_mmq_sources} PROPERTIES LANGUAGE HIP)
@@ -25,15 +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** | 90 tokens/s | 67 tokens/s | 1,840 tokens/s |
-| **IQ2_XS** | 74 tokens/s | 60 tokens/s | 1,800 tokens/s |
-| **IQ3_XXS** | 62 tokens/s | 46 tokens/s | 1,550 tokens/s |
-| **IQ3_S** | 52 tokens/s | 41 tokens/s | 1,500 tokens/s |
-| **Coder** (IQ1_M) | 51 tokens/s | 44 tokens/s | 1,870 tokens/s |
+| **Q2_0** | 93 tokens/s | 74 tokens/s | 2,170 tokens/s |
+| **IQ2_XS** | 79 tokens/s | 63 tokens/s | 2,090 tokens/s |
+| **IQ3_XXS** | 62 tokens/s | 49 tokens/s | 1,750 tokens/s |
+| **IQ3_S** | 53 tokens/s | 46 tokens/s | 1,620 tokens/s |
+| **Coder** (IQ1_M) | 55 tokens/s | 43 tokens/s | 2,180 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), measured on a
-  32K-token prompt; a 4K prompt reads at 890-1,520 tokens/s. A 32K prompt takes about 17 seconds with Q2_0.
+  32K-token prompt; a 4K prompt reads at 910-1,580 tokens/s. A 32K prompt takes about 15 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
@@ -47,7 +47,7 @@ asks whether to share the model across them (recommended when two can). An insta
 its next start. Or choose yourself: `START-HERE.bat --gpus 0,2` (both, remembered), `--gpus all`, or `--gpu 0` (one
 card, this start only). Each card keeps the experts of its own layers, and prompts flow through the cards in a
 pipeline: on an RTX 5080 + RTX 3090 prompts were read 18-20% faster than on the 5080 alone, decoding on par.
-Every card must be an RTX 30 series or newer with 8 GB or more. See [docs/MULTI_GPU.md](docs/MULTI_GPU.md).
+Every card must be an RTX 20 series or newer with 8 GB or more. See [docs/MULTI_GPU.md](docs/MULTI_GPU.md).

 ## Which model should I pick?

@@ -90,7 +90,8 @@ yet). Details: [AMD HIP](docs/AMD_HIP.md).

 ## Install

-**You need:** an NVIDIA RTX 30, 40 or 50 card with 12 GB of VRAM or more, enough RAM for the size you pick (above),
+**You need:** an NVIDIA RTX 20, 30, 40 or 50 card with 12 GB of VRAM or more (RTX 20 since 0.1.27), enough RAM for the size you pick (above;
+a big GPU makes up for less RAM - the [low-RAM mode](docs/DETAILS.md)),
 ~80 GB of free disk space (an SSD makes the first start much faster), and Windows 10/11 or Linux. The only thing you
 install yourself is a current **NVIDIA driver** ([nvidia.com/drivers](https://www.nvidia.com/drivers) or the NVIDIA
 App). Everything else - Python, the engine, the model - is set up for you.
@@ -0,0 +1,36 @@
+# Speed with engine 0.1.26
+
+RTX 5070 12 GB (PCIe 5.0 x16), Ryzen 5 7600, 64 GB DDR5-5200, Windows 10. The ready-made 0.1.26 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, images
+off, no timing marks. The same code-agent prompts as `2026-09-29-speed-0122`. Every run: [`matrix.json`](matrix.json).
+
+## Prompt (tokens/s)
+
+| Model | 1K | 4K | 32K | 64K | 128K |
+| --- | ---: | ---: | ---: | ---: | ---: |
+| **Q2_0** | 536 | 1,299 | 2,171 | 2,126 | 2,107 |
+| **IQ2_XS** | 534 | 1,256 | 2,092 | 1,754 | 1,752 |
+| **IQ3_XXS** | 482 | 1,007 | 1,745 | 1,609 | 1,602 |
+| **IQ3_S** | 427 | 913 | 1,624 | 1,640 | 1,443 |
+| **Coder** | 656 | 1,583 | 2,177 | 2,236 | 2,208 |
+
+## Output (tokens/s)
+
+| Model | 1K | 4K | 32K | 64K | 128K |
+| --- | ---: | ---: | ---: | ---: | ---: |
+| **Q2_0** | 87.3 | 93.0 | 81.8 | 76.2 | 73.7 |
+| **IQ2_XS** | 79.6 | 78.6 | 76.3 | 63.7 | 62.7 |
+| **IQ3_XXS** | 61.9 | 61.6 | 58.5 | 57.2 | 49.0 |
+| **IQ3_S** | 52.4 | 53.3 | 48.3 | 46.3 | 45.5 |
+| **Coder** | 58.9 | 55.1 | 54.9 | 53.2 | 43.0 |
+
+- Against 0.1.22 (`2026-09-29-speed-0122`), prompts at 32K-128K are 8-28% faster: Q2_0 1,844 -> 2,171 at 32K and
+  1,682 -> 2,107 at 128K, the Coder 1,779 -> 2,208 at 128K. The gains come from:
+  - 0.1.24's QSA selection on tensor cores;
+  - 0.1.25's grouping tables in mapped memory and fused norms;
+  - 0.1.26's batched draft-layer pass.
+- 262K was not run again (0.1.22: Q2_0 1,304, IQ2_XS 1,181 tokens/s).
+- Output speed is the same decode as before. One run per cell moves several percent with the text (the share of
+  accepted drafts, `spec_accept` in `matrix.json`). For example, the Coder's 128K answer accepted 55% of its drafts
+  and its 64K answer 71%.
@@ -0,0 +1,352 @@
+[
+ {
+  "model": "Q2_0",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 23.9,
+  "decode_tok_s": 87.33,
+  "decoded": 256,
+  "prefill_tok_s": 536.42,
+  "ttft_s": 3.4,
+  "spec_accept": 0.681,
+  "tokens_per_round": 2.72,
+  "vram_slots": 4174
+ },
+ {
+  "model": "Q2_0",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 27.6,
+  "decode_tok_s": 92.97,
+  "decoded": 256,
+  "prefill_tok_s": 1299.44,
+  "ttft_s": 4.21,
+  "spec_accept": 0.838,
+  "tokens_per_round": 3.23,
+  "vram_slots": 4121
+ },
+ {
+  "model": "Q2_0",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 40.8,
+  "decode_tok_s": 81.75,
+  "decoded": 256,
+  "prefill_tok_s": 2170.85,
+  "ttft_s": 16.02,
+  "spec_accept": 0.788,
+  "tokens_per_round": 2.99,
+  "vram_slots": 3832
+ },
+ {
+  "model": "Q2_0",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 57.3,
+  "decode_tok_s": 76.23,
+  "decoded": 256,
+  "prefill_tok_s": 2126.03,
+  "ttft_s": 32.06,
+  "spec_accept": 0.707,
+  "tokens_per_round": 2.8,
+  "vram_slots": 3778
+ },
+ {
+  "model": "Q2_0",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 87.7,
+  "decode_tok_s": 73.73,
+  "decoded": 256,
+  "prefill_tok_s": 2106.85,
+  "ttft_s": 62.48,
+  "spec_accept": 0.679,
+  "tokens_per_round": 2.71,
+  "vram_slots": 3695
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 42.2,
+  "decode_tok_s": 79.62,
+  "decoded": 256,
+  "prefill_tok_s": 533.83,
+  "ttft_s": 3.51,
+  "spec_accept": 0.774,
+  "tokens_per_round": 2.89,
+  "vram_slots": 4359
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 40.6,
+  "decode_tok_s": 78.6,
+  "decoded": 256,
+  "prefill_tok_s": 1256.02,
+  "ttft_s": 4.31,
+  "spec_accept": 0.81,
+  "tokens_per_round": 3.29,
+  "vram_slots": 4308
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 53.6,
+  "decode_tok_s": 76.34,
+  "decoded": 256,
+  "prefill_tok_s": 2092.02,
+  "ttft_s": 16.57,
+  "spec_accept": 0.85,
+  "tokens_per_round": 3.2,
+  "vram_slots": 4029
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 78.3,
+  "decode_tok_s": 63.69,
+  "decoded": 256,
+  "prefill_tok_s": 1753.65,
+  "ttft_s": 38.08,
+  "spec_accept": 0.636,
+  "tokens_per_round": 2.59,
+  "vram_slots": 3977
+ },
+ {
+  "model": "IQ2_XS",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 115.1,
+  "decode_tok_s": 62.72,
+  "decoded": 256,
+  "prefill_tok_s": 1751.79,
+  "ttft_s": 74.85,
+  "spec_accept": 0.696,
+  "tokens_per_round": 2.6,
+  "vram_slots": 3896
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 46.9,
+  "decode_tok_s": 61.94,
+  "decoded": 256,
+  "prefill_tok_s": 482.04,
+  "ttft_s": 3.67,
+  "spec_accept": 0.845,
+  "tokens_per_round": 3.57,
+  "vram_slots": 3279
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 47.4,
+  "decode_tok_s": 61.64,
+  "decoded": 256,
+  "prefill_tok_s": 1007.48,
+  "ttft_s": 5.03,
+  "spec_accept": 0.801,
+  "tokens_per_round": 3.19,
+  "vram_slots": 3242
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 65.6,
+  "decode_tok_s": 58.49,
+  "decoded": 256,
+  "prefill_tok_s": 1745.16,
+  "ttft_s": 19.63,
+  "spec_accept": 0.789,
+  "tokens_per_round": 2.85,
+  "vram_slots": 3007
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 88.5,
+  "decode_tok_s": 57.21,
+  "decoded": 256,
+  "prefill_tok_s": 1609.31,
+  "ttft_s": 41.45,
+  "spec_accept": 0.74,
+  "tokens_per_round": 2.87,
+  "vram_slots": 2963
+ },
+ {
+  "model": "IQ3_XXS",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 127.0,
+  "decode_tok_s": 48.95,
+  "decoded": 256,
+  "prefill_tok_s": 1602.33,
+  "ttft_s": 81.71,
+  "spec_accept": 0.643,
+  "tokens_per_round": 2.37,
+  "vram_slots": 2899
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 54.2,
+  "decode_tok_s": 52.4,
+  "decoded": 256,
+  "prefill_tok_s": 426.6,
+  "ttft_s": 4.03,
+  "spec_accept": 0.822,
+  "tokens_per_round": 3.17,
+  "vram_slots": 2617
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 58.9,
+  "decode_tok_s": 53.3,
+  "decoded": 256,
+  "prefill_tok_s": 912.99,
+  "ttft_s": 5.58,
+  "spec_accept": 0.85,
+  "tokens_per_round": 3.38,
+  "vram_slots": 2579
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 72.8,
+  "decode_tok_s": 48.31,
+  "decoded": 256,
+  "prefill_tok_s": 1624.07,
+  "ttft_s": 21.07,
+  "spec_accept": 0.827,
+  "tokens_per_round": 3.0,
+  "vram_slots": 2383
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 93.3,
+  "decode_tok_s": 46.25,
+  "decoded": 256,
+  "prefill_tok_s": 1640.34,
+  "ttft_s": 40.77,
+  "spec_accept": 0.693,
+  "tokens_per_round": 2.56,
+  "vram_slots": 2346
+ },
+ {
+  "model": "IQ3_S",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 144.6,
+  "decode_tok_s": 45.51,
+  "decoded": 256,
+  "prefill_tok_s": 1443.21,
+  "ttft_s": 90.72,
+  "spec_accept": 0.697,
+  "tokens_per_round": 2.73,
+  "vram_slots": 2291
+ },
+ {
+  "model": "Coder",
+  "tier": "1k",
+  "prompt_tokens": 1017,
+  "exit": 0,
+  "wall_s": 37.0,
+  "decode_tok_s": 58.91,
+  "decoded": 256,
+  "prefill_tok_s": 656.36,
+  "ttft_s": 2.48,
+  "spec_accept": 0.743,
+  "tokens_per_round": 2.91,
+  "vram_slots": 2649
+ },
+ {
+  "model": "Coder",
+  "tier": "4k",
+  "prompt_tokens": 3562,
+  "exit": 0,
+  "wall_s": 18.3,
+  "decode_tok_s": 55.1,
+  "decoded": 256,
+  "prefill_tok_s": 1583.1,
+  "ttft_s": 3.09,
+  "spec_accept": 0.713,
+  "tokens_per_round": 2.88,
+  "vram_slots": 2614
+ },
+ {
+  "model": "Coder",
+  "tier": "32k",
+  "prompt_tokens": 31566,
+  "exit": 0,
+  "wall_s": 30.4,
+  "decode_tok_s": 54.89,
+  "decoded": 256,
+  "prefill_tok_s": 2177.17,
+  "ttft_s": 15.33,
+  "spec_accept": 0.731,
+  "tokens_per_round": 2.75,
+  "vram_slots": 2417
+ },
+ {
+  "model": "Coder",
+  "tier": "64k",
+  "prompt_tokens": 64162,
+  "exit": 0,
+  "wall_s": 44.8,
+  "decode_tok_s": 53.22,
+  "decoded": 256,
+  "prefill_tok_s": 2236.49,
+  "ttft_s": 29.51,
+  "spec_accept": 0.709,
+  "tokens_per_round": 2.68,
+  "vram_slots": 2382
+ },
+ {
+  "model": "Coder",
+  "tier": "128k",
+  "prompt_tokens": 128478,
+  "exit": 0,
+  "wall_s": 75.9,
+  "decode_tok_s": 43.04,
+  "decoded": 256,
+  "prefill_tok_s": 2208.27,
+  "ttft_s": 59.02,
+  "spec_accept": 0.553,
+  "tokens_per_round": 2.23,
+  "vram_slots": 2324
+ }
+]
@@ -12,7 +12,7 @@ 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, engine 0.1.22 (prompts) / 0.1.14 (output) with the settings setup writes
+RTX 5070 **12 GB**, Ryzen 5 7600 (6 cores), 64 GB DDR5-5200, Windows, engine 0.1.26 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.
@@ -21,24 +21,27 @@ measured with Swift 1.5's IQ2_XS, which runs at the original's speed.

 | Model | 1K | 4K | 32K | 64K | 128K | 262K |
 | --- | ---: | ---: | ---: | ---: | ---: | ---: |
-| **Q2_0** | 519 | 1,226 | 1,844 | 1,836 | 1,682 | 1,304 |
-| **IQ2_XS** | 524 | 1,196 | 1,799 | 1,611 | 1,495 | 1,181* |
-| **IQ3_XXS** | 472 | 974 | 1,555 | 1,449 | 1,386 | - |
-| **IQ3_S** | 419 | 893 | 1,499 | 1,285 | 1,245 | - |
-| **Coder** | 660 | 1,522 | 1,871 | 1,938 | 1,779 | 1,034** |
+| **Q2_0** | 536 | 1,299 | 2,171 | 2,126 | 2,107 | 1,304† |
+| **IQ2_XS** | 534 | 1,256 | 2,092 | 1,754 | 1,752 | 1,181*† |
+| **IQ3_XXS** | 482 | 1,007 | 1,745 | 1,609 | 1,602 | - |
+| **IQ3_S** | 427 | 913 | 1,624 | 1,640 | 1,443 | - |
+| **Coder** | 656 | 1,583 | 2,177 | 2,236 | 2,208 | 1,034** |

-Engine 0.1.22 (the prompt path of 0.1.23 is the same); `bench/results/2026-09-29-speed-0122`. \* measured with
-images on (the image encoder's VRAM reserve leaves fewer experts cached). \*\* not measured again: 0.1.14.
+Engine 0.1.26; `bench/results/2026-09-29-speed-0126`. At 32K-128K that is 8-28% faster than 0.1.22. † not measured
+again: 0.1.22. \* measured with images on (the image encoder's VRAM reserve leaves fewer experts cached). \*\* not
+measured again: 0.1.14.

 ### Output (tokens/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 | - |
-| **Coder** | 53.3 | 50.6 | 53.3 | 50.8 | 44.0 | 42.8 |
+| **Q2_0** | 87.3 | 93.0 | 81.8 | 76.2 | 73.7 | 60.3† |
+| **IQ2_XS** | 79.6 | 78.6 | 76.3 | 63.7 | 62.7 | 52.8† |
+| **IQ3_XXS** | 61.9 | 61.6 | 58.5 | 57.2 | 49.0 | - |
+| **IQ3_S** | 52.4 | 53.3 | 48.3 | 46.3 | 45.5 | - |
+| **Coder** | 58.9 | 55.1 | 54.9 | 53.2 | 43.0 | 42.8† |
+
+Engine 0.1.26, the same runs. † not measured again: 0.1.14.

 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,
@@ -67,6 +70,22 @@ VRAM. RTX 3090, the Coder at 198K context: 99 instead of 85 tokens/s output, the
 slower. It does not stream its KV cache (KV streaming is on by default from 64K), so it pays off mostly on large
 cards at long contexts.

+**The draft layer's tokens (0.1.27, `--draft-vocab`):** the MTP draft layer can only propose tokens from a subset
+of the vocabulary (`mtp/rt/draft_vocab.bin`). Since 0.1.27 the subset includes every Chinese, Japanese and Korean
+token (106,299 ids), so answers in those languages are 15-38% faster (Q2_0, RTX 5070). Its head takes ~180 MiB of
+VRAM, which the expert cache leaves free for it (0.1.28). `START-HERE.bat --setup --draft-vocab en` keeps the
+English/code subset from before (40,525 ids, ~110 MiB less VRAM, English answers 1-2% faster; CJK answers get
+almost no drafts). `tools/draft_vocab.py` builds and inspects subsets.
+
+**Low-RAM mode (engine 0.1.26, chosen by setup):** normally all of a model's experts are copied into RAM (23-50 GB,
+pinned) and the GPU holds a copy of the most-used ones. On a PC whose RAM cannot hold them beside the system (the
+experts plus ~10 GB), setup instead maps them from one file in the model's folder (`--mmap-experts`, the pack's
+`experts.bin`, +23-50 GB of disk). The OS file cache holds what the GPU does not, and it can give that memory back.
+On the Coder the engine's committed memory drops from 36 to ~13 GB, with the same answers. With a big GPU (an RTX
+5090 holds all of the Coder's experts, most of Q2_0's) it runs at nearly the usual speed. With a small one, most
+experts come from the SSD and it is much slower (setup says so). `START-HERE.bat --setup --low-ram on|off` overrides
+the choice.
+
 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).

@@ -157,7 +176,7 @@ You need **only an NVIDIA driver** (version 580 or newer; update it with the NVI

 | | |
 | --- | --- |
-| GPU | NVIDIA **RTX 30, 40 or 50 series**, **12 GB VRAM or more** (8 GB runs, slowly). Measured on an RTX 5070; RTX 30/40 are untested. |
+| GPU | NVIDIA **RTX 20, 30, 40 or 50 series**, **12 GB VRAM or more** (8 GB runs, slowly). Measured on an RTX 5070 and an RTX 3090; RTX 20 (Turing, since 0.1.27) was tested by a contributor on an RTX 2070. |
 | RAM | **64 GB** recommended (see the table above). |
 | CPU | x86-64 with AVX2 (any Intel/AMD desktop CPU from the last ~8 years). AVX-512 (Ryzen 7000/9000) is a bit faster. |
 | 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. |
@@ -169,7 +188,7 @@ elsewhere) finds them and sets itself up the same way. The place is remembered p
 `~/.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
+(from pip, ~0.4 GB), the ready-made Strata engine for RTX 20/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
 Windows, `build-essential` + CUDA on Ubuntu) and compiles the engine for your GPU (asks first; 20-40 minutes once).

@@ -18,7 +18,7 @@ Strata can use it:
 ```
  Your NVIDIA GPUs:
    GPU 0: NVIDIA GeForce RTX 5080, 16 GB VRAM - can be used
-    GPU 1: NVIDIA GeForce RTX 2080 Ti, 22 GB VRAM - not supported - older than the RTX 30 series (compute capability 7.5; Strata needs 8.0 or newer)
+    GPU 1: NVIDIA GeForce GTX 1080 Ti, 11 GB VRAM - not supported - older than the RTX 20 series (compute capability 6.1; Strata needs 7.5 or newer)
    GPU 2: NVIDIA GeForce RTX 3090, 24 GB VRAM - can be used
  ...
  1) GPU 0 (NVIDIA GeForce RTX 5080, 16 GB) + GPU 2 (NVIDIA GeForce RTX 3090, 24 GB) together   (recommended)
@@ -41,7 +41,7 @@ now on; the answer is kept.
 ```

 **Not supported** (setup says so and names the cards that can be used instead):
- a card older than the RTX 30 series (compute capability below 8.0: RTX 20, GTX 16/10);
+- a card older than the RTX 20 series (compute capability below 7.5: GTX 10 and older);
 - a card with less than 8 GB of VRAM, together with others (each card holds a copy of the dense weights and its
  own prompt buffers);
 - AMD and Intel GPUs, and a mix of NVIDIA with them.
@@ -96,8 +96,8 @@ into the card that owns the layer.
  capped to leave room for them.
 - On Windows only 8 GiB of the expert arena is pinned (more, mapped into two GPU contexts, leaves WDDM refusing
  allocations); the PCIe share covers those layers.
- Every card needs compute capability 8.0 (RTX 30 or newer). A Turing card (RTX 20, sm_75) builds only with the
-  experimental `-DSTRATA_EXPERIMENTAL_SM75=ON`.
+- Every card needs compute capability 7.5 (RTX 20 or newer). The pre-sm_80 QSA scorer path is fp32 FMAs, so a
+  Turing card runs the same kernels instead of the tensor-core prompt attention.

 ## Measured

@@ -53,6 +53,9 @@ public:
    /// KV streaming: refill the ring of the drafter's window from its host copy for a sequence that continues at
    /// `upto` (a conversation-cache resume). No-op unless the drafter's K/V is a ring.
    void kv_restore(int64_t upto);
+    /// The VRAM bind() will allocate for a native head of `head_row_bytes` per vocabulary row: the draft logits and
+    /// the draft head over rt/draft_vocab.bin's subset.  The expert cache is sized before bind(), so it reserves this.
+    uint64_t bind_bytes(uint64_t head_row_bytes, int64_t n_vocab) const;
    /// The main model's embedding and head, and the verify window's final residuals (T rows, hc*n_embd each).
    bool bind(const WeightTable& wt, const NativeHead* head, const float* window_R, std::string& err);

@@ -73,6 +76,19 @@ public:
    double ms_draft = 0, ms_prefill = 0;
    int64_t rounds = 0;

+    /// E-9: the prompt path computes this layer's prompt K/V in batches (Prefill::draft_kv): its tensors, its K/V
+    /// state, the first cell a round can still read, its device, and a wait for its own stream.
+    const float* tensor_f32(const char* name) const { return f32(name); }
+    const uint16_t* tensor_bf16(const char* name) const { return bf16(name); }
+    const void* tensor_q8(const char* name) const { return q8(name); }
+    QsaState& kv_state_rw() { return st_; }
+    int64_t first_needed() const { return (window_ > 0 && prompt_len_ > 0) ? prompt_len_ - window_ - 64 : 0; }
+    int device() const { return device_; }
+    bool idle(std::string& err) {
+        if (cs_ && cudaStreamSynchronize(cs_) != cudaSuccess) { err = "mtp: its stream failed"; return false; }
+        return true;
+    }
+
 private:
    bool record_forward(int T, int step_row0, cudaStream_t cs, std::string& err);
    bool capture_prefill(int T, std::string& err);
@@ -34,8 +34,19 @@ struct PrefillStats {
    double ms_ple = 0;
 };

+}  // namespace strata::prefill
+namespace strata::core { class MtpDrafter; }
+namespace strata::prefill {
+
 class Prefill {
 public:
+    /// E-9: the draft layer's K/V for prompt cells [cell0, cell0 + n) from their final residual rows `R_rows`
+    /// (device) and `next_tokens` (host: the token at cell+1), in batches through this path's GEMMs and its idle
+    /// scratch - call it from on_chunk.  false with `err` empty: not applicable here (a ring or hybrid K/V, another
+    /// device, too little scratch; STRATA_MTP_BATCH=0), the caller runs the drafter's own pass.  Not bit-identical to
+    /// that pass (FP16 GEMMs instead of Q8_1 activations): the drafts may differ, never the target's tokens' logits.
+    bool draft_kv(core::MtpDrafter& mtp, const float* R_rows, const int32_t* next_tokens, int64_t n, int64_t cell0,
+                  std::string& err);
    Prefill();
    ~Prefill();
    Prefill(const Prefill&) = delete;
@@ -242,6 +242,61 @@ CALL_START = "<tool_call>"
 CALL_END = "</tool_call>"


+PARAM_END = "</parameter>"
+FUNC_END = "</function>"
+
+
+def param_end(text: str, final: bool = False) -> int:
+    """Where a parameter value in `text` ends: the first `</parameter>` followed (after whitespace) by the next
+    `<parameter=` or `</function>` - the same text inside a value (a file that documents the call format, #210) is
+    part of the value.  -1: none yet; -2: a candidate whose follower has not arrived (streaming; `final` accepts it)."""
+    at = text.find(PARAM_END)
+    while at >= 0:
+        after = text[at + len(PARAM_END):].lstrip()
+        if after.startswith(("<parameter=", FUNC_END)):
+            return at
+        if not after or "<parameter=".startswith(after) or FUNC_END.startswith(after):
+            return at if final else -2
+        at = text.find(PARAM_END, at + 1)
+    return -1
+
+
+def call_end(text: str) -> int:
+    """Where a tool call's body ends (#210): the `</tool_call>` after the call's own `</function>`, found by walking
+    its parameters with param_end, so a value may contain either tag.  -1: not complete yet.  A body that is not in
+    the call format ends at the first `</tool_call>`, as before."""
+    s = text.lstrip()
+    pos = len(text) - len(s)
+    if not s.startswith("<function="):
+        return text.find(CALL_END) if not "<function=".startswith(s) else -1
+    gt = text.find(">", pos)
+    if gt < 0:
+        return -1
+    pos = gt + 1
+    while True:
+        rest = text[pos:]
+        s = rest.lstrip()
+        pos += len(rest) - len(s)
+        if s.startswith("<parameter="):
+            gt = text.find(">", pos)
+            if gt < 0:
+                return -1
+            end = param_end(text[gt + 1:])
+            if end < 0:
+                return -1
+            pos = gt + 1 + end + len(PARAM_END)
+        elif s.startswith(FUNC_END):
+            rest = text[pos + len(FUNC_END):]
+            s = rest.lstrip()
+            if s.startswith(CALL_END):
+                return pos + len(FUNC_END) + len(rest) - len(s)
+            return -1 if CALL_END.startswith(s) else text.find(CALL_END, pos)
+        elif not s or "<parameter=".startswith(s) or FUNC_END.startswith(s):
+            return -1
+        else:
+            return text.find(CALL_END, pos)
+
+
 def parse_tool_call(body: str, schema: dict | None = None) -> ToolCall:
    """`<function=NAME>\\n<parameter=P>\\nVALUE\\n</parameter>...</function>` -> ToolCall. Values are JSON-decoded
    when the tool's schema says the parameter is not a string (or, without a schema, when they parse as JSON
@@ -257,9 +312,9 @@ def parse_tool_call(body: str, schema: dict | None = None) -> ToolCall:
        rest = rest[rest.index("<parameter=") + len("<parameter="):]
        pname = rest[:rest.index(">")]
        rest = rest[rest.index(">") + 1:]
-        end = rest.find("</parameter>")
+        end = param_end(rest, final=True)
        value = rest[:end] if end >= 0 else rest
-        rest = rest[end + len("</parameter>"):] if end >= 0 else ""
+        rest = rest[end + len(PARAM_END):] if end >= 0 else ""
        if value.startswith("\n"):
            value = value[1:]
        if value.endswith("\n"):
@@ -351,16 +406,17 @@ class OutputParser:
                        self.sp += 1
                        rest = rest[1:]
                    self.sval_started = True
-                end = rest.find("</parameter>")
+                end = param_end(rest)
                if end >= 0:
                    value = rest[:end]
                    if value.endswith("\n"):
                        value = value[:-1]
                    args(json.dumps(value)[1:-1] + '"')
-                    self.sp += end + len("</parameter>")
+                    self.sp += end + len(PARAM_END)
                    self.ss = "between"
                    continue
-                safe = len(rest) - self._hold(rest, ("</parameter>",))
+                # an undecided </parameter> (-2) is held from its start, like a tag still arriving
+                safe = rest.find(PARAM_END) if end == -2 else len(rest) - self._hold(rest, (PARAM_END,))
                if safe > 0 and rest[safe - 1] == "\n":   # may be the trailing newline before </parameter>
                    safe -= 1
                if safe > 0:
@@ -368,7 +424,7 @@ class OutputParser:
                    self.sp += safe
                return out
            elif self.ss == "raw":
-                end = rest.find("</parameter>")
+                end = param_end(rest)
                if end < 0:
                    return out
                value = rest[:end]
@@ -381,7 +437,7 @@ class OutputParser:
                except ValueError:
                    v = value
                args(json.dumps(v, ensure_ascii=False))
-                self.sp += end + len("</parameter>")
+                self.sp += end + len(PARAM_END)
                self.ss = "between"
            else:
                return out
@@ -449,7 +505,7 @@ class OutputParser:
                self.buf = self.buf[i + len(CALL_START):]
                self.state = "call"
            else:
-                i = self.buf.find(CALL_END)
+                i = call_end(self.buf)
                if self.stream_tools:
                    if i >= 0:
                        whole, self.buf = self.buf, self.buf[:i]     # scan only the body
@@ -24,6 +24,7 @@ import argparse
 import collections
 import base64
 import hashlib
+import hmac
 import codecs
 import json
 import os
@@ -49,6 +50,7 @@ from serve.mcp import McpCancelled, hub_from_config  # noqa: E402

 IM_END = "<|im_end|>"
 IMAGE_PAD = "<|image_pad|>"
+VISION_START = "<|vision_start|>"
 CTX_SLACK = 8               # `strata --serve` rejects prompt + max_new + 8 > context: keep the same margin here
 # The live tok/s is a rate over a window, not a mean since the first token: a mean reads ~1/elapsed at the first
 # token (the Monitor showed five-digit numbers) and then undershoots for the first second of every answer.
@@ -215,6 +217,11 @@ class StrataEngine:
            if "issue #29" in line:
                return ("The engine stopped itself because it had stopped making progress - a hang it caught. Its log "
                        "line: " + line.strip() + " - please report it at github.com/Niko1221/Strata/issues.")
+        rc = self.proc.poll()
+        last = next((x.strip() for x in reversed(tail.splitlines()) if x.strip().startswith(("strata", "ERR"))), "")
+        if rc is not None and rc >= 0 and last:          # it ended by itself: its own last words say why (#215)
+            return (f"The engine exited (code {rc}). Its last log line: {last} - if that does not explain it, please "
+                    "report it at github.com/Niko1221/Strata/issues with the log.")
        return ("The usual cause is running out of RAM: Linux then ends the biggest program (check: sudo dmesg | "
                "grep -i -E 'killed process|out of memory'); Windows slows down instead. Close other programs or use a "
                "smaller model (Q2_0 / IQ2_XS).")
@@ -720,17 +727,24 @@ class Service:
                raise ValueError("this server was started without the vision encoder (run setup again and choose "
                                 "'vision'), so it cannot read images")
            pad = self.tok.encode(IMAGE_PAD, parse_special=True)[0]
+            start = self.tok.encode(VISION_START, parse_special=True)[0]
            # Encode only while the engine is idle: the engine and the image encoder (a separate process) must not
            # run on the GPU at the same time - an encode during a running request left that request stuck at
            # "reading the prompt" with CPU and GPU busy, for good (reproduced).  So encoding takes its turn in the
            # same FIFO as the requests.
            with self.fifo:
                encoded = [self.vision.encode(src) for src in images]
+            # one <|image_pad|> per image -> one per image token.  Only the markers the template writes for an image
+            # (right after <|vision_start|>) are images: the same text inside a message (an agent reading these docs,
+            # #150) is kept as plain text, or it took an image's place and the counts no longer matched.
+            literal = self.tok.encode(IMAGE_PAD, parse_special=False)
            out, k = [], 0
-            for t in ids:                               # one <|image_pad|> per image -> one per image token
-                if t == pad and k < len(encoded):
+            for j, t in enumerate(ids):
+                if t == pad and j > 0 and ids[j - 1] == start and k < len(encoded):
                    out += [pad] * encoded[k][1]
                    k += 1
+                elif t == pad:
+                    out += literal
                else:
                    out.append(t)
            if k != len(encoded):
@@ -906,6 +920,8 @@ class Service:
                    if os.environ.get("STRATA_DEBUG") and raw_ids:
                        print(f"[strata] raw: {self.tok.decode(raw_ids)!r}", flush=True)
                self.status["busy"] = False
+                self.status.pop("tail", None)            # #212: the answer's end is not kept once it is done
+                self.status.pop("tool", None)
        for ev in parser.finish():
            yield "event", ev
        yield "done", {"finish": finish, "completion_tokens": n, "reused": (timings or {}).get("cache_n", 0),
@@ -1237,7 +1253,7 @@ def make_handler(svc: Service):
                return True
            auth = self.headers.get("Authorization", "")
            given = auth[7:].strip() if auth.lower().startswith("bearer ") else self.headers.get("x-api-key", "")
-            if given == svc.api_key:
+            if hmac.compare_digest(given.encode(), svc.api_key.encode()):   # #213: constant-time
                return True
            self._json(401, {"error": {"type": "authentication_error", "message": "missing or wrong API key"}})
            return False
@@ -1302,6 +1318,8 @@ def make_handler(svc: Service):
                self._json(200, {"status": "ok", "max_context": svc.engine.max_context, "model": svc.model,
                                 "images": svc.vision is not None, "api_key": bool(svc.api_key)})
            elif path == "/status":
+                if not self._authorized():                  # #212: it shows the end of the last answer
+                    return
                with svc.status_lock:
                    s = dict(svc.status)
                now = time.time()
@@ -1493,6 +1511,10 @@ class Server(ThreadingHTTPServer):
    # either one (a forgotten second start of run-<model>.bat).  Without it the second start fails loudly instead.
    allow_reuse_address = os.name != "nt"

+    def handle_error(self, request, client_address):
+        if not isinstance(sys.exc_info()[1], ConnectionError):   # a client that hangs up needs no stack trace
+            super().handle_error(request, client_address)
+

 def warn_tight_ram(arena_mib) -> None:
    """The model's experts live in RAM (INFO arena_mib, engine 0.1.10+).  With less than ~6 GB left beside them for the
@@ -1716,6 +1738,12 @@ def main() -> int:
                  model_name=cfg.get("model_name", "qwen3.8-flash-next"), vision=vision,
                  sampling_defaults=sampling_defaults,
                  fit_max_tokens=a.fit_max_tokens or cfg.get("fit_max_tokens") is True)
+    if ("STRATA_API_KEY" in os.environ and not os.environ["STRATA_API_KEY"].strip()) or             any(x == "--api-key" and i + 1 < len(sys.argv) and not sys.argv[i + 1].strip() or x.strip() == "--api-key="
+                for i, x in enumerate(sys.argv)):
+        # #213: an empty key would switch authentication off without a word
+        print("[strata] an API key was given but it is empty: set a key, or leave --api-key / STRATA_API_KEY out",
+              file=sys.stderr)
+        return 2
    svc.api_key = a.api_key or cfg.get("api_key", "")
    svc.gpu_index = (gpu_list(cfg) or [0])[0]           # the Monitor reads the card the engine runs on (issue #51)
    svc.gpu_indices = gpu_list(cfg)                     # ... or every card of a layer split (issue #112)
@@ -1760,15 +1788,19 @@ def main() -> int:
        import webbrowser
        webbrowser.open(f"http://{'127.0.0.1' if a.host in ('0.0.0.0', '') else a.host}:{a.port}/")
    try:
-        threading.Event().wait()
+        while True:
+            time.sleep(1)                               # Windows never delivers Ctrl+C to an untimed Event.wait()
    except KeyboardInterrupt:
-        httpd.shutdown()
-        if hasattr(engine, "close"):
-            engine.close()
-        if vision:
-            vision.close()
-        if hub is not None:
-            hub.close()
+        print("\n[strata] stopping (Ctrl+C again to end the engine at once) ...", flush=True)
+        closers = [httpd.shutdown, getattr(engine, "close", None), vision.close if vision else None,
+                   hub.close if hub is not None else None]
+        for close in filter(None, closers):
+            try:
+                close()
+            except KeyboardInterrupt:                   # a second Ctrl+C: don't wait for the engine to free its memory
+                if getattr(engine, "proc", None):
+                    engine.proc.kill()
+        print("[strata] stopped", flush=True)
    return 0


@@ -167,6 +167,90 @@ class FitMaxTokens(unittest.TestCase):
        self.assertIn("no room to answer", b["error"]["message"])


+class ImageMarkers(unittest.TestCase):
+    """#150: the text "<|image_pad|>" inside a message is text, not an image's place."""
+
+    class FakeVision:
+        def __init__(self, d):
+            self.dir = Path(d)
+            self.rows = self.dir / "img.sve"
+            self.rows.write_bytes(b"rows")
+
+        def encode(self, source):
+            return self.rows, 3
+
+    def test_literal_marker_with_an_image(self):
+        import tempfile
+        tok = ByteTokenizer()
+        with tempfile.TemporaryDirectory() as d:
+            svc = Service(MockEngine(tok, "ok", max_context=CTX), tok, ChatTemplate(ROOT / "serve/chat_template.jinja"),
+                          vision=self.FakeVision(d))
+            pad = tok.encode("<|image_pad|>", parse_special=True)[0]
+            for text in ("the docs say <|image_pad|> marks an image", "plain"):
+                with self.subTest(text=text):
+                    msgs = [{"role": "user", "content": [{"type": "text", "text": text},
+                                                         {"type": "image", "source": "x.png"}]}]
+                    ids, _, _ = svc.prepare(msgs, None, {})
+                    self.assertEqual(ids.count(pad), 3)          # the image's three rows, nothing else
+                    self.assertIn("<|image_pad|> marks" if "docs" in text else "plain", tok.decode(ids))
+            svc.embeddings.path.unlink(missing_ok=True)
+
+
+class StatusNeedsTheKey(unittest.TestCase):
+    """#212: /status shows the end of the answer being written, so it needs the key like /v1/*."""
+
+    def test_status(self):
+        tok = ByteTokenizer()
+        svc = Service(MockEngine(tok, "ok", max_context=CTX), tok, ChatTemplate(ROOT / "serve/chat_template.jinja"))
+        svc.api_key = "k3y"
+        httpd = serve(svc, port=0)
+        base = f"http://127.0.0.1:{httpd.server_address[1]}/status"
+        try:
+            with self.assertRaises(urllib.error.HTTPError) as e:
+                urllib.request.urlopen(base, timeout=10)
+            self.assertEqual(e.exception.code, 401)
+            e.exception.close()
+            req = urllib.request.Request(base, headers={"Authorization": "Bearer k3y"})
+            with urllib.request.urlopen(req, timeout=10) as r:
+                self.assertEqual(r.status, 200)
+                self.assertNotIn("tail", json.loads(r.read()))
+        finally:
+            httpd.shutdown()
+            httpd.server_close()
+
+
+class ToolCallTerminators(unittest.TestCase):
+    """#210: a value that contains </parameter> or </tool_call> (a file documenting the call format) is kept whole."""
+    CONTENT = ("Close each value with </parameter> and the call with </function></tool_call>.\n"
+               "<parameter=x>\nnot a parameter\n</parameter>\nend")
+    SCHEMA = [{"name": "write", "parameters": {"properties": {"path": {"type": "string"},
+                                                              "content": {"type": "string"}}}}]
+
+    def run_parser(self, stream_tools, step):
+        from serve.frontend import OutputParser
+        text = ("</think>\n\n<tool_call>\n<function=write>\n<parameter=path>\ndoc.md\n</parameter>\n"
+                f"<parameter=content>\n{self.CONTENT}\n</parameter>\n</function>\n</tool_call>")
+        p = OutputParser(thinking=True, tools=self.SCHEMA, stream_tools=stream_tools)
+        evs = []
+        for i in range(0, len(text), step):
+            evs += p.feed(text[i:i + step])
+        evs += p.finish()
+        return evs
+
+    def test_values_keep_the_terminators(self):
+        for stream_tools in (False, True):
+            for step in (1, 7, 10_000):
+                with self.subTest(stream_tools=stream_tools, step=step):
+                    evs = self.run_parser(stream_tools, step)
+                    calls = [e.call for e in evs if e.kind == "tool_call"]
+                    self.assertEqual(len(calls), 1)
+                    self.assertEqual(calls[0].arguments, {"path": "doc.md", "content": self.CONTENT})
+                    self.assertFalse([e for e in evs if e.kind == "content" and e.text.strip()])
+                    if stream_tools:
+                        streamed = "".join(e.text for e in evs if e.kind == "tool_args")
+                        self.assertEqual(json.loads(streamed), {"path": "doc.md", "content": self.CONTENT})
+
+
 class ClientShapes(unittest.TestCase):
    """What real clients send: Claude Code posts /v1/messages?beta=true (issue #55) and puts hook context into the
    conversation as a mid-conversation system message (issue #56); some OpenAI clients send a late developer message."""
@@ -13,7 +13,7 @@ What the first run does (each step is skipped when it is already done):
  1. checks your PC: NVIDIA GPU and driver, RAM, CPU, free disk space
  2. asks the questions
  3. installs the Python packages it needs into .venv (numpy, jinja2, ..., and NVIDIA's CUDA libraries)
-  4. gets the Strata engine: a ready-made build for RTX 30/40/50 cards (no compiler needed); if none fits your PC,
+  4. gets the Strata engine: a ready-made build for RTX 20/30/40/50 cards (no compiler needed); if none fits your PC,
     it installs the build tools (asks first) and compiles the engine for your GPU
  5. downloads the model from Hugging Face (resumable), and the vision encoder if you want images
  6. prepares the model for Strata and fetches the MTP draft layer (~5 GB, from the original Qwen checkpoint)
@@ -58,11 +58,13 @@ 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, 25)                # v0.1.25: faster prompts (grouping off the copy engine, fused hyper-connection kernels), AMD HIP backend, --kv k8v4; v0.1.24: long prompts faster (QSA select on tensor cores); v0.1.23: image requests honor sampling, 8 GB cards start, batched verify window; v0.1.22: faster prompts (tensor-core attention), multi-GPU across images/steering/KV streaming; v0.1.21: multi-GPU layer split (--gpus); v0.1.20: system-prompt checkpoint, PCIe probe, hit rate; v0.1.19: penalties
+MIN_ENGINE = (0, 1, 28)                # v0.1.28: the expert cache reserves the draft head, a cancelled request no longer fails the next; v0.1.27: RTX 20 (sm_75) in the ready-made engine, the HIP build without CUDA headers; v0.1.26: the draft layer's prompt pass in batches; v0.1.25: faster prompts (grouping off the copy engine, fused hyper-connection kernels), AMD HIP backend, --kv k8v4; v0.1.24: long prompts faster (QSA select on tensor cores); v0.1.23: image requests honor sampling, 8 GB cards start, batched verify window; v0.1.22: faster prompts (tensor-core attention), multi-GPU across images/steering/KV streaming; v0.1.21: multi-GPU layer split (--gpus); v0.1.20: system-prompt checkpoint, PCIe probe, hit rate; v0.1.19: penalties
 PY_PACKAGES = ["numpy", "jinja2", "regex", "pyyaml", "tqdm", "requests", "cmake", "ninja", "pillow", "psutil"]

 MODELS = {
-    "Q2_0": {"about": "2-bit, the fastest", "download_gb": 66.4, "ram_gb": 48, "arena_gb": 34.0},
+    # the original model only for now: Swift 1.5's Q2_0 files split one layer's experts across the two shards, which
+    # the pack tool (tools/iq_pack.py) cannot prepare yet (#171)
+    "Q2_0": {"about": "2-bit, the fastest", "download_gb": 66.4, "ram_gb": 48, "arena_gb": 34.0, "families": ("qwen",)},
    "IQ2_XS": {"about": "2-bit i-quant, a little better quality, close in speed", "download_gb": 68.0, "ram_gb": 48,
               "arena_gb": 35.5},
    "IQ3_XXS": {"about": "3-bit i-quant, better quality, slower (more CPU work per token)", "download_gb": 75.8,
@@ -221,7 +223,7 @@ def cpu_info():
    name, avx2, avx512 = platform.processor() or "unknown CPU", False, False
    if WIN:
        pf = ctypes.windll.kernel32.IsProcessorFeaturePresent
-        avx2 = bool(pf(40))                       # PF_AVX2_INSTRUCTIONS_AVAILABLE
+        avx2 = bool(pf(40)) or _cpuid_avx2()      # PF_AVX2_INSTRUCTIONS_AVAILABLE, else the CPU itself (#159)
        n = out(["powershell", "-NoProfile", "-Command", "(Get-CimInstance Win32_Processor).Name"]).strip()
        name = n or name
        avx512 = bool(pf(41)) and _cpuid_avx512_full()
@@ -259,6 +261,47 @@ def _cpuid_avx512_full() -> bool:
        return False


+def _run_stub(code: bytes, *args) -> None:
+    """Runs a few bytes of x64 machine code (Windows calling convention: the arguments in rcx, rdx)."""
+    k32 = ctypes.windll.kernel32
+    k32.VirtualAlloc.restype = ctypes.c_void_p
+    k32.VirtualFree.argtypes = (ctypes.c_void_p, ctypes.c_size_t, ctypes.c_uint32)
+    buf = k32.VirtualAlloc(None, len(code), 0x3000, 0x40)
+    if not buf:
+        raise OSError("VirtualAlloc failed")
+    try:
+        ctypes.memmove(buf, code, len(code))
+        ctypes.CFUNCTYPE(None, *[ctypes.c_void_p] * len(args))(buf)(*args)
+    finally:
+        k32.VirtualFree(buf, 0, 0x8000)
+
+
+def _cpuid_avx2() -> bool:
+    """AVX2 asked from the CPU (CPUID leaf 7 EBX bit 5), with the OS saving the YMM registers (OSXSAVE + XCR0):
+    Windows' IsProcessorFeaturePresent(PF_AVX2) says no on some PCs whose CPU has it (a Ryzen 9 3950X, #159)."""
+    try:
+        def cpuid(leaf):
+            regs = (ctypes.c_uint32 * 4)()
+            _run_stub(bytes([0x53, 0x49, 0x89, 0xC8, 0x89, 0xD0, 0x31, 0xC9, 0x0F, 0xA2,      # push rbx; r8=rcx; eax=edx; ecx=0; cpuid
+                             0x41, 0x89, 0x00, 0x41, 0x89, 0x58, 0x04, 0x41, 0x89, 0x48, 0x08,  # [r8]=eax, [r8+4]=ebx, [r8+8]=ecx
+                             0x41, 0x89, 0x50, 0x0C, 0x5B, 0xC3]),                             # [r8+12]=edx; pop rbx
+                      ctypes.addressof(regs), leaf)
+            return list(regs)
+        if cpuid(0)[0] < 7:
+            return False
+        ecx1 = cpuid(1)[2]
+        if not (ecx1 >> 27) & 1 or not (ecx1 >> 28) & 1:             # OSXSAVE, AVX
+            return False
+        xcr0 = (ctypes.c_uint32 * 2)()
+        _run_stub(bytes([0x49, 0x89, 0xC8, 0x31, 0xC9, 0x0F, 0x01, 0xD0,                        # r8=rcx; ecx=0; xgetbv
+                         0x41, 0x89, 0x00, 0x41, 0x89, 0x50, 0x04, 0xC3]), ctypes.addressof(xcr0))
+        if xcr0[0] & 6 != 6:                                           # the OS saves XMM and YMM
+            return False
+        return bool((cpuid(7)[1] >> 5) & 1)
+    except Exception:
+        return False
+
+
 def gpus():
    """Every NVIDIA GPU, numbered as nvidia-smi numbers them (by PCI bus, the order the engine is told to use)."""
    s = out(["nvidia-smi", "--query-gpu=index,name,memory.total,compute_cap,driver_version",
@@ -285,8 +328,8 @@ def cc(g) -> str:

 def gpu_problem(g, together=False):
    """Why Strata cannot use this card, in plain words (None: it can)."""
-    if int(g["arch"]) < 80:
-        return (f"not supported - older than the RTX 30 series (compute capability {cc(g)}; Strata needs 8.0 or "
+    if int(g["arch"]) < 75:
+        return (f"not supported - older than the RTX 20 series (compute capability {cc(g)}; Strata needs 7.5 or "
                "newer)")
    if together and g["vram_gb"] < SPLIT_MIN_VRAM_GB - 0.5:
        return (f"not supported together with other GPUs - {g['vram_gb']:.0f} GB of VRAM (a card sharing the model "
@@ -349,7 +392,7 @@ def check_gpus(sel, found, what="") -> None:
        ones = " or ".join(f"--gpu {x['index']}" for x in single)
        both = "--gpus " + ",".join(str(x["index"]) for x in can) if can else ""
        hint = ((f"use these together: {both}" + (f" (or one card: {ones})" if not together else "")) if can else
-                f"use one card: {ones}" if single else "Strata needs an NVIDIA RTX 30 series or newer card")
+                f"use one card: {ones}" if single else "Strata needs an NVIDIA RTX 20 series or newer card")
        fail(f"GPU {i}{'' if g is None else ' (' + g['name'] + ')'} {what}cannot be used: {p}", hint)


@@ -384,7 +427,7 @@ def choose_gpus(a, found) -> list:
    single = sorted([g for g in found if gpu_problem(g) is None], key=lambda x: (-round(x["vram_gb"]), x["index"]))
    if not single:
        gpu_table(found)
-        fail("none of your GPUs can run Strata", "it needs an NVIDIA RTX 30 series or newer (compute capability 8.0+)")
+        fail("none of your GPUs can run Strata", "it needs an NVIDIA RTX 20 series or newer (compute capability 7.5+)")
    can = together_ok(found)
    if not can:
        return [single[0]["index"]]
@@ -566,6 +609,18 @@ def download(url, dst: Path, what=None):
    ok(f"{what or dst.name} downloaded")


+def whole_shard(s: Path) -> bool:
+    """A shard as long as its own tensor directory says (check_shards' test, without stopping setup)."""
+    sys.path.insert(0, str(ROOT / "tools"))
+    from gguf_reader import GGUFFile
+    try:
+        g = GGUFFile(s)
+        return s.stat().st_size >= g.data_start + max((t.offset + (t.expected_bytes() or 0) for t in g.tensors),
+                                                     default=0)
+    except (OSError, ValueError, struct.error):
+        return False
+
+
 def check_shards(shards):
    """Every shard present and whole, or setup stops naming the file and the numbers.  Whole means as long as
    its own tensor directory says (the header is read, the data is not): a truncated copy (--gguf-dir, a .part
@@ -597,7 +652,9 @@ def get_llama_cpp():
    tmp = ROOT / "third_party" / "_unpack"
    shutil.rmtree(tmp, ignore_errors=True)
    with zipfile.ZipFile(z) as f:
-        f.extractall(tmp)
+        # llama.cpp's own web UI (tools/ui) is not used, and its deep paths passed Windows' 260-character limit in a
+        # folder like Downloads\Strata-main\Strata-main (#206)
+        f.extractall(tmp, [m for m in f.namelist() if "/tools/ui/" not in m])
    top = next(tmp.iterdir())
    shutil.rmtree(llama, ignore_errors=True)
    # PR #63: on Windows a rename can fail with PermissionError while an antivirus scanner still holds a file of the
@@ -773,7 +830,16 @@ def get_prebuilt(url_base, gpu, vision, updating=False) -> Path | None:
    if info.exists() and (eng / EXE).exists():
        meta = json.loads(info.read_text())
        ver = tuple(int(x) for x in str(meta.get("version", "0")).split(".")[:3] if x.isdigit())
-        if meta.get("source") == "local" or ver >= MIN_ENGINE:
+        if meta.get("source") == "local":              # compiled here: build_engine checks its source and cards
+            return None
+        have = [int(a) for a in meta.get("archs", [])]
+        miss = [int(x) for x in gpu.get("archs", [gpu["arch"]])
+                if have and int(x) not in have and not (meta.get("ptx") and int(x) > max(have))]
+        if miss:                                       # a card it has no code for (#128): compiled here instead
+            warn(f"the installed engine is built for {', '.join(str(a) for a in have)}; your GPU is "
+                 f"{', '.join(str(x) for x in miss)}: compiling instead")
+            return None
+        if ver >= MIN_ENGINE:
            ok("ready-made engine already installed")
            return eng
        say(f"  Updating the ready-made engine ({meta.get('version')} -> {'.'.join(map(str, MIN_ENGINE))} or newer) ...")
@@ -898,7 +964,8 @@ def update_installed_engine(url_base) -> None:
 def install_build_tools(gpu, yes):
    """The compiler and the CUDA toolkit, installed for the user (asks once).  Returns (nvcc, vcvars)."""
    nvcc, cuda_v = find_nvcc()
-    need_cuda = (12, 8) if max(int(x) for x in gpu.get("archs", [gpu["arch"]])) >= 120 else (12, 0)
+    # RTX 50 (sm_120): CUDA 13.0 - an engine built with 12.8 crashed in the prompt path on Linux (#220)
+    need_cuda = (13, 0) if max(int(x) for x in gpu.get("archs", [gpu["arch"]])) >= 120 else (12, 0)
    vcvars = find_vcvars() if WIN else None
    have_cc = vcvars is not None if WIN else shutil.which("g++") is not None
    missing = []
@@ -1004,16 +1071,24 @@ def build_engine(gpu, vision, yes, llama) -> Path:
    want_vision = vision != "none"
    local = meta.get("source") == "local"
    src, vsrc = source_hash(ENGINE_SOURCES), source_hash(VISION_SOURCES)
-    engine_ok = local and (eng / EXE).exists() and meta.get("src") == src
+    archs = sorted({int(x) for x in gpu.get("archs", [gpu["arch"]])})    # every card the model runs on
+    built = {int(x) for x in meta.get("archs", [])}
+    # a card the engine has no code for (a GPU added with --gpus, #128) needs a compile even when the source is the
+    # same; the compile keeps the generations it was built for
+    new_arch = local and not set(archs) <= built
+    engine_ok = local and (eng / EXE).exists() and meta.get("src") == src and not new_arch
    vision_ok = not want_vision or ((eng / VEXE).exists() and (not local or meta.get("vision_src") == vsrc))
    if engine_ok and vision_ok:
        ok("engine already built for this PC")
        return eng
-    nvcc, vcvars = install_build_tools(gpu, yes)
-    archs = sorted({int(x) for x in gpu.get("archs", [gpu["arch"]])})    # every card the model runs on
+    if local:
+        archs = sorted(built | set(archs))
+    nvcc, vcvars = install_build_tools({**gpu, "archs": archs}, yes)
    cuda_archs = ";".join(str(x) for x in archs)
    if not engine_ok:
-        say("  The engine's source changed: compiling it again (only what changed, a few minutes) ..."
+        say("  Compiling the engine for " + ", ".join(f"sm_{x}" for x in archs) + " (a card it had no code for; "
+            "10-20 minutes, once) ..." if new_arch else
+            "  The engine's source changed: compiling it again (only what changed, a few minutes) ..."
            if local and (eng / EXE).exists() else "  Compiling the Strata engine for your GPU (10-20 minutes, once) ...")
        cmake_build(ROOT, ROOT / "build", "strata",
                    ["-DSTRATA_ENABLE_CUDA=ON", "-DSTRATA_BUILD_TESTS=OFF", f"-DCMAKE_CUDA_ARCHITECTURES={cuda_archs}",
@@ -1043,6 +1118,26 @@ def build_engine(gpu, vision, yes, llama) -> Path:
 DATA_ITEMS = ("models", "packs", "mtp")


+LOW_RAM_HEADROOM_GB = 10   # RAM beside the experts: the OS, the engine's other buffers, the server
+
+
+def low_ram_needed(model, ram) -> bool:
+    """The model's experts do not fit this PC's RAM with room left for the rest: they are then mapped from the pack's
+    experts.bin instead of copied into RAM (the low-RAM mode)."""
+    return ram < MODELS[model]["arena_gb"] + LOW_RAM_HEADROOM_GB
+
+
+def low_ram_gpu_share(model, vram_gb) -> float:
+    """About how much of the model's experts the GPU holds (VRAM minus ~5 GB for the dense weights, KV, buffers)."""
+    return max(0.0, min(1.0, (vram_gb - 5) / MODELS[model]["arena_gb"]))
+
+
+def low_ram_fits(model, ram, vram_gb) -> bool:
+    """In the low-RAM mode: the experts the GPU does not hold fit the RAM left beside the rest (as file cache)."""
+    arena = MODELS[model]["arena_gb"]
+    return ram - 6 + max(0.0, vram_gb - 5) >= arena
+
+
 def settings_path() -> Path:
    if WIN:
        return Path(os.environ.get("APPDATA") or Path.home() / "AppData" / "Roaming") / "Strata" / "settings.json"
@@ -1165,7 +1260,18 @@ def data_folder(requested: str | None) -> tuple:
        warn(f"cannot use {dest} for the model files ({e}): keeping them in {ROOT}")
        dest = ROOT
    elsewhere = []
-    for folder in [ROOT, *other_installs(settings)]:
+    # #198: the data folder remembered before (a --data-dir to a new place) is a source too, and so is a Strata-data
+    # folder nested in any of them (an install that kept its models one level down)
+    sources = [ROOT, *other_installs(settings)]
+    if settings.get("data_dir") and Path(settings["data_dir"]) != dest:
+        sources.append(Path(settings["data_dir"]))
+    sources += [f / "Strata-data" for f in list(sources) if (f / "Strata-data") != dest]
+    seen = set()
+    for folder in sources:
+        key = os.path.normcase(str(folder))
+        if key in seen:
+            continue
+        seen.add(key)
        if folder == dest or not has_data(folder):
            continue
        if not same_drive(folder, dest):
@@ -1338,12 +1444,21 @@ def upgrade_config(cfg_path: Path, cfg: dict) -> dict:


 def start(cfg_path: Path, port: int | None, gpu: int | list | None = None, open_browser=True, yes=False,
-          layer_split=None) -> int:
+          layer_split=None, keep=None) -> int:
+    """keep: settings given on this start that the model keeps from now on (--host, --api-key, --draft-vocab)."""
    cfg = upgrade_config(cfg_path, json.loads(cfg_path.read_text(encoding="utf-8-sig")))
    missing = [p for p in [cfg["exe"], *[a for a in cfg["args"] if a.endswith(".gguf")]] if not Path(p).exists()]
    if missing:
        fail(f"{cfg_path.name} refers to missing files: {missing[0]}", "run it again with --setup to repair")
+    keep = {k: v for k, v in (keep or {}).items() if v is not None}
+    if keep and any(cfg.get(k) != v for k, v in keep.items()):   # #179: a --host/--api-key on a start was ignored
+        cfg.update(keep)
+        cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
+        ok("saved for this model: " + ", ".join("api key" if k == "api_key" else f"{k.replace('_', ' ')} {v}"
+                                                for k, v in keep.items()))
    cfg_path.touch()                                     # the most recently used model
+    if "--mtp" in cfg["args"][:-1]:
+        refresh_draft_vocab(Path(cfg["args"][cfg["args"].index("--mtp") + 1]), cfg.get("draft_vocab", "cjk"))
    cmd = [sys.executable, str(ROOT / "serve" / "server.py"), "--engine", "strata", "--config", str(cfg_path),
           "--port", str(port or cfg.get("port", 8080))]
    if cfg.get("backend") == "hip":                    # AMD: one card, numbered as HIP numbers them
@@ -1378,11 +1493,13 @@ def start(cfg_path: Path, port: int | None, gpu: int | list | None = None, open_
    elif isinstance(use, list):
        check_gpus(use, found, "(chosen for this model) ")
        byid = {g["index"]: g for g in found}
+        cfg = ensure_engine_for([byid[i] for i in use], cfg_path, cfg, yes)
        ok("GPUs: " + " + ".join(gpu_name(byid[i]) for i in use) + f" together (layers split {cfg.get('layer_split') or 'auto'})")
    elif found:
        g = next((x for x in found if x["index"] == use), None) if use is not None else max(
            found, key=lambda x: (round(x["vram_gb"]), -x["index"]))
        if g is not None:
+            cfg = ensure_engine_for([g], cfg_path, cfg, yes)
            ok("GPU: " + gpu_name(g))
    if open_browser:
        cmd.append("--open")
@@ -1403,13 +1520,60 @@ def start(cfg_path: Path, port: int | None, gpu: int | list | None = None, open_
    return subprocess.call(cmd)


+# the draft subsets setup copied before (sha256): replaced by the current one, a subset made by hand is kept
+OLD_DRAFT_VOCABS = {"369151522226a5edaa5f12cfd1e2ae7db8f4fbdbd222f3dcf327dced9597fb25"}   # to 0.1.26: 27 Han tokens
+
+
+DRAFT_VOCABS = {"cjk": "draft_vocab.bin", "en": "draft_vocab_en.bin"}
+
+
+def refresh_draft_vocab(rt: Path, choice: str = "cjk") -> None:
+    """The draft layer's token subset in the MTP folder: `cjk` (data/draft_vocab.bin, since 0.1.27, #137) or `en`
+    (data/draft_vocab_en.bin, the English/code subset before it: ~110 MiB less VRAM, English answers 1-2% faster).
+    Copied when missing or when a shipped subset other than the chosen one is there; a subset made by hand is kept."""
+    new, dst = ROOT / "data" / DRAFT_VOCABS.get(choice, "draft_vocab.bin"), rt / "draft_vocab.bin"
+    if not new.exists() or not rt.is_dir():
+        return
+    if dst.exists():
+        old = hashlib.sha256(dst.read_bytes()).hexdigest()
+        shipped = OLD_DRAFT_VOCABS | {hashlib.sha256((ROOT / "data" / f).read_bytes()).hexdigest()
+                                      for f in DRAFT_VOCABS.values() if (ROOT / "data" / f).exists()}
+        if old not in shipped or old == hashlib.sha256(new.read_bytes()).hexdigest():
+            return
+        ok("draft layer: the token subset " + ("with Chinese, Japanese and Korean" if choice == "cjk" else
+                                               "for English and code (less VRAM)"))
+    shutil.copyfile(new, dst)
+
+
+def ensure_engine_for(cards, cfg_path: Path, cfg: dict, yes: bool) -> dict:
+    """The installed engine must have code for every card the model starts on: a card added later (--gpus with an
+    older or newer generation, #128) or a new GPU in the PC otherwise stops the start with 'no kernel image'.  Such a
+    card gets the engine compiled for all of them, before the start."""
+    missing = [g for g in cards if not engine_runs_on(g)]
+    if not missing:
+        return cfg
+    info = ROOT / "engine" / "BUILD.json"
+    meta = json.loads(info.read_text())
+    say()
+    say("  The installed engine has no code for " + ", ".join(f"{g['name']} (sm_{g['arch']})" for g in missing) +
+        ": it is compiled for " + ("these cards" if len(cards) > 1 else "it") + " now.")
+    main = gpu_info(cards[0]["index"])
+    archs = sorted({int(x) for x in meta.get("archs", [])} | {int(g["arch"]) for g in cards})
+    vision = meta.get("vision") or ("gpu" if (ROOT / "engine" / VEXE).exists() else "none")
+    build_engine({**main, "archs": archs}, vision, yes, get_llama_cpp())
+    dirs = json.loads(info.read_text()).get("cuda_dirs") or []
+    cfg["lib_dirs"] = dirs + [d for d in cfg.get("lib_dirs") or [] if d not in dirs]
+    cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
+    return cfg
+
+
 def write_run_script(model, cfg_path, port):
    serve = [sys.executable, str(ROOT / "serve" / "server.py"), "--engine", "strata", "--config", str(cfg_path),
             "--port", str(port), "--open"]
    if WIN:
        script = ROOT / f"run-{model.lower()}.bat"
        script.write_text("@echo off\r\ntitle Strata " + model + "\r\ncd /d \"" + str(ROOT) + "\"\r\n" +
-                          " ".join(f'"{x}"' for x in serve) + "\r\npause\r\n", encoding="utf-8")
+                          " ".join(f'"{x}"' for x in serve) + "\r\nif errorlevel 1 pause\r\n", encoding="utf-8")
    else:
        script = ROOT / f"run-{model.lower()}.sh"
        script.write_text("#!/bin/sh\ncd \"" + str(ROOT) + "\"\nexec " + " ".join(f'"{x}"' for x in serve) + "\n",
@@ -1456,6 +1620,12 @@ def main() -> int:
    ap.add_argument("--check", action="store_true", help="only check this PC and exit")
    ap.add_argument("--calibrate", action="store_true",
                    help="tune the engine's settings for this PC (about 5-10 minutes), then start the model")
+    ap.add_argument("--draft-vocab", choices=list(DRAFT_VOCABS),
+                    help="the draft layer's tokens: cjk = with Chinese, Japanese and Korean (default), en = English "
+                         "and code only (~110 MiB less VRAM, English answers 1-2%% faster)")
+    ap.add_argument("--low-ram", choices=["auto", "on", "off"], default="auto",
+                    help="map the model's experts from its folder instead of copying them into RAM (for a PC with a big "
+                         "GPU and little RAM); auto: when the experts would not fit the RAM")
    ap.add_argument("--backend", choices=["cuda", "hip"],
                    help="cuda = NVIDIA (default), hip = AMD RX 7900 XT/XTX on Linux (experimental; chosen by itself "
                         "when the PC has no NVIDIA card Strata can use)")
@@ -1512,29 +1682,44 @@ def main() -> int:
                say(f"  {i}) {json.loads(c.read_text(encoding='utf-8-sig')).get('model_name', c.stem)}")
            pick_cfg = have[int(ask("Tune which one?", [str(i) for i in range(1, len(have) + 1)], "1", a.yes)) - 1]
        calibrate_config(pick_cfg)
-        return 0 if a.no_start else start(pick_cfg, a.port, run_gpu, yes=a.yes, layer_split=a.layer_split)
+        return 0 if a.no_start else start(pick_cfg, a.port, run_gpu, yes=a.yes, layer_split=a.layer_split,
+                     keep={"host": a.host, "api_key": a.api_key, "draft_vocab": a.draft_vocab})
    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], a.port, run_gpu, yes=a.yes, layer_split=a.layer_split)
+            return start(have[0], a.port, run_gpu, yes=a.yes, layer_split=a.layer_split,
+                     keep={"host": a.host, "api_key": a.api_key, "draft_vocab": a.draft_vocab})
        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], a.port, run_gpu, yes=a.yes, layer_split=a.layer_split)
+            return start(have[pick - 1], a.port, run_gpu, yes=a.yes, layer_split=a.layer_split,
+                     keep={"host": a.host, "api_key": a.api_key, "draft_vocab": a.draft_vocab})

    # ---- 1. the PC
    step(1, "checking your PC")
    found = gpus()
    amd = [] if WIN else amd_gpus()
-    hip = a.backend == "hip" or (a.backend is None and not any(gpu_problem(g) is None for g in found)
-                                 and any(amd_problem(g) is None for g in amd))
+    nv_ok = any(gpu_problem(g) is None for g in found)
+    amd_ok = [g for g in amd if amd_problem(g) is None]
+    hip = a.backend == "hip" or (a.backend is None and not nv_ok and bool(amd_ok))
+    if a.backend is None and nv_ok and amd_ok:
+        # both kinds of card: asked (a first run on such a PC used to take NVIDIA without mentioning the Radeon)
+        say()
+        say("  This PC has NVIDIA and AMD cards Strata can use:")
+        say("  1) NVIDIA: " + ", ".join(f"{g['name']} ({g['vram_gb']:.0f} GB)" for g in found if gpu_problem(g) is None)
+            + "   (recommended)")
+        say("  2) AMD: " + ", ".join(f"{g['name']} ({g['vram_gb']:.0f} GB)" for g in amd_ok)
+            + "   (experimental: compiled here, one GPU, no images - docs/AMD_HIP.md)")
+        hip = ask("Which cards?", ["1", "2"], "1", a.yes or a.check) == "2"
+        if a.check and not hip:
+            say("  (the AMD card: ./setup.sh --backend hip)")
    if hip:                                            # AMD (experimental): one card, compiled here
        if WIN:
-            fail("Strata's AMD backend runs on Linux only", "use an NVIDIA RTX 30 series or newer card on Windows")
+            fail("Strata's AMD backend runs on Linux only", "use an NVIDIA RTX 20 series or newer card on Windows")
        say("  Your AMD GPUs:" if amd else "  No AMD GPU found (the amdgpu driver's KFD topology is empty).")
        for g in amd:
            say(f"    GPU {g['index']}: {g['name']}, {g['vram_gb']:.0f} GB VRAM - " + (amd_problem(g) or "can be used"))
@@ -1578,13 +1763,15 @@ def main() -> int:
    ram = ram_gb()
    cpu, avx2, avx512 = cpu_info()
    need = min(d["ram_gb"] for d in MODELS.values())
-    if ram < need - 4 and not a.check:
+    low_ok = low_ram_fits("IQ1_M", ram, gpu["vram_gb"]) and a.low_ram != "off"   # the smallest model, mapped
+    if ram < need - 4 and not a.check and not low_ok:
        # 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 (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"RAM: {ram:.0f} GB" if ram >= need - 4 else f"RAM: {ram:.0f} GB (less than the {need} GB the smallest model needs)"
+       + ("; the GPU's VRAM makes up for it (the low-RAM mode)" if ram < need - 4 and low_ok else ""))
    pf = page_file_gb()
    if pf is not None and pf < 4:
        warn(f"Windows' page file is {pf:.1f} GB: the graphics card's memory needs room there too (issue #60), so "
@@ -1597,6 +1784,9 @@ def main() -> int:
        say()
        for m, d in MODELS.items():
            verdict = "fits" if ram >= d["ram_gb"] else "tight" if ram >= d["ram_gb"] - 8 else "does not fit"
+            if low_ram_needed(m, ram) and low_ram_fits(m, ram, gpu["vram_gb"]) and a.low_ram != "off":
+                verdict = (f"fits in the low-RAM mode (the GPU holds ~{100 * low_ram_gpu_share(m, gpu['vram_gb']):.0f}% "
+                           "of its experts)")
            say(f"  {m:8s} needs ~{d['ram_gb']} GB RAM: {verdict}")
        say("\nThis PC can run Strata. Run it again without --check to install.")
        return 0
@@ -1622,10 +1812,23 @@ def main() -> int:
    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}"
+        if low_ram_needed(m, ram) and low_ram_fits(m, ram, gpu["vram_gb"]) and a.low_ram != "off":
+            fit = f"   <- fits in the low-RAM mode (the GPU holds ~{100 * low_ram_gpu_share(m, gpu['vram_gb']):.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 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:
+    low_ram = a.low_ram == "on" or (a.low_ram == "auto" and low_ram_needed(model, ram))
+    if low_ram and multi:
+        warn("the low-RAM mode runs on one GPU: using " + gpu_name(gpu) + " only")
+        multi, sel, chosen = [], [gpu["index"]], [gpu]
+    if low_ram:
+        share = low_ram_gpu_share(model, gpu["vram_gb"])
+        ok(f"low-RAM mode: {model}'s experts ({MODELS[model]['arena_gb']:.0f} GB) are read from the model folder "
+           f"through the OS file cache instead of a copy in RAM ({ram:.0f} GB); the GPU holds ~{100 * share:.0f}% of them")
+        if share < 0.6:
+            warn("most of the experts are read from the SSD while it answers: expect it to be much slower than with "
+                 "enough RAM (a faster SSD and a smaller size help)")
+    elif ram < MODELS[model]["ram_gb"] - 4:
        # #125: a warning and a question, not a stop: the user may accept paging (asked, "no" by default, so an
        # unattended --yes install still stops here)
        need_gb, arena = MODELS[model]["ram_gb"], MODELS[model]["arena_gb"]
@@ -1706,9 +1909,13 @@ def main() -> int:
                models_dir, shards = cand[0].parent, cand
                ok(f"model files found in {models_dir}")
                break
+    for s in shards:                                   # #173: a whole file copied in by hand has no finish mark
+        if s.exists() and not done(s) and whole_shard(s):
+            mark(s, "whole (checked against its own tensor directory)")
    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)
+        (40 if model == "Q2_0" and avx512 and family == "qwen" else 0) + (1 if vision != "none" else 0) + \
+        (MODELS[model]["arena_gb"] + 1 if low_ram and not (model == "Q2_0" and avx512 and family == "qwen") else 0)
    if free_gb(models_dir) < need:
        fail(f"not enough free disk space in {models_dir}: need ~{need:.0f} GB", "use --models-dir on a bigger drive")

@@ -1726,6 +1933,12 @@ def main() -> int:
        if vision != "none" and not (eng / VEXE).exists():
            warn("the ready-made engine has no image encoder: compiling it")
            eng = None
+        elif vision == "gpu":                          # the encoder can cover fewer cards than the engine (RTX 20)
+            m = json.loads((eng / "BUILD.json").read_text())
+            va = [int(x) for x in m.get("vision_archs", m.get("archs", []))]
+            if va and int(gpu["arch"]) not in va and not (m.get("ptx") and int(gpu["arch"]) > max(va)):
+                warn(f"the ready-made image encoder has no code for your GPU (sm_{gpu['arch']}): compiling it")
+                eng = None
    if eng is None:
        eng = build_engine_hip(gpu, llama) if hip else build_engine(gpu, vision, a.yes, llama)
    meta = json.loads((eng / "BUILD.json").read_text())
@@ -1779,6 +1992,10 @@ def main() -> int:
    elif not (pack / "native_experts.txt").exists() or not (pack / "tokenizer" / "vocab.json").exists():
        # 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)
+    if low_ram and not (pack / "experts.bin").exists():
+        say(f"  Writing the experts into one file for the low-RAM mode (one time, {MODELS[model]['arena_gb']:.0f} GB) ...")
+        run([sys.executable, str(ROOT / "tools" / "iq_pack.py"), "--gguf", str(shards[0]), "--out", str(pack),
+             "--experts-bin"], env=env)
    ok(f"model prepared: {pack}")
    mtp = (find_in(roots, "mtp/rt/experts.bin") or data / "mtp/rt/experts.bin").parent.parent
    rt = mtp / "rt"
@@ -1790,8 +2007,7 @@ def main() -> int:
             "--out", str(mtp / "mtp-q2_0.gguf")], env=env)
        run([sys.executable, str(ROOT / "tools" / "mtp_rt.py"), "--gguf", str(mtp / "mtp-q2_0.gguf"), "--out", str(rt)],
            env=env)
-    if not (rt / "draft_vocab.bin").exists():
-        shutil.copyfile(ROOT / "data" / "draft_vocab.bin", rt / "draft_vocab.bin")
+    refresh_draft_vocab(rt, a.draft_vocab or "cjk")
    ok(f"MTP draft layer: {rt}")

    # ---- 7. the start script
@@ -1807,6 +2023,8 @@ def main() -> int:
            "--max-context", str(ctx)]
    if ctx > 8192:
        args += ["--kv", kv]
+    if low_ram:
+        args += ["--mmap-experts"]   # the experts from the pack's experts.bin, not copied into RAM
    # KV streaming: from 64K up the whole KV cache lives in RAM and only the part the attention reads (32K positions
    # 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.
@@ -1849,6 +2067,8 @@ def main() -> int:
        cfg["host"] = a.host
    if a.api_key:
        cfg["api_key"] = a.api_key
+    if a.draft_vocab:
+        cfg["draft_vocab"] = a.draft_vocab
    if vision != "none":
        cfg["vision"] = {"exe": str(eng / VEXE), "mmproj": str(mmproj), "model": str(shards[0]),
                         "gpu": vision == "gpu", "max_tokens": VISION[vision]["max_tokens"]}
@@ -33,7 +33,7 @@ DeviceInfo device_info(int ordinal) {
 #if defined(STRATA_USE_HIP)
        throw CudaError("no HIP device is present; this backend targets gfx1100 wave32", -1);
 #else
-        throw CudaError("no CUDA device is present; Strata needs an NVIDIA GPU of compute capability 8.0 or newer", -1);
+        throw CudaError("no CUDA device is present; Strata needs an NVIDIA GPU (RTX 20 series or newer)", -1);
 #endif
    }
    if (ordinal < 0 || ordinal >= count) {
@@ -60,19 +60,20 @@ DeviceInfo device_info(int ordinal) {
    check(cudaDriverGetVersion(&d.driver_version), "cudaDriverGetVersion");
    check(cudaRuntimeGetVersion(&d.runtime_version), "cudaRuntimeGetVersion");

-    // The kernels need sm_80 or newer (tf32 mma in the attention scorer, bf16 math) - the same floor the
-    // arch guard in CMakeLists enforces at build time (RTX 30 / 40 / 50).  Anything older is caught here,
-    // because a binary can be carried to a machine with an older card and would otherwise silently take
-    // whatever path the driver chose.  The HIP backend is validated on gfx1100 (wave32) only.
+    // The engine supports compute capability 7.5 and newer (Turing: the QSA scorer's tf32 mma has a portable
+    // fp32-FMA fallback below sm_80, the tensor-core prompt kernels refuse and fall back).  Compiling for a
+    // supported arch is enforced by CMake; RUNNING on an older card is caught here, because a binary can be carried
+    // to a machine with an older card and would otherwise silently take whatever path the driver chose.  The HIP
+    // backend is validated on gfx1100 (wave32) only.
 #if defined(STRATA_USE_HIP)
    if (std::strncmp(p.gcnArchName, "gfx1100", 7) != 0 || p.warpSize != 32) {
        throw CudaError("HIP backend requires validated gfx1100 wave32 hardware", -1);
    }
 #else
-    if (d.cc_major < 8) {
+    if (d.cc_major * 10 + d.cc_minor < 75) {
        throw CudaError("device " + d.name + " reports compute capability " + std::to_string(d.cc_major) +
                            "." + std::to_string(d.cc_minor) +
-                            "; Strata needs an NVIDIA GPU of compute capability 8.0 or newer (RTX 30 / 40 / 50)",
+                            "; Strata needs compute capability 7.5 or newer (RTX 20 / 30 / 40 / 50 series)",
                        -1);
    }
 #endif
@@ -292,6 +292,19 @@ bool MtpDrafter::load(const std::string& rt_dir, const ModelGeometry& g, Session
    return true;
 }

+uint64_t MtpDrafter::bind_bytes(uint64_t head_row_bytes, int64_t n_vocab) const {
+    uint64_t bytes = head_logits_ ? 0 : (uint64_t) max_t_ * (uint64_t) n_vocab * sizeof(float);
+    if (dhead_ == nullptr) {
+        if (FILE* f = std::fopen((rt_dir_ + "/draft_vocab.bin").c_str(), "rb")) {
+            std::fseek(f, 0, SEEK_END);
+            const long size = std::ftell(f);
+            std::fclose(f);
+            if (size >= 4 && size % 4 == 0) bytes += (uint64_t) (size / 4) * head_row_bytes + (uint64_t) size;
+        }
+    }
+    return bytes;
+}
+
 bool MtpDrafter::bind(const WeightTable& wt, const NativeHead* head, const float* window_R, std::string& err) {
    const OnDevice on_device(device_);
    wt_ = &wt;
@@ -148,7 +148,8 @@ bool Verifier::init(const WeightTable& wt, const ModelGeometry& g, SessionState&
        return false;
    }
    if (hits.d_res == nullptr || hits.cache_base == nullptr || hits.blob <= 0) {
-        err = "verify: needs the profile-filled VRAM expert tier (--expert-profile and --expert-cache)";
+        err = "verify: needs the profile-filled VRAM expert tier (--expert-profile and --expert-cache); with "
+              "--expert-cache auto, no VRAM was left for it: lower --max-context, use --kv k8v4 or images on the CPU";
        return false;
    }
    std::string why;
@@ -850,8 +851,10 @@ bool Verifier::capture(int T, std::string& err) {
            if (ty == cudaGraphNodeTypeKernel) {
                cudaKernelNodeParams kp{};
                if (cudaGraphKernelNodeGetParams(nd, &kp) == cudaSuccess) {
+#if CUDART_VERSION >= 12030   // cudaFuncGetName arrived in CUDA 12.3
                    const char* fn = nullptr;
                    if (cudaFuncGetName(&fn, kp.func) == cudaSuccess && fn) name = fn;
+#endif
                }
            } else if (ty == cudaGraphNodeTypeMemcpy) name = "memcpy";
            else if (ty == cudaGraphNodeTypeMemset) name = "memset";
@@ -325,19 +325,40 @@ void fused_gr_read_multi(const FusedGrArgs* a, int n_tok, float* xn_scratch, voi
    // the shared-memory opt-in is a per-DEVICE setting: once per device, not once per process (a layer split
    // runs this kernel on two cards)
    static bool attr[64] = {};
+    static int chunk[64] = {};   // Turing port: tokens the down kernel may carry in one launch on this card
    int dev = 0;
    cudaGetDevice(&dev);
-    if (dev < 0 || dev >= 64 || !attr[dev]) {
+    if (dev >= 0 && dev < 64 && !attr[dev]) {
        // at most what the card allows (Turing: 64 KB - enough for windows of up to 6 tokens)
        int optin = 0;
        cudaDeviceGetAttribute(&optin, cudaDevAttrMaxSharedMemoryPerBlockOptin, dev);
        int want = (int) (kFusedGrMaxT * TILE * sizeof(float));
        if (optin > 0 && want > optin) want = optin;
        cudaFuncSetAttribute(gr_down_multi_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, want);
-        cudaGetLastError();
-        if (dev >= 0 && dev < 64) attr[dev] = true;
+        cudaGetLastError();      // drop any error the attempt left behind
+        // Turing port: the down kernel stages n_tok*TILE floats of dynamic shared memory - 80 KB at the full
+        // 8 tokens.  A card whose opt-in is below that (Turing: 64 KB, so 7+ tokens fail to launch as
+        // "invalid argument") processes the tokens in slices that fit; a card that reports no opt-in gets what
+        // fits the 48 KB default (4 tokens of the CUDA tile; all 8 of HIP's smaller tile).  The down kernel's
+        // outputs (lo, inject_out) are strictly per-token, so the chunk boundaries are safe, and the up kernel
+        // below still sees every token of the batch in one launch.
+        const int capacity = (optin > 0 ? optin : 48 * 1024) / (int) (TILE * sizeof(float));
+        chunk[dev] = capacity < 1 ? 1 : (capacity > kFusedGrMaxT ? kFusedGrMaxT : capacity);
+        attr[dev] = true;
+    }
+    const int chunk_tok = (dev >= 0 && dev < 64 && chunk[dev]) ? chunk[dev] : kFusedGrMaxT;
+    if (chunk_tok >= n_tok) {
+        gr_down_multi_kernel<<<DOWN_BLOCKS + 1, THREADS, (size_t) n_tok * TILE * sizeof(float), st>>>(m);
+    } else {
+        for (int c0 = 0; c0 < n_tok; c0 += chunk_tok) {
+            const int ct = n_tok - c0 < chunk_tok ? n_tok - c0 : chunk_tok;
+            GrMulti c{};
+            c.xn = xn_scratch + (size_t) c0 * D;
+            c.T = ct;
+            for (int k = 0; k < ct; ++k) c.a[k] = a[c0 + k];
+            gr_down_multi_kernel<<<DOWN_BLOCKS + 1, THREADS, (size_t) ct * TILE * sizeof(float), st>>>(c);
+        }
    }
-    gr_down_multi_kernel<<<DOWN_BLOCKS + 1, THREADS, (size_t) n_tok * TILE * sizeof(float), st>>>(m);
    if (stamp_buf) gpu_stamp(stamp_buf, stamp_i0 + 1, stream);
    gr_up_multi_kernel<<<UPM_BLOCKS, THREADS, 0, st>>>(m);
    const cudaError_t e = cudaGetLastError();
@@ -5,7 +5,7 @@

 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
-#include <math_constants.h>
+#include <cmath>

 #include <cstdio>
 #include <cstdlib>
@@ -117,7 +117,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_kernel(const float* __res
        S.qh[h][d] = hi;
        S.ql[h][d] = __float2half_rn(x - __half2float(hi));
    }
-    if (t < 16) { S.mrow[t] = -CUDART_INF_F; S.lsum[t] = 0.0f; }
+    if (t < 16) { S.mrow[t] = -INFINITY; S.lsum[t] = 0.0f; }

    float acc[8][4];
 #pragma unroll
@@ -244,17 +244,17 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_kernel(const float* __res
                sc[3] = fmaf(tg[3], s1, sc[3]);
            }
            const int c = cb + 2 * tig;
-            S.s[gid][c] = c < nh ? sc[0] * qdown : -CUDART_INF_F;
-            S.s[gid][c + 1] = c + 1 < nh ? sc[1] * qdown : -CUDART_INF_F;
-            S.s[gid + 8][c] = c < nh ? sc[2] * qdown : -CUDART_INF_F;
-            S.s[gid + 8][c + 1] = c + 1 < nh ? sc[3] * qdown : -CUDART_INF_F;
+            S.s[gid][c] = c < nh ? sc[0] * qdown : -INFINITY;
+            S.s[gid][c + 1] = c + 1 < nh ? sc[1] * qdown : -INFINITY;
+            S.s[gid + 8][c] = c < nh ? sc[2] * qdown : -INFINITY;
+            S.s[gid + 8][c + 1] = c + 1 < nh ? sc[3] * qdown : -INFINITY;
        }
        __syncthreads();
        // online softmax: row t/8, 4 cells per thread, 8 threads per row (lanes 8r..8r+7 of a warp)
        {
            constexpr int PER = CH / 8;
            const int r = t >> 3, sub = t & 7;
-            float x[PER], mx = -CUDART_INF_F;
+            float x[PER], mx = -INFINITY;
 #pragma unroll
            for (int j = 0; j < PER; ++j) { x[j] = S.s[r][sub * PER + j]; mx = fmaxf(mx, x[j]); }
 #pragma unroll
@@ -264,7 +264,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_kernel(const float* __res
            float sum = 0.0f;
 #pragma unroll
            for (int j = 0; j < PER; ++j) {
-                const float e = x[j] == -CUDART_INF_F ? 0.0f : exp2f(x[j] - m_new);
+                const float e = x[j] == -INFINITY ? 0.0f : exp2f(x[j] - m_new);
                S.s[r][sub * PER + j] = e;
                sum += e;
            }
@@ -272,7 +272,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_kernel(const float* __res
            for (int o = 1; o < 8; o <<= 1) sum += __shfl_xor_sync(0xffffffffu, sum, o);
            __syncwarp();
            if (sub == 0) {
-                const float a = m_old == -CUDART_INF_F ? 0.0f : exp2f(m_old - m_new);
+                const float a = m_old == -INFINITY ? 0.0f : exp2f(m_old - m_new);
                S.alpha[r] = a;
                S.lsum[r] = fmaf(S.lsum[r], a, sum);
                S.mrow[r] = m_new;
@@ -419,7 +419,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_i8_kernel(const float* __
 #pragma unroll
    for (int o = 16; o > 0; o >>= 1) qm = fmaxf(qm, __shfl_xor_sync(0xffffffffu, qm, o));
    if (lane == 0) S.qmax[warp] = qm;
-    if (t < 16) { S.mrow[t] = -CUDART_INF_F; S.lsum[t] = 0.0f; }
+    if (t < 16) { S.mrow[t] = -INFINITY; S.lsum[t] = 0.0f; }
    __syncthreads();
    qm = fmaxf(fmaxf(S.qmax[0], S.qmax[1]), fmaxf(S.qmax[2], S.qmax[3]));
    int qe = 0;
@@ -511,12 +511,12 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_i8_kernel(const float* __
        // online softmax over the four groups' sum (fixed order): row t/8, 4 cells per thread
        {
            const int r = t >> 3, sub = t & 7;
-            float x[4], mx = -CUDART_INF_F;
+            float x[4], mx = -INFINITY;
 #pragma unroll
            for (int j = 0; j < 4; ++j) {
                const int c = sub * 4 + j;
                x[j] = c0 + c < n ? (((S.part[0][r][c] + S.part[1][r][c]) + S.part[2][r][c]) + S.part[3][r][c]) * qdown
-                                  : -CUDART_INF_F;
+                                  : -INFINITY;
                mx = fmaxf(mx, x[j]);
            }
 #pragma unroll
@@ -526,7 +526,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_i8_kernel(const float* __
            float sum = 0.0f;
 #pragma unroll
            for (int j = 0; j < 4; ++j) {
-                const float e = x[j] == -CUDART_INF_F ? 0.0f : exp2f(x[j] - m_new);
+                const float e = x[j] == -INFINITY ? 0.0f : exp2f(x[j] - m_new);
                S.p[r][sub * 4 + j] = e;
                sum += e;
            }
@@ -534,7 +534,7 @@ __global__ void __launch_bounds__(THREADS) prompt_attn_i8_kernel(const float* __
            for (int o = 1; o < 8; o <<= 1) sum += __shfl_xor_sync(0xffffffffu, sum, o);
            __syncwarp();
            if (sub == 0) {
-                const float a = m_old == -CUDART_INF_F ? 0.0f : exp2f(m_old - m_new);
+                const float a = m_old == -INFINITY ? 0.0f : exp2f(m_old - m_new);
                S.alpha[r] = a;
                S.lsum[r] = fmaf(S.lsum[r], a, sum);
                S.mrow[r] = m_new;
@@ -93,6 +93,7 @@ bool supported(int t) {
    switch ((ggml_type) t) {
        case GGML_TYPE_Q2_0: case GGML_TYPE_IQ2_XXS: case GGML_TYPE_IQ2_XS: case GGML_TYPE_IQ2_S:
        case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ3_S: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS:
+        case GGML_TYPE_Q8_0:   // the draft layer's dense matrices (E-9)
            return true;
        default:
            return false;
@@ -141,6 +142,7 @@ void Context::run(const Product& p, void* stream) {
        case GGML_TYPE_IQ3_S: mul_mat_q_case<GGML_TYPE_IQ3_S>(ctx, a, s); break;
        case GGML_TYPE_IQ4_NL: mul_mat_q_case<GGML_TYPE_IQ4_NL>(ctx, a, s); break;
        case GGML_TYPE_IQ4_XS: mul_mat_q_case<GGML_TYPE_IQ4_XS>(ctx, a, s); break;
+        case GGML_TYPE_Q8_0: mul_mat_q_case<GGML_TYPE_Q8_0>(ctx, a, s); break;
        default:
            std::fprintf(stderr, "prefill mmq: type %d is not covered\n", (int) t);
            std::exit(1);
@@ -1,5 +1,6 @@
 // src/prefill/prefill.cpp - see include/strata/prefill/prefill.hpp.
 #include "strata/prefill/prefill.hpp"
+#include "strata/core/mtp.hpp"
 #include "strata/core/progress.hpp"
 #include "strata/core/on_device.hpp"

@@ -673,6 +674,151 @@ bool Prefill::relayout(int64_t chunk, void* borrow, uint64_t borrow_bytes, std::
 }

 int64_t Prefill::chunk() const { return impl_->T; }
+
+bool Prefill::draft_kv(core::MtpDrafter& mtp, const float* R_rows, const int32_t* next_tokens, int64_t n, int64_t cell0,
+                       std::string& err) {
+    Impl& m = *impl_;
+    static const bool off = [] { const char* v = std::getenv("STRATA_MTP_BATCH"); return v != nullptr && v[0] == '0'; }();
+    core::QsaState& st = mtp.kv_state_rw();
+    if (off || n <= 0 || m.g == nullptr || m.region == nullptr || st.kv_mode != 0 || st.kv_hybrid ||
+        mtp.device() != m.device)
+        return false;
+    const auto t0 = Clock::now();
+    const core::ModelGeometry& g = *m.g;
+    const int64_t Nn = g.n_embd, HCN = g.hc * g.n_embd, KV = g.n_head_kv * g.head_dim;
+    constexpr int kQ8_0 = 8;   // GGML_TYPE_Q8_0
+    const float* w_ne = mtp.tensor_f32("pre_fc_norm_embedding.weight");
+    const void* w_fe = mtp.tensor_q8("fc_embedding.weight");
+    const float* w_nh = mtp.tensor_f32("pre_fc_norm_hidden.weight");
+    const void* w_fh = mtp.tensor_q8("fc_hidden.weight");
+    const float* w_hn = mtp.tensor_f32("attn_hyper_connection.hc_norm.weight");
+    const uint16_t* w_dn = mtp.tensor_bf16("attn_hyper_connection.input_mix_weight_down.weight");
+    const uint16_t* w_up = mtp.tensor_bf16("attn_hyper_connection.input_mix_weight_up.weight");
+    const void* w_k = mtp.tensor_q8("self_attn.k_proj.weight");
+    const void* w_v = mtp.tensor_q8("self_attn.v_proj.weight");
+    const float* w_kn = mtp.tensor_f32("self_attn.k_norm.weight");
+    if (!w_ne || !w_fe || !w_nh || !w_fh || !w_hn || !w_dn || !w_up || !w_k || !w_v || !w_kn) return false;
+    const core::NativeEmbed* nemb = core::native_embed();
+    const core::WeightRef* wemb = nemb ? nullptr : m.wt->find("token_embd.weight");
+    if (!nemb && (wemb == nullptr || wemb->codebook_iq4nl || wemb->ne0 != g.n_embd || wemb->group_elems <= 0 ||
+                  (wemb->code_bits != 2 && wemb->code_bits != 4 && wemb->code_bits != 8)))
+        return false;
+    // the cells the drafter's window can still reach
+    const int64_t r0 = std::max<int64_t>(0, mtp.first_needed() - cell0);
+    if (r0 >= n) return true;
+    // per row: emb/e2 (N), en16 (N half), hn/h2/Rm/gated (HCN), hn16/xn16 (HCN half), lo (LR) + lo16, grs, mixed (N) +
+    // mixed_h, K and V (KV each), the token id
+    // E-9: the drafter's Q8_0 matrices through Q8_1 x Q8_0 MMQ - its own pass's integer dot products (mmvq), so
+    // its K/V stay close to what the drafter computes itself; STRATA_MTP_BATCH_F16=1: FP16 GEMMs (the A/B)
+    static const bool f16_only = [] { const char* v = std::getenv("STRATA_MTP_BATCH_F16"); return v && v[0] == '1'; }();
+    const bool q8 = !f16_only && mmq::built() && mmq::supported(kQ8_0);
+    const uint64_t per_row = 4 * (2 * Nn + 4 * HCN + LR + HC + Nn + 2 * KV + 1) + 2 * (Nn + 2 * HCN + LR + Nn) + 64 +
+                             (q8 ? (uint64_t) mmq::q8_bytes(g.hc, Nn) + 4 * g.hc : 0);
+    const int64_t B = std::min<int64_t>(n - r0, (int64_t) (m.region_bytes / per_row) & ~(int64_t) 63);
+    if (B < 64) return false;
+    uint8_t* q = m.region;
+    auto carve = [&](size_t bytes) { void* p = q; q += (bytes + 255) & ~(size_t) 255; return p; };
+    float* emb = (float*) carve((size_t) B * Nn * 4);
+    float* e2 = (float*) carve((size_t) B * Nn * 4);
+    uint16_t* en16 = (uint16_t*) carve((size_t) B * Nn * 2);
+    float* hn = (float*) carve((size_t) B * HCN * 4);
+    float* h2 = (float*) carve((size_t) B * HCN * 4);
+    float* Rm = (float*) carve((size_t) B * HCN * 4);
+    float* gated = (float*) carve((size_t) B * HCN * 4);
+    uint16_t* hn16 = (uint16_t*) carve((size_t) B * HCN * 2);
+    uint16_t* xn16 = (uint16_t*) carve((size_t) B * HCN * 2);
+    float* lo = (float*) carve((size_t) B * LR * 4);
+    uint16_t* lo16 = (uint16_t*) carve((size_t) B * LR * 2);
+    float* grs = (float*) carve((size_t) B * HC * 4);
+    float* mixed = (float*) carve((size_t) B * Nn * 4);
+    uint16_t* mixed_h = (uint16_t*) carve((size_t) B * Nn * 2);
+    float* Kc = (float*) carve((size_t) B * KV * 4);
+    float* Vc = (float*) carve((size_t) B * KV * 4);
+    int32_t* tok = (int32_t*) carve((size_t) B * 4);
+    void* xq = q8 ? carve(mmq::q8_bytes(B * g.hc, Nn)) : nullptr;
+    int32_t* ident = q8 ? (int32_t*) carve((size_t) B * g.hc * 4) : nullptr;
+    int32_t* bnd = q8 ? (int32_t*) carve(16) : nullptr;
+    if ((uint64_t) (q - m.region) > m.region_bytes) return false;
+    if (!mtp.idle(err)) return false;   // the drafter's own stream (its graph uploads) before this writes its K/V
+    strata::kernels::QsaShapes s = strata::kernels::qsa_real_shapes();
+    s.n_head = g.n_head; s.n_head_kv = g.n_head_kv; s.head_dim = g.head_dim; s.idx_n_head = g.idx_q_heads;
+    s.idx_dim = g.idx_key_dim;
+    std::vector<int32_t> tk((size_t) B);
+    if (q8) {
+        if (!m.mmq_ctx) m.mmq_ctx = std::make_unique<mmq::Context>();
+        mmq::iota(ident, B * g.hc, m.cs);
+    }
+    // y[rows, n_out] = x[rows, k] . w^T for a Q8_0 matrix: MMQ from the FP32 rows (bounds slot 0: rows, 1: rows*hc)
+    auto proj = [&](const float* x, const uint16_t* x16, const void* w, float* y, int64_t rows, int64_t n_out,
+                    int64_t k, int slot) {
+        if (!q8) { m.gemm.native(x16, kQ8_0, w, y, rows, n_out, k); return; }
+        mmq::quantize(x, nullptr, xq, kQ8_0, k, k, rows, m.cs);
+        mmq::Product p;
+        p.w = w; p.type = kQ8_0; p.w_rows = n_out; p.w_cols = k; p.expert_bytes = mmq::matrix_bytes(kQ8_0, n_out, k);
+        p.n = 1; p.xq = xq; p.bounds = bnd + 2 * slot; p.ids = ident; p.total_rows = rows; p.max_rows = rows;
+        p.dst = y; p.ld_dst = n_out;
+        m.mmq_ctx->run(p, m.cs);
+    };
+    for (int64_t b0 = r0; b0 < n; b0 += B) {
+        const int64_t nb = std::min(B, n - b0), c0 = cell0 + b0;
+        for (int64_t i = 0; i < nb; ++i) tk[(size_t) i] = next_tokens[b0 + i];
+        const int32_t bh[4] = {0, (int32_t) nb, 0, (int32_t) (nb * g.hc)};
+        if (q8 && cudaMemcpyAsync(bnd, bh, sizeof bh, cudaMemcpyHostToDevice, m.cs) != cudaSuccess) {
+            err = "prefill: the draft bounds' upload failed";
+            return false;
+        }
+        if (cudaMemcpyAsync(tok, tk.data(), (size_t) nb * 4, cudaMemcpyHostToDevice, m.cs) != cudaSuccess) {
+            err = "prefill: the draft tokens' upload failed";
+            return false;
+        }
+        // the input branches: the next token's embedding, and this cell's final residual rows
+        if (nemb) {
+            nemb->gather_dev(tok, nb, emb, m.cs);
+        } else {
+            const auto* codes = (const uint8_t*) wemb->data;
+            const auto* scales = (const float*) (codes + wemb->codes_bytes);
+            const auto* offsets = wemb->has_offset ? (const float*) (codes + wemb->codes_bytes + wemb->scales_bytes)
+                                                   : nullptr;
+            strata::kernels::embedding_gather_dev(codes, scales, offsets, tok, (int) nb, wemb->ne0, wemb->code_bits,
+                                                  wemb->code_bias, wemb->group_elems,
+                                                  (uint64_t) (wemb->ne0 / (8 / wemb->code_bits)),
+                                                  (uint64_t) (wemb->ne0 / wemb->group_elems), emb, m.cs);
+        }
+        rms_rows(emb, w_ne, nb, Nn, Nn, EPS, m.cs);
+        if (!q8) to_f16(emb, en16, nb * Nn, m.cs);
+        proj(emb, en16, w_fe, e2, nb, Nn, Nn, 0);
+        cudaMemcpyAsync(hn, R_rows + (size_t) b0 * HCN, (size_t) nb * HCN * 4, cudaMemcpyDeviceToDevice, m.cs);
+        rms_rows(hn, w_nh, nb, HCN, HCN, EPS, m.cs);
+        if (!q8) to_f16(hn, hn16, nb * HCN, m.cs);
+        proj(hn, hn16, w_fh, h2, nb * g.hc, Nn, Nn, 1);   // every stream through fc_hidden
+        strata::kernels::add_streams_broadcast(h2, e2, Rm, Nn, (int) g.hc, (int) nb, m.cs);
+        // the attention hyper-connection's read (its mixed input only: this pass writes nothing back)
+        gr_norm_rs(Rm, w_hn, EPS, grs, xn16, nb, m.cs);
+        m.gemm.bf16(xn16, w_dn, lo, nb, LR, HCN);
+        gr_silu(lo, lo16, nb, m.cs);
+        m.gemm.bf16(lo16, w_up, gated, nb, HCN, LR);
+        gr_mix_r(Rm, grs, w_hn, gated, mixed, nullptr, nb, m.cs, mixed_h);
+        // K and V into the drafter's cache, as the prompt path's QSA layers append theirs
+        proj(mixed, mixed_h, w_k, Kc, nb, KV, Nn, 0);
+        proj(mixed, mixed_h, w_v, Vc, nb, KV, Nn, 0);
+        rms_rows(Kc, w_kn, nb * g.n_head_kv, g.head_dim, g.head_dim, EPS, m.cs);
+        rope(Kc, nb, g.n_head_kv, g.head_dim, KV, c0, (float) strata::kernels::qsa_freq_base(), m.cs);
+        if (st.kv_q4) {
+            strata::kernels::fwht256_inplace_cuda(Kc, nb * g.n_head_kv, m.cs);
+            strata::kernels::fwht256_inplace_cuda(Vc, nb * g.n_head_kv, m.cs);
+            strata::kernels::kv_append_q4(st.k_q4, st.v_q4, st.page_table, c0, nb, Kc, Vc, s, m.cs, &st.host);
+        } else {
+            kv_append(Kc, Vc, nb, c0, st.page_table, s.page_size, st.kv_int8 ? nullptr : st.k_pool,
+                      st.kv_int8 ? nullptr : st.v_pool, st.k_q, st.v_q, st.k_scale, st.v_scale, m.cs, &st.host);
+        }
+    }
+    if (cudaStreamSynchronize(m.cs) != cudaSuccess) {
+        err = std::string("prefill: the draft layer's K/V: ") + cudaGetErrorString(cudaGetLastError());
+        return false;
+    }
+    mtp.ms_prefill += ms_since(t0);
+    return true;
+}
 void Prefill::set_pinned_share(double share) { g_pinned_share = share; }
 double Prefill::pinned_share() { return g_pinned_share; }

@@ -781,6 +927,7 @@ struct PfTimer {
 }  // namespace

 bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& err) {
+    err.clear();
    Impl& m = *impl_;
    const core::OnDevice on_device(m.device);
    const core::ModelGeometry& g = *m.g;
@@ -81,6 +81,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
+#include <filesystem>
 #include <fstream>
 #include <sstream>
 #include <string>
@@ -1763,9 +1764,14 @@ int main(int argc, char** argv) {
    strata::core::ArenaExpertSource arena_src;
    strata::core::ExpertSource* srcp = nullptr;
    if (o.mmap_experts) {
-        if (native_pack) {   // FileExpertSource maps the canonical pack's experts.bin; a native pack has none
-            std::fprintf(stderr, "strata generate: --mmap-experts needs a canonical pack (experts.bin); %s is a native "
-                                 "(IQ) pack, whose experts are loaded into the arena\n", o.pack.c_str());
+        // FileExpertSource maps the pack's experts.bin: a canonical pack has it; a native (IQ) pack has it when
+        // built with `tools/iq_pack.py --experts-bin` (the per-layer blob sizes of its layout, PR #121).  The low-RAM
+        // mode: the experts come from the file through the OS cache instead of a pinned copy in RAM, for a PC whose
+        // GPU holds most of them but whose RAM cannot hold them all.
+        if (native_pack && !std::filesystem::exists(std::filesystem::path(o.pack) / "experts.bin")) {
+            std::fprintf(stderr, "strata generate: --mmap-experts needs the pack's experts.bin; %s is a native (IQ) pack "
+                                 "built without it: python tools/iq_pack.py --gguf <shard 1> --out %s --experts-bin\n",
+                         o.pack.c_str(), o.pack.c_str());
            return 2;
        }
        if (!src.open(o.pack, g.n_layers, g.n_expert, err)) {
@@ -1987,12 +1993,20 @@ int main(int argc, char** argv) {
        // (with borrowing - the default with a profile - the prompt path lends cache slots instead)
        const bool borrow = !o.no_prefill_borrow && !o.expert_profile.empty();
        const int64_t prefill_mib = (o.prefill_chunk > 0 && !borrow) ? 160 + (o.prefill_chunk * 680) / 1024 : 0;
-        const int64_t reserve = ((int64_t) o.vram_reserve_mib + prefill_mib) << 20;
+        // the draft layer's head and logits are allocated when it binds, after this: 0.1.27's CJK subset made them
+        // ~110-180 MiB larger, and out of the reserve they left 16 GB cards below the stall line (#199)
+        const int64_t mtp_bind = (!o.mtp.empty() && native_head.loaded())
+                                     ? (int64_t) mtp.bind_bytes(native_head.row_bytes(), n_vocab) : 0;
+        const int64_t reserve = (((int64_t) o.vram_reserve_mib + prefill_mib) << 20) + mtp_bind;
        int64_t slots = ((int64_t) free_b - reserve) / (int64_t) strata::kernels::cpu::expert_layout().max_blob;
        if (!profile.empty()) slots = std::min<int64_t>(slots, (int64_t) profile.size());
        o.expert_cache = (int) std::max<int64_t>(slots, 0);
-        std::fprintf(stderr, "strata generate: expert cache auto: %.2f GiB free, %d MiB reserved -> %d slots\n",
-                     (double) free_b / 1073741824.0, o.vram_reserve_mib, o.expert_cache);
+        std::fprintf(stderr, "strata generate: expert cache auto: %.2f GiB free, %d MiB reserved (+%lld MiB for the "
+                             "draft head) -> %d slots\n",
+                     (double) free_b / 1073741824.0, o.vram_reserve_mib, (long long) (mtp_bind >> 20), o.expert_cache);
+        if (o.expert_cache == 0)   // the verify window cannot start without it (#174): say what makes room
+            std::fprintf(stderr, "strata generate: no VRAM is left for the expert cache: lower --max-context, use "
+                                 "--kv k8v4, run images on the CPU, or close other programs that use the GPU\n");
    } else if (multi_gpu && o.expert_cache > 0) {
        // a layer split's prompt path has its own buffers (it borrows no slots): an explicit cache size leaves room
        // for them and the reserve, or the first prompt fails with "device buffers ... do not fit"
@@ -2058,6 +2072,7 @@ int main(int argc, char** argv) {
        // STRATA_TEST_CACHE_FAIL=N: the first N opens fail as an out-of-commit cudaMalloc does (tests the retry)
        int fake_fails = std::getenv("STRATA_TEST_CACHE_FAIL") ? std::atoi(std::getenv("STRATA_TEST_CACHE_FAIL")) : 0;
        int failed = 0;
+        int zero_reads = 0;
        for (int attempt = 0;; ++attempt) {
            bool ok = false;
            if (fake_fails > 0) {
@@ -2101,9 +2116,11 @@ int main(int argc, char** argv) {
            cudaMemGetInfo(&free_b, &total_b);
            const int64_t want = (int64_t) o.vram_reserve_mib << 20;
            if ((int64_t) free_b >= want - (64ll << 20)) break;
-            // short by (want - free); a figure of 0 only says "at least", so then give back a quarter as well
+            // short by (want - free); a figure of 0 only says "at least": the first two such reads give back 1 GiB
+            // each (under WDDM the free figure read before the allocation runs ~0.7 GiB high), later ones a quarter
            int64_t give = want - (int64_t) free_b + (64ll << 20);
-            if (free_b < ((size_t) 16 << 20)) give = std::max<int64_t>(give, xcache.bytes() / 4);
+            if (free_b < ((size_t) 16 << 20))
+                give = std::max<int64_t>(give, ++zero_reads <= 2 ? 1ll << 30 : xcache.bytes() / 4);
            const int64_t keep_bytes = xcache.bytes() - give;
            std::fprintf(stderr, "strata generate: only %lld MiB free once the slots are written (reserve %d MiB); "
                                 "shrinking the expert cache\n", (long long) (free_b >> 20), o.vram_reserve_mib);
@@ -3266,7 +3283,9 @@ int main(int argc, char** argv) {
        sp.on_chunk = [&](const float* R_rows, int64_t T, int64_t p0, std::string& e) -> bool {
            std::vector<int32_t> nxt((size_t) T);
            for (int64_t t = 0; t < T; ++t) nxt[(size_t) t] = (int32_t) cur[(size_t) (p0 + t + 1)];
-            if (!mtp.prefill(R_rows, nxt.data(), T, p0, e)) return false;
+            // E-9: batched through the prompt path when it can (one GPU: a layer split's drafter is on the last stage)
+            const bool batched = !multi_gpu && sp.draft_kv(mtp, R_rows, nxt.data(), T, p0, e);
+            if (!e.empty() || (!batched && !mtp.prefill(R_rows, nxt.data(), T, p0, e))) return false;
            // progress for the server window: PP <position reached> <prompt tokens> <ms> <fresh tokens/s>
            const int64_t done = p0 + T;
            const double ms = std::chrono::duration<double, std::milli>(Clock::now() - pp_t0).count();
@@ -3551,6 +3570,7 @@ int main(int argc, char** argv) {
                ~BusyScope() { strata::core::progress().busy.store(false); strata::core::progress_at("idle"); }
            } busy_scope;
            stop_req.store(false);   // a STOP that arrived between requests is stale
+            err.clear();
            const bool geni = line.rfind("GENI ", 0) == 0;
            if (!geni && line.rfind("GEN ", 0) != 0) {
                std::printf("ERR expected: GEN <max_new> <id,id,...> or GENI <max_new> <file> <id,id,...>\n");
@@ -3836,6 +3856,7 @@ int main(int argc, char** argv) {
            };
            // tokens [a, b) through the windows: commit all of them, then give the draft layer their residuals
            auto read_windows = [&](int64_t a, int64_t b, std::string& e) -> bool {
+                strata::core::progress_at("reading the prompt (verify windows), from token", a);   // #217: not "batched"
                // every token is committed and the picks are discarded: no head sampling (see set_head_sampling)
                struct NoHeadSampling {
                    strata::core::Verifier& v;
@@ -3968,6 +3989,7 @@ int main(int argc, char** argv) {
            int64_t at = read_from;
            for (const int64_t to : {reread_to, root_at, turn_at, n - 1}) {
                if (to <= at) continue;
+                err.clear();
                const bool win = windows_ok(at, to);
                if (win && !refill(err)) {
                    std::printf("ERR refilling a lent slot failed: %s\n", err.c_str());
@@ -4371,7 +4393,8 @@ int main(int argc, char** argv) {
                // cell i pairs R_i with the token at i + 1 (every such token is in the prompt)
                std::vector<int32_t> nxt((size_t) T);
                for (int64_t t = 0; t < T; ++t) nxt[(size_t) t] = (int32_t) o.tokens[(size_t) (p0 + t + 1)];
-                return mtp.prefill(R_rows, nxt.data(), T, p0, e);
+                if (prefill.draft_kv(mtp, R_rows, nxt.data(), T, p0, e)) return true;   // E-9
+                return e.empty() && mtp.prefill(R_rows, nxt.data(), T, p0, e);
            };
        }
        const Clock::time_point tp0 = Clock::now();
@@ -0,0 +1,94 @@
+"""The draft head's token subset (rt/draft_vocab.bin): the vocabulary rows the MTP draft layer may propose.
+
+The engine reads rt/draft_vocab.bin (int32 token ids) when the draft layer binds and computes the draft logits over
+those rows only: a smaller head is faster and takes less VRAM, but a token outside the subset can never be drafted.
+The shipped subset was built from English and code and had 27 of the vocabulary's 55,328 Han tokens, so an answer in
+Chinese drafted almost nothing (#137).  This adds whole scripts to a subset:
+
+    python tools/draft_vocab.py --gguf <model>-00001-of-0000N.gguf --base data/draft_vocab.bin --add cjk \
+        --out data/draft_vocab.bin
+
+--add takes han, kana, hangul, cjk_punct, or cjk (all four).  The base ids keep their order; the added ones follow in
+id order.  --stats prints what a subset holds.
+"""
+from __future__ import annotations
+
+import argparse
+import sys
+from array import array
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parent))
+import strata_tokenizer as ST  # noqa: E402
+
+SCRIPTS = {
+    "han": [(0x3400, 0x4DBF), (0x4E00, 0x9FFF), (0xF900, 0xFAFF), (0x20000, 0x2FA1F)],
+    "kana": [(0x3040, 0x30FF), (0x31F0, 0x31FF), (0xFF66, 0xFF9F)],
+    "hangul": [(0x1100, 0x11FF), (0x3130, 0x318F), (0xAC00, 0xD7AF)],
+    "cjk_punct": [(0x3000, 0x303F), (0xFF00, 0xFF65), (0xFFA0, 0xFFEF)],
+}
+GROUPS = {"cjk": ["han", "kana", "hangul", "cjk_punct"]}
+
+
+def scripts_of(text: str) -> set[str]:
+    found = set()
+    for ch in text:
+        c = ord(ch)
+        for name, ranges in SCRIPTS.items():
+            if any(a <= c <= b for a, b in ranges):
+                found.add(name)
+    return found
+
+
+def token_text(tok, i: int) -> str | None:
+    """A token's text, or None when its bytes are not whole UTF-8 characters (a piece of one) or it is special."""
+    try:
+        return tok.token_bytes(i).decode("utf-8")
+    except (KeyError, UnicodeDecodeError, IndexError):
+        return None
+
+
+def main() -> int:
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--gguf", required=True, help="the model's (first) GGUF file: its vocabulary")
+    ap.add_argument("--base", help="the subset to extend (int32 ids)")
+    ap.add_argument("--add", default="", help="comma list: han, kana, hangul, cjk_punct, cjk")
+    ap.add_argument("--out", help="where to write the new subset")
+    ap.add_argument("--stats", action="store_true", help="print what the base (and the new) subset holds")
+    a = ap.parse_args()
+
+    tok = ST.Tokenizer.from_gguf(a.gguf)
+    n = len(tok.tokens)
+    kinds = [scripts_of(t) if (t := token_text(tok, i)) is not None else set() for i in range(n)]
+    base = list(array("i", Path(a.base).read_bytes())) if a.base else []
+    if any(i < 0 or i >= n for i in base):
+        sys.exit(f"{a.base} has ids outside the vocabulary ({n})")
+
+    want = set()
+    for x in filter(None, (s.strip() for s in a.add.split(","))):
+        names = GROUPS.get(x, [x])
+        if any(nm not in SCRIPTS for nm in names):
+            sys.exit(f"--add {x}: expected one of {', '.join([*SCRIPTS, *GROUPS])}")
+        want.update(names)
+    have = set(base)
+    added = [i for i in range(n) if i not in have and kinds[i] & want]
+    ids = base + added
+
+    def stats(label, sel):
+        counts = {nm: sum(1 for i in sel if nm in kinds[i]) for nm in SCRIPTS}
+        full = {nm: sum(1 for k in kinds if nm in k) for nm in SCRIPTS}
+        print(f"{label}: {len(sel)} ids; " + ", ".join(f"{nm} {counts[nm]}/{full[nm]}" for nm in SCRIPTS))
+
+    if a.stats or not a.out:
+        if base:
+            stats("base", base)
+        if added:
+            stats("new", ids)
+    if a.out:
+        Path(a.out).write_bytes(array("i", ids).tobytes())
+        print(f"wrote {a.out}: {len(ids)} ids ({len(added)} added)")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
Author	SHA1	Message	Date
Niko1221	10409135dd	Engine 0.1.28: the expert cache reserves the draft head's VRAM (0.1.27 left 16 GB cards below the stall line), a cancelled request no longer fails the next one, the API key protects /status; setup requires it	2026-09-30 11:18:54 +02:00
Niko1221	03939f1c06	docs: the draft layer's token subset and --draft-vocab	2026-09-30 11:18:54 +02:00
Niko1221	75e6efeb77	setup: --host/--api-key/--draft-vocab on a start are kept for the model (#179 ); --draft-vocab en keeps the English/code subset (~110 MiB less VRAM); CUDA 13.0 for RTX 50 builds (#220 ); llama.cpp's web UI is not unpacked (Windows path limit, #206 ); Swift Q2_0 hidden (#171 ); the earlier data folder and nested Strata-data move into --data-dir (#198 ); whole shards without a finish mark count (#173 )	2026-09-30 10:54:45 +02:00
Niko1221	fa59b3f3a9	server: /status needs the API key and forgets the answer's end (#212 ); constant-time key check, an explicitly empty key refuses to start (#213 ); an engine exit quotes its last log line (#215 ); tool calls end at their own </parameter> and </tool_call>, so values may contain them (#210 )	2026-09-30 10:54:45 +02:00
Niko1221	eef241178c	engine: the expert cache reserves the draft head's VRAM (0.1.27's CJK subset came out of the 700 MiB reserve, #199 ); a clear message when no VRAM is left for the cache (#174 ); the prompt's verify-window stage has its own label (#217 )	2026-09-30 10:54:45 +02:00
Oluwabori Olaleye	7980d3719f	Expert cache on WDDM: 1 GiB steps before quarter steps (~4 GiB more cache on a 32 GB card); Ctrl+C stops the server on Windows; no stack trace when a client hangs up (#153 ) Squashed from PR #153.	2026-09-30 10:46:55 +02:00
GioStrives	0425829226	verify: cudaFuncGetName only with CUDA 12.3+ (an sm_75 build with CUDA 12.0 failed) It names kernels only for the STRATA_VERIFY_NODES debug dump.	2026-09-30 10:46:55 +02:00
Pravesh Khatana	481b69246f	fix(engine): clear stale error/cancellation state between requests and before prompt prefill When a request is cancelled during prompt reading (e.g. client disconnect or timeout), err is set to 'cancelled'. Because err is declared in main() and was never cleared between requests, any subsequent request in v0.1.26+ evaluated 'if (!e.empty())' in sp.on_chunk on its first chunk, immediately failing with 'strata serve: cancelled' and terminating the engine with code 1. Clearing err at the start of each request, before prompt segment runs, and on entry to Prefill::run prevents state pollution across requests.	2026-09-30 10:46:55 +02:00
Niko1221	99b40973c6	Engine 0.1.27: CJK answers faster (the draft subset with Chinese, Japanese and Korean), RTX 20 cards, the HIP build without CUDA headers; setup requires it	2026-09-29 22:58:59 +02:00
Niko1221	e227683eb9	data/draft_vocab.bin: the draft subset with all Chinese, Japanese and Korean tokens (40,525 -> 106,299 ids; #137 ). Q2_0 on CJK answers +15-38% decode, English -1-2%, the same expert slots on 12 GB	2026-09-29 22:58:59 +02:00
Niko1221	d62579d482	RTX 20 in the docs and setup's texts; the ready-made image encoder is compiled here when it has no code for the card (BUILD.json vision_archs)	2026-09-29 22:58:59 +02:00
Niko1221	717971d33b	tools/draft_vocab.py: the draft head's token subset, with whole scripts added (#137 ); setup replaces an older shipped subset at setup and at start, a hand-made one is kept	2026-09-29 22:26:45 +02:00
Niko1221	d92c2a3034	setup: AVX2 from the CPU itself when Windows says no (#159 ); the engine is compiled again for a card it has no code for, at setup and at start (#128 )	2026-09-29 22:23:55 +02:00
Niko1221	ed01b4de03	server: only the image markers the template writes are images; the same text in a message stays text (#150 )	2026-09-29 22:23:55 +02:00
Hozák Martin	ddfb98a6ea	Turing port: run on RTX 20 (sm_75) and newer Rebased onto main (0.1.22): the portable pre-sm_80 QSA scorer and the sm_80 guard on the tensor-core prompt attention are already in main, so this is the rest of the port. - the runtime floor: device_info refuses anything older than compute capability 7 (main's runtime check still requires sm_120, so a binary built for sm_75 would never start), - fused_gr: the down kernel carries the tokens in slices that fit the card's shared-memory opt-in (Turing, 64 KB: 6-token slices; 4 tokens when the card reports no opt-in, under the 48 KB default), merged on top of main's per-device opt-in - full 8-token windows no longer fail on Turing with "invalid argument", - the setup/CMake floor is 7.5 (RTX 20 / 30 / 40 / 50) without an experimental flag: CMake's arch refusals, setup.py's gpu_problem and the GPU hints and error texts, - docs/MULTI_GPU.md says what a Turing card actually runs. Verified end-to-end on RTX 2070 (sm_75, 0.1.20 base): model loads, OpenAI/ Anthropic API serving works, greedy output identical to the sm_120 build. Merged (PR #87) with two changes: the runtime floor is 7.5 exactly (major*10+minor < 75; the PR let Volta 7.0/7.2 through), kept beside main's HIP check; and a card reporting no shared-memory opt-in slices the tokens to what fits the 48 KB default for its own tile (4 tokens on CUDA, all 8 with HIP's smaller tile, so gfx1100 keeps one launch).	2026-09-29 22:23:41 +02:00
samuelishida	c81819a87c	build: make qsa_prompt_attn.cu compile on HIP-only toolchains The tensor-core prompt-attention kernel is already refused at runtime on AMD (line 695), but the file unconditionally included <math_constants.h> for CUDART_INF_F, which is not present on a ROCm-only build. Replace it with the portable INFINITY macro from <cmath> so the HIP build succeeds without a CUDA toolkit installed.	2026-09-29 22:19:51 +02:00
Niko1221	4832dccd09	docs: speed tables for engine 0.1.26 (prompts 8-28% faster at 32K-128K than 0.1.22)	2026-09-29 22:09:44 +02:00
Niko1221	78db15b630	docs: the low-RAM mode	2026-09-29 21:23:21 +02:00
Niko1221	cfc73be8b9	Low-RAM mode: --mmap-experts works with native (IQ) packs built with experts.bin; setup uses it when the model's experts do not fit the RAM beside the OS (a 32 GB PC with a big GPU: the experts come from the file cache, the GPU holds the most-used)	2026-09-29 20:35:16 +02:00
Niko1221	2f8d9be954	Engine 0.1.26: prompts up to 32K 4-10% faster (the draft layer's prompt pass in batches); setup requires it	2026-09-29 20:01:05 +02:00
Niko1221	c9077a7669	setup: a PC with NVIDIA and AMD cards asks which to use (a first run took NVIDIA without mentioning the Radeon)	2026-09-29 19:59:11 +02:00
Niko1221	a402eefdd1	E-9: the draft layer's prompt pass in batches through the prompt path's kernels (its Q8_0 matrices as Q8_1 x Q8_0 MMQ, the drafter's own integer dot products): 0.9 s -> 0.07 s of a 32K prompt, the same draft acceptance; STRATA_MTP_BATCH=0 switches back	2026-09-29 19:59:11 +02:00