Prompt lookup that pays, sampling fixes, shallow-path setup fix (engine 0.1.7)

- Prompt lookup on by default: the suffix drafter's windows are taken only where a learned policy (measured
  acceptance per match length, measured window cost) expects more tokens per ms than the MTP window; lookup windows
  up to --spec + 2. Code edits +6-11%, other text unchanged, verified output identical. --suffix-draft 0 = off.
- Per-request sampling (#19) keeps greedy requests on the in-graph sampler (no extra launch per window), and draws
  row t of a window with Philox(seed, pos0 + t), so a seed replays the same text whatever the drafts were.
- tools: a shallow install (G:\Strata) no longer crashes the MTP packer and the pack tools (IndexError, #20).
- setup: says at step 1 when the PC has less RAM than the smallest model needs; requires engine 0.1.7.

CMakeLists.txt

@@ -177,6 +177,7 @@ if(STRATA_ENABLE_CUDA)
                           src/kernels/cuda/native_qsa_score.cu
                           src/kernels/cuda/ple.cu
                           src/kernels/cuda/kv_q8.cu
+                          src/kernels/cuda/kv_stream.cu
                           src/kernels/cuda/fused_gr.cu
                           src/kernels/cuda/qsa_decode_attn.cu
                           src/kernels/cuda/fused_gdn.cu
@@ -250,7 +251,7 @@ if(STRATA_ENABLE_CUDA)
 
   add_executable(strata src/program/generate.cpp)
   target_include_directories(strata PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
-  target_link_libraries(strata PRIVATE strata_engine strata_prefill CUDA::cudart)
+  target_link_libraries(strata PRIVATE strata_engine strata_prefill strata_spec CUDA::cudart)
 
   add_executable(dequant_s2_parity src/kernels/dequant_s2_parity.cpp)
   # The tests are HOST C++ files that call cudaMalloc/cudaMemcpy themselves, so they need the toolkit headers
@@ -356,6 +357,12 @@ if(STRATA_ENABLE_CUDA)
   target_link_libraries(kv_q8_parity PRIVATE strata_kernels CUDA::cudart)
   add_test(NAME kv_q8_parity COMMAND kv_q8_parity)
 
+  # KV streaming (docs/kv-streaming-design.md): streamed vs resident attention, bitwise (GPU, synthetic, no model).
+  add_executable(kv_stream_parity src/kernels/kv_stream_parity.cpp)
+  target_include_directories(kv_stream_parity PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+  target_link_libraries(kv_stream_parity PRIVATE strata_kernels CUDA::cudart)
+  add_test(NAME kv_stream_parity COMMAND kv_stream_parity)
+
   # Plan v0.3 P3: native MMVQ for 1..8 tokens (bitwise selftest + --bench timing); synthetic data, no model.
   if(STRATA_BUILD_TESTS)
   add_executable(native_mmvq_multi bench/micro/native_mmvq_multi.cpp)
@@ -417,7 +424,7 @@ endif()
 #     worse opt-out.  `cpu_require_expert_support()` therefore refuses at startup on a CPU without the
 #     features, and the fallback exists for tests, exactly as P2.S3 specifies.
 # Plan v0.3 P6: speculation components that need no GPU.
-add_library(strata_spec STATIC src/spec/suffix_drafter.cpp src/spec/controller.cpp)
+add_library(strata_spec STATIC src/spec/suffix_drafter.cpp src/spec/controller.cpp src/spec/draft_policy.cpp)
 target_include_directories(strata_spec PUBLIC include)
 if(STRATA_BUILD_TESTS)
   add_executable(suffix_drafter_test src/spec/suffix_drafter_test.cpp)
@@ -426,6 +433,9 @@ if(STRATA_BUILD_TESTS)
   add_executable(controller_test src/spec/controller_test.cpp)
   target_link_libraries(controller_test PRIVATE strata_spec)
   add_test(NAME controller_test COMMAND controller_test)
+  add_executable(draft_policy_test src/spec/draft_policy_test.cpp)
+  target_link_libraries(draft_policy_test PRIVATE strata_spec)
+  add_test(NAME draft_policy_test COMMAND draft_policy_test)
 endif()
 
 add_library(strata_kernels_cpu STATIC src/kernels/cpu/expert.cpp src/kernels/cpu/pool.cpp

README.md

@@ -39,7 +39,7 @@ model - is set up for you.
 2. Double-click **`START-HERE.bat`**.
 3. Answer 4 questions - or just press Enter each time for the recommended choice:
    - **Which model?** The original, or Swift 1.5 (a version that thinks shorter and answers sooner)
-   - **Which size?** Q2_0, IQ2_XS or IQ3_XXS - see [which model](#which-model-should-i-pick)
+   - **Which size?** Q2_0, IQ2_XS, IQ3_XXS or IQ3_S - see [which model](#which-model-should-i-pick)
    - **How much context?** How much text it can keep in mind at once (it suggests one for your card)
    - **Images?** Whether it should also read pictures
 
@@ -60,6 +60,7 @@ Measured on an RTX 5070 (12 GB), a Ryzen 5 7600 and 64 GB of RAM:
 | **Q2_0** | 95 tokens/s | 65 tokens/s | 539 tokens/s |
 | **IQ2_XS** | 78 tokens/s | 52 tokens/s | 463 tokens/s |
 | **IQ3_XXS** | 66 tokens/s | 45 tokens/s | 410 tokens/s |
+| **IQ3_S** | 54 tokens/s | 42 tokens/s | 374 tokens/s |
 
 - **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).
@@ -76,7 +77,8 @@ A card with more VRAM is faster, because more of the model fits on the GPU: an R
 | --- | ---: | --- | --- | --- |
 | **Q2_0** | 66 GB | fastest | good | you want speed |
 | **IQ2_XS** | 68 GB | fast | better | you want a good all-rounder (**recommended**) |
-| **IQ3_XXS** | 76 GB | slower | best | you want the best answers (uses 43 GB of your 64 GB RAM) |
+| **IQ3_XXS** | 76 GB | slower | great | you want better answers (uses 43 GB of your 64 GB RAM) |
+| **IQ3_S** | 84 GB | slowest | best: matches the full model on the published tests | you want the very best answers (original model only; uses 50 GB of your 64 GB RAM, so close other big programs) |
 
 **The version:**
 
@@ -98,8 +100,8 @@ Not sure? Take **IQ2_XS**. You can add another one later with `START-HERE.bat --
 - **Pictures:** in the chat page click **Picture**; in `chat.py` type `/image <path>`; in apps just attach them.
 - **From your phone or another PC:** see the [details](docs/DETAILS.md#using-it) (set an API key first).
 
-**Good to know:** it answers one request at a time, and it re-reads the whole conversation for every answer. So in
-very long chats you wait longer before it starts writing: about 1 minute per 30,000 tokens of conversation.
+**Good to know:** it answers one request at a time. The first message of a chat is read in full (about 1 minute per
+30,000 tokens); after that it keeps the conversation and reads only what is new, so follow-ups start in seconds.
 
 ## Something went wrong?
 

docs/DETAILS.md

@@ -21,6 +21,7 @@ tokens, MTP speculative decoding on. "262K" is the model's full context window (
 | **Q2_0** | 389 | 539 | 571 | 561 | 543 | 496 |
 | **IQ2_XS** | 332 | 463 | 495 | 486 | 472 | 437 |
 | **IQ3_XXS** | 285 | 410 | 435 | 427 | 414 | - |
+| **IQ3_S** | 260 | 374 | 397 | - | 378 | - |
 
 ### Output (tokens/s)
 
@@ -29,9 +30,17 @@ tokens, MTP speculative decoding on. "262K" is the model's full context window (
 | **Q2_0** | 88.7 | 94.6 | 87.5 | 76.4 | 65.1 | 56.3 |
 | **IQ2_XS** | 82.0 | 78.0 | 65.3 | 63.7 | 52.0 | 48.0 |
 | **IQ3_XXS** | 64.6 | 65.6 | 57.3 | 54.4 | 44.8 | - |
+| **IQ3_S** | 51.2 | 54.4 | 51.1 | - | 42.2 | - |
 
-IQ3_XXS at 262K is not measured: with its 43 GB of experts, a 260K-token context brings a 64 GB PC to its memory
-limit. Use up to 128K with IQ3_XXS on 64 GB.
+IQ3_XXS and IQ3_S at 262K are not measured: with their 43 / 50 GB of experts, a 260K-token context brings a 64 GB PC
+to its memory limit. Use up to 128K with them on 64 GB (setup caps it). IQ3_S (engine 0.1.4 or newer) is only published
+for the original model, not for Swift 1.5.
+
+**KV streaming (engine 0.1.5):** at 64K and more, setup keeps the context's KV cache in RAM and only the part the
+attention reads in VRAM (`--kv-resident 32768`), so more experts fit on the GPU. Q2_0 at 262K: 50.9 -> 62.6 tokens/s
+(1,589 -> 3,872 experts in VRAM); at 128K about +6%. The attention reads exactly the same values (only where the KV lives changes); it
+costs ~13.7 KB of RAM per context token (1.7 GB at 128K). Existing installs: run `START-HERE.bat --setup` once to turn
+it on.
 
 Time to first token is prompt length / prompt speed: about 7 s at 4K, 55 s at 32K, 4 minutes at 128K and 9 minutes at
 262K. The raw numbers: [`bench/results/`](../bench/results/). The [paper](paper/Strata-Paper.pdf) explains every number.
@@ -200,9 +209,21 @@ print(r.choices[0].message.content)
   add `"api_key": "some-long-secret"` to `strata-<model>.json` (or set the `STRATA_API_KEY` environment variable);
   clients then send it as their API key.
 
-**Current limits (v1):** one request at a time; greedy decoding (temperature is ignored); every request processes its
-whole prompt again (no conversation cache yet, so long chats have a long time-to-first-token); images only when set up
-with them (below); no video.
+**Conversation cache.** A request that continues a chat reads only the part after what the engine already holds: the
+live session, or one of the checkpoints it keeps in RAM (up to 6, ~118 MB each, taken at the start of each new
+assistant turn and every 16K prompt tokens). A checkpoint is used only when the prompt starts with exactly its tokens
+and pictures. Engine options: `--prompt-cache N` (0 = off), `--prompt-cache-every N`, `--turn-token ID`.
+
+**Current limits (v1):** one request at a time, and one conversation cached at a time (switching between two chats
+re-reads the other one); images only when set up with them (below); no video. **Temperature / top_p / top_k / min_p /
+seed** are honored per request (OpenAI and Anthropic fields); with the default adaptive expert tier a sampled result
+is not reproducible run to run - for seed-reproducible output add `--adapt-every 100000` (static residency) to the
+engine arguments. The run config's optional `sampling` block sets the defaults for requests that leave the fields out
+(`"sampling": {"temperature": 1.0, "top_p": 0.95, "top_k": 20}`); a request's own fields always win, and with no
+block at all a request without sampling keys decodes greedy. The penalties (`presence_penalty`, `frequency_penalty`,
+`repetition_penalty`, with `penalty_last_n` capping how many recent tokens they count over, default 64 when any
+penalty is set) ride the same path; they count the tokens the request has consumed, so a repetition penalty
+suppresses what the model itself just said, not the prompt alone.
 
 ---
 
@@ -305,13 +326,15 @@ test images.
 <p align="center"><img src="paper/tiers.svg" width="760" alt="memory tiers"></p>
 
 - **GPU (VRAM):** attention and DeltaNet mixers, the gated-residual weights, routers, shared experts, output head, the MTP
-  draft layer, the KV cache, and an **expert cache** that fills the rest of VRAM with the most-used experts (it adapts to
+  draft layer, the KV cache (from 64K: only its most-read part, the rest streams from RAM), and an **expert cache** that fills the rest of VRAM with the most-used experts (it adapts to
   the conversation while you chat).
 - **RAM:** all 24,576 experts, pinned. The CPU computes the experts that are not on the GPU **in place**, at the same time
   as the GPU works on the cached ones (AVX-512 / AVX2 kernels, ggml's for the i-quants).
 - **SSD:** the 28.8 GB n-gram table, read a few rows per token through the OS cache.
 - **Speculation:** the model's own MTP layer drafts up to 3 tokens; one pass over all 48 layers checks them. 2.4-3.2
-  tokens per pass on average.
+  tokens per pass on average. When the reply repeats the context (code edits, quoted text), **prompt lookup** (engine
+  0.1.7) drafts up to 5 tokens from the earlier copy, but only where its measured acceptance and cost say it pays:
+  code edits 6-11% faster, other text unchanged. The drafts are checked like the MTP's, so the output is the same.
 - **Prompts** are processed in 2,048-token chunks with the experts streamed to the GPU over PCIe.
 
 The full story, with measurements, bottlenecks and what comes next: **[docs/paper/Strata-Paper.pdf](paper/Strata-Paper.pdf)**.

include/strata/core/layer.hpp

@@ -36,6 +36,7 @@
 #include "strata/core/weights.hpp"
 
 #include "strata/kernels/gr.hpp"
+#include "strata/kernels/kv_stream.hpp"
 #include "strata/kernels/ngram.hpp"
 #include "strata/kernels/ple.hpp"
 
@@ -206,10 +207,20 @@ struct QsaState {
     int8_t* v_q = nullptr;
     uint16_t* k_scale = nullptr;
     uint16_t* v_scale = nullptr;
-    int32_t* page_table = nullptr;   ///< (n_pages,) logical page -> physical page
+    int32_t* page_table = nullptr;   ///< (n_pages,) logical page -> physical page (-1: not resident, streamed)
     int64_t n_pages = 0;
     int64_t max_cells = 0;
 
+    /// KV STREAMING (docs/kv-streaming-design.md, `kv_stream.hpp`). `kv_mode` 0: every page in VRAM, identity
+    /// table (n_slots == n_pages, no host copy). 1: streamed - the authoritative K/V in `host`, `n_slots` pages
+    /// in VRAM, `map` resolves the selection's blocks on device. 2: a ring (the MTP drafter's window) - `n_slots`
+    /// pages at `block % n_slots`, with the host copy for a resume. The pool pointers above point at the slots.
+    int kv_mode = 0;
+    int64_t n_slots = 0;
+    strata::kernels::KvHostPools host;
+    strata::kernels::KvStreamMap map;
+    int64_t idx_pooled_rows = 0;     ///< rows of `idx_pooled` (a ring, which has no indexer, keeps 2)
+
     float* idx_tail = nullptr;       ///< (idx_block - 1, idx_dim): the raw tail of the block being filled
     float* idx_dead = nullptr;       ///< (idx_dim,): the spare slot's key, CONSTANT for the sequence
     float* idx_pooled = nullptr;     ///< (max_cells/idx_block + 2, idx_dim)
@@ -242,13 +253,28 @@ struct QsaState {
 /// Plan v0.3 P7: the RoPE cos/sin table (max_cells x n_rot/2 x 2 floats, 64 MiB at 262K) is identical in every
 /// QSA layer. `with_rope = false` sizes a state that borrows it; `share_rope` points `st` at another state's table
 /// instead of building a copy (the session builds it once, in the first QSA layer).
-uint64_t qsa_state_bytes(const ModelGeometry& g, int64_t max_cells, bool with_rope = true);
+uint64_t qsa_state_bytes(const ModelGeometry& g, int64_t max_cells, bool with_rope = true, int64_t ring_cells = 0);
+/// KV streaming: keep `cells` cells of each QSA layer in VRAM and the rest in pinned host memory (0: all in VRAM,
+/// the default). Set before sizing and initializing the session; a context that fits in `cells` is not streamed.
+/// Also puts the MTP drafter's K/V in a ring of its window (`ring_cells` of qsa_state_bytes/init; -1 forces a fully
+/// resident state).
+void qsa_set_kv_resident(int64_t cells);
+int64_t qsa_kv_resident();
+/// The fewest resident cells a streamed layer may have: one verify window's selections (8 queries x 2,051 cells
+/// in whole blocks) must fit at once, with room to spare.
+int64_t qsa_kv_resident_min();
+/// Pinned host bytes the streamed states hold (their host copies).
+uint64_t qsa_kv_host_bytes();
 /// Plan v0.3 P7: store K/V as INT8 with FP16 scales per 64 values (half the VRAM of FP16). Set before sizing and
 /// initializing the session; default off until gate G-C accepts it.
 void qsa_set_kv_int8(bool enabled);
 bool qsa_kv_int8();
 uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, QsaState& st,
-                        const QsaState* share_rope = nullptr);
+                        const QsaState* share_rope = nullptr, int64_t ring_cells = 0);
+/// KV streaming: the pools a reader sees (the VRAM slots) and, when streamed, make the selection's blocks resident.
+strata::kernels::QsaAttnPools qsa_attn_pools(const QsaState& st);
+void qsa_kv_resolve(const QsaState& st, const ModelGeometry& g, const int32_t* ids, const int32_t* steps, int64_t n_q,
+                    int64_t cap, void* stream);
 /// Zeroes the pools AND the indexer, so a fresh sequence matches the reference's own `zeros()`.  The KV pool
 /// matters even for cells that are never attended, because `kv_gather` reads whatever the selection names.
 void qsa_state_zero(const QsaState& st, const ModelGeometry& g, void* stream);

include/strata/core/mtp.hpp

@@ -45,7 +45,14 @@ public:
               int64_t window = 32768);
     /// The prompt's length: prefill() skips the cells the attention window can never reach again.
     void set_prompt_len(int64_t n) { prompt_len_ = n; }
+    /// At most this many drafts per round (below max_t - 1): a window longer than the MTP's comes from elsewhere.
+    void set_max_drafts(int k) { max_drafts_ = k; }
     uint64_t vram_bytes() const { return vram_; }
+    /// The draft layer's K/V state (read-only: --serve's STRATA_STATE_HASH check hashes it)
+    const QsaState& kv_state() const { return st_; }
+    /// 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 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);
 
@@ -82,6 +89,7 @@ private:
     const NativeHead* head_ = nullptr;
     const float* window_R_ = nullptr;
     int max_t_ = 0;
+    int max_drafts_ = 1 << 30;
     int64_t n_vocab_ = 0;
     uint64_t vram_ = 0;
     cudaStream_t cs_ = nullptr;

include/strata/core/verify.hpp

@@ -24,6 +24,7 @@
 #include "strata/core/expert_source.hpp"
 #include "strata/core/layer.hpp"
 #include "strata/core/session.hpp"
+#include "strata/kernels/sampler.hpp"
 
 #include <cuda_runtime.h>
 
@@ -58,6 +59,21 @@ public:
     /// One window: `tokens[0..T)` at positions pos0.., the pool served per layer; `out[t]` = argmax after token t.
     /// The PLE rows are gathered here from `ss.ple_prev` and the tokens.  Captures the T-token graph on first use.
     bool run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool, void* user, int32_t* out, std::string& err);
+    /// The sampling the verify window's head applies (temperature / top_p / top_k / seed).  Set per
+    /// request; greedy by default.  The sampling itself runs OUTSIDE the captured graph - its
+    /// parameters would otherwise be baked forever - so this can change between requests freely.
+    void set_sampling(const strata::kernels::SamplerParams& sp) {
+        sampling_ = sp;   // row t of a window at pos0 draws Philox(seed, pos0 + t): see run()
+    }
+
+    /// The penalty-history row for `sampling_.penalty_last_n`: ONE row of `history_len` int32 slots, most
+    /// recent token LAST, unused front slots -1 (the kernel reads only the tail window).  Null disables the
+    /// penalties entirely - the neutral run's sampling call is byte-for-byte what it was.  The engine
+    /// re-uploads the request's tail before every window; the pointer must stay alive across the request.
+    void set_history(const int32_t* history, int history_len) {
+        hist_d_ = history;
+        hist_len_ = history_len;
+    }
 
     /// Keep the first `n_keep` (1..T) tokens of the last window; advances `ss.ple_prev` by them.
     bool commit(int n_keep, std::string& err);
@@ -83,6 +99,14 @@ public:
 
 private:
     bool capture(int T, std::string& err);
+    strata::kernels::SamplerParams sampling_ = [] {
+        strata::kernels::SamplerParams s;
+        s.greedy = true;
+        s.temperature = 0.0f;
+        return s;
+    }();   ///< greedy by default; per-request via set_sampling
+    const int32_t* hist_d_ = nullptr;   ///< penalty-history row (set_history); null = no penalties apply
+    int hist_len_ = 0;
     bool capture_commit(std::string& err);
     bool record_window(int T, cudaStream_t cs, std::string& err);
 

include/strata/kernels/cpu/pool.hpp

@@ -23,7 +23,10 @@
 #include "strata/kernels/cpu/native_expert.hpp"
 
 #include <atomic>
+#include <chrono>
+#include <condition_variable>
 #include <cstdint>
+#include <mutex>
 #include <thread>
 #include <vector>
 
@@ -136,10 +139,19 @@ public:
         repark = ms_repark_;
     }
 
+    /// **A PARKED WORKER SPINS FOR THIS LONG, THEN SLEEPS.**  The park is a `_mm_pause` spin because a layer's
+    /// batches are microseconds apart and a wake-up from the OS costs more than that.  But a spin that never ends
+    /// keeps every worker's core at 100% while the engine waits for a request - issue #4, "CPU 50% even when
+    /// doing nothing" (7 of the 5700X's 16 threads).  Inside a request the gaps between batches are far below
+    /// this, so the token path still never sleeps; between requests the workers block on `sleep_cv_`.
+    static constexpr std::chrono::milliseconds kSpinBeforeSleep{20};
+
 private:
     void worker(int id);
     void drain(int id, ExpertScratch& scratch);
     void run_phase(int mode, int n_tasks);
+    /// Bump `epoch_`, and wake the workers that went to sleep.  Every publish goes through here.
+    void publish();
 
     int n_ = 0;
     bool host_works_ = true;
@@ -169,6 +181,11 @@ private:
     alignas(64) std::atomic<uint32_t> parked_{0};
     alignas(64) std::atomic<uint32_t> epoch_{0};
     alignas(64) std::atomic<bool> stop_{false};
+    // The sleep after `kSpinBeforeSleep`.  `sleepers_` is how `publish` knows whether anyone needs waking, so the
+    // token path pays one uncontended load per publish and never takes the mutex while the workers spin.
+    alignas(64) std::atomic<uint32_t> sleepers_{0};
+    std::mutex sleep_mu_;
+    std::condition_variable sleep_cv_;
     std::vector<std::thread> threads_;
     std::vector<ExpertScratch> scratch_;   // one per worker: no allocation, no false sharing of the hot data
     // run_split state: mode 0 = whole experts, 1 = gate/up row parts, 2 = down row parts

include/strata/kernels/kv_q8.hpp

@@ -12,6 +12,7 @@
 // the FP16 pools: codes `[page][kv_head][page_size][head_dim]`, scales `[page][kv_head][page_size][head_dim/64]`.
 #pragma once
 
+#include "strata/kernels/kv_stream.hpp"
 #include "strata/kernels/qsa.hpp"
 
 #include <cstdint>
@@ -25,9 +26,11 @@ inline uint64_t kv_q8_bytes_per_cell(const QsaShapes& s) {
     return (uint64_t) s.n_head_kv * s.head_dim * 2 + (uint64_t) s.n_head_kv * (s.head_dim / KV_Q8_GROUP) * 2 * 2;
 }
 
-/// Append the cell at step[kStepPos] (graph-capturable: position and page come from device memory).
+/// Append the cell at step[kStepPos] (graph-capturable: position and page come from device memory). With a host
+/// copy (KV streaming, `kv_stream.hpp`) the cell is written there too, and to VRAM only if its block is resident.
 void kv_append_q8_step(int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale, const int32_t* page_table,
-                       const int32_t* step, const float* kcur, const float* vcur, const QsaShapes& s, void* stream);
+                       const int32_t* step, const float* kcur, const float* vcur, const QsaShapes& s, void* stream,
+                       const KvHostPools* host = nullptr);
 
 /// Gather step[kStepWidth] cells named by `ids` into FP16 scratch `[id][kv_head][head_dim]`; the grid is sized by
 /// `max_ids` (capacity), the kernel reads the real count from `step`.

include/strata/kernels/kv_stream.hpp

@@ -0,0 +1,92 @@
+// include/strata/kernels/kv_stream.hpp - KV streaming for the QSA layers (docs/kv-streaming-design.md).
+//
+// A streamed QSA layer keeps its AUTHORITATIVE K/V in pinned, device-mapped host memory, laid out exactly as a
+// fully resident pool with the identity page table would be. VRAM holds a pool of `n_slots` pages (one page =
+// `page_size` = 4 cells = one indexer block, all KV heads, K and V: 4,224 B in int8), and the page table becomes a
+// RESIDENCY MAP: `page_table[block]` is the slot holding that block, or -1.
+//
+// Every reader (`qsa_decode_attn_*`, `kv_gather_*`) is unchanged: it goes through the page table as before. What
+// changes is that after the selection and before the attention, `kv_stream_resolve` makes every block the
+// selection names resident - on device, inside the captured graphs:
+//   1. one block of 1,024 threads walks the selected ids: a resident block gets this call's epoch and its clock
+//      reference bit, a missing block is claimed once (page_table -1 -> -2) into a miss list;
+//   2. the same block picks one victim slot per miss with a CLOCK sweep (second chance) that never takes a slot
+//      this call uses, re-points the page table (old block -> -1, new block -> slot);
+//   3. a copy kernel reads the missed blocks from the host copy (zero-copy, over PCIe) into their slots.
+// Writers (`kv_append_*`) write the host copy always, and the slot only if the block is resident, so the slots
+// never go stale and nothing needs invalidating - not after a rejected draft, not on a conversation-cache resume.
+//
+// A RING (the MTP drafter, which only ever reads its last `window` cells) uses the same host copy with a static
+// table `block -> block % n_slots` and no resolve: the ring is refilled from the host copy on a resume.
+#pragma once
+
+#include "strata/kernels/qsa.hpp"
+#include "strata/kernels/qsa_decode_attn.hpp"
+
+#include <cstdint>
+
+namespace strata::kernels {
+
+/// The host copy of one layer's K/V: device-mapped pointers (UVA) into pinned memory, identity layout
+/// `[block][kv_head][page_size][head_dim]`. All null when the layer is fully resident.
+struct KvHostPools {
+    uint16_t* k_pool = nullptr;   ///< fp16 mode
+    uint16_t* v_pool = nullptr;
+    int8_t* k_q = nullptr;        ///< int8 mode: codes
+    int8_t* v_q = nullptr;
+    uint16_t* k_scale = nullptr;  ///< int8 mode: fp16 scale per 64 values
+    uint16_t* v_scale = nullptr;
+    bool present() const { return k_pool != nullptr || k_q != nullptr; }
+};
+
+/// The residency map of a streamed layer, all device memory at fixed addresses (the graphs bake them in).
+struct KvStreamMap {
+    int32_t* page_table = nullptr;  ///< (n_blocks,) block -> slot, or -1
+    int32_t* slot_block = nullptr;  ///< (n_slots,) slot -> block, or -1
+    int32_t* slot_stamp = nullptr;  ///< (n_slots,) the resolve epoch that last used the slot
+    int32_t* slot_ref = nullptr;    ///< (n_slots,) clock reference bit
+    int32_t* ctl = nullptr;         ///< kKvCtlInts: epoch, hand, misses of the last call, overflow, u64 counters
+    int32_t* miss_block = nullptr;  ///< (n_slots,)
+    int32_t* miss_slot = nullptr;   ///< (n_slots,)
+    int64_t n_blocks = 0;
+    int64_t n_slots = 0;
+};
+
+inline constexpr int kKvCtlInts = 16;   ///< [0] epoch [1] hand [2] misses [3] overflow; u64 at [4] misses, [6] lookups, [8] calls
+/// Bytes of the map's arrays besides the page table (which the state already has): 5 per-slot ints + ctl.
+inline uint64_t kv_stream_map_bytes(int64_t n_slots) { return (uint64_t) n_slots * 4 * 5 + kKvCtlInts * 4; }
+
+/// Every block evicted: page table all -1, slots free, counters zero.
+void kv_stream_reset(const KvStreamMap& m, void* stream);
+
+/// Make every block named by the selections of `n_q` queries resident (ids [n_q][cap], width from
+/// steps[q * kStepCount + kStepWidth]). Capturable. `n_slots` must hold the distinct blocks of one call
+/// (n_q x (cap / page_size + 2)); a call that cannot sets ctl[3] (see `kv_stream_counters`).
+void kv_stream_resolve(const KvStreamMap& m, const QsaAttnPools& slots, const KvHostPools& host, bool int8,
+                       const int32_t* ids, const int32_t* steps, int64_t n_q, int64_t cap, const QsaShapes& s,
+                       void* stream);
+
+/// The static ring table `block -> block % n_slots`.
+void kv_ring_table(int32_t* page_table, int64_t n_blocks, int64_t n_slots, void* stream);
+
+/// Copy blocks [b0, b1) of the host copy into the pool pages `page_table` names (host-side table `phys(b)`
+/// computed as `b % n_slots`: the ring restore). Not capturable.
+void kv_ring_restore(const QsaAttnPools& slots, const KvHostPools& host, bool int8, int64_t b0, int64_t b1,
+                     int64_t n_slots, const QsaShapes& s, void* stream);
+
+/// Copy the first `n_blocks` blocks of the host copy into a fully resident (identity-layout) pool: the prompt
+/// path's staging of one layer. Not capturable (DMA).
+void kv_stage_from_host(const QsaAttnPools& stage, const KvHostPools& host, bool int8, int64_t n_blocks,
+                        const QsaShapes& s, void* stream);
+
+struct KvStreamCounters {
+    uint64_t misses = 0, lookups = 0, calls = 0;
+    bool overflow = false;
+};
+/// Synchronous read of the counters (debug and the end-of-request summary).
+KvStreamCounters kv_stream_counters(const KvStreamMap& m);
+
+/// Bytes of one block (page) of K and V together.
+uint64_t kv_block_bytes(const QsaShapes& s, bool int8);
+
+}  // namespace strata::kernels

include/strata/kernels/qsa.hpp

@@ -110,7 +110,9 @@ inline QsaShapes qsa_real_shapes() {
     QsaShapes s;
     s.n_head = 24; s.n_head_kv = 2; s.head_dim = 256; s.n_rot = 64;
     s.idx_n_head = 4; s.idx_dim = 128; s.idx_block = 4; s.idx_top_k = 2048;
-    s.page_size = 512;
+    // one page = one indexer block (4 cells): the granule KV streaming keeps resident (kv_stream.hpp). The readers
+    // resolve a row per cell anyway, so the page size costs them nothing (measured: identical output and speed).
+    s.page_size = 4;
     return s;
 }
 
@@ -353,8 +355,12 @@ void qsa_attend_step(const float* q, const uint16_t* k_scratch, const uint16_t*
                      const int32_t* step, int64_t max_ids, const QsaShapes& s, float* attn, float* weights,
                      void* stream);
 
-/// The cell's position comes from `step`, so this is the form a graph may contain.
+struct KvHostPools;   // kv_stream.hpp
+
+/// The cell's position comes from `step`, so this is the form a graph may contain. With a host copy (KV streaming)
+/// the cell is written there too, and to VRAM only if its block is resident.
 void kv_append_step(uint16_t* k_pool, uint16_t* v_pool, const int32_t* page_table, const int32_t* step,
-                    const float* kcur, const float* vcur, const QsaShapes& s, void* stream);
+                    const float* kcur, const float* vcur, const QsaShapes& s, void* stream,
+                    const KvHostPools* host = nullptr);
 
 }  // namespace strata::kernels

include/strata/kernels/sampler.hpp

@@ -1,12 +1,10 @@
 // include/strata/kernels/sampler.hpp - the sampler chain, host-callable (P2.S2).
 //
-//     penalties  ->  top_k  ->  top_p  ->  temperature  ->  pick
+//     penalties -> top_k -> min_p -> top_p -> temperature -> penalties -> pick
 //
-// The ORDER is specified in docs/sampling.md, transcribed from llama.cpp's own chain.  Two facts there are
-// easy to get backwards: temperature comes AFTER the truncation filters, and penalties come AFTER temperature.
-//
-// PENALTIES ARE NOT IMPLEMENTED YET: they need the token history, which is not in this signature, so a caller
-// passing a non-zero `penalty_last_n` gets a loud failure rather than a silently unpenalised sample.
+// The ORDER is specified in docs/sampling.md, transcribed from llama.cpp's own chain.  Three facts there are
+// easy to get backwards: temperature comes AFTER the truncation filters, penalties come AFTER temperature as
+// well as before the selection, and min_p cuts BEFORE top_p (both on the pre-temperature logits).
 #pragma once
 
 #include <cstdint>
@@ -16,6 +14,7 @@ namespace strata::kernels {
 struct SamplerParams {
     int top_k = 20;              // 0 disables the filter
     float top_p = 0.95f;         // 1.0 disables the filter
+    float min_p = 0.0f;          // 0 disables; keeps tokens with p >= min_p * p_max (a prefix of the top_k list)
     float temperature = 1.0f;    // <= 0 means greedy
     int min_keep = 1;            // top_p keeps at least this many
     int penalty_last_n = 0;      // 0 disables; the window over `history` to count occurrences in

include/strata/prefill/kernels.hpp

@@ -5,6 +5,8 @@
 // the kernel walks the chunk's tokens in order inside one launch.
 #pragma once
 
+#include "strata/kernels/kv_stream.hpp"
+
 #include <cstdint>
 
 namespace strata::prefill {
@@ -63,11 +65,14 @@ void gate_attn(const float* attn, const float* q_full, uint16_t* out16, int64_t
 /// K, V: [T, 2, 256].
 void kv_append(const float* K, const float* V, int64_t T, int64_t pos0, const int32_t* page_table, int64_t page_size,
                uint16_t* k_pool, uint16_t* v_pool, int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale,
-               void* stream);
+               void* stream, const strata::kernels::KvHostPools* host = nullptr,
+               const strata::kernels::KvHostPools* stage = nullptr);
 
 /// fp32 -> fp16 bits and fp32 -> bf16, n elements (the two activation images of the prompt GEMMs).
 void to_f16(const float* x, uint16_t* y, int64_t n, void* stream);
 void to_bf16(const float* x, uint16_t* y, int64_t n, void* stream);
+/// y = fp32(fp16(x)): what an FP16 store of x would read back (the FP16 indexer-key experiment)
+void round_f16(const float* x, float* y, int64_t n, void* stream);
 /// Expert blob -> FP16 (Q2_0 values are exact in FP16).
 void blob_dequant_f16(const uint8_t* blob, uint16_t* gu16, uint16_t* down16, void* stream);
 

include/strata/spec/draft_policy.hpp

@@ -0,0 +1,52 @@
+// include/strata/spec/draft_policy.hpp - per verify round: the MTP's window, or a lookup (suffix) window?
+//
+// The suffix drafter (prompt lookup) proposes the tokens that followed an earlier repeat of the context. Taken
+// whenever it proposes more than the MTP, it lost 2-8% on ordinary text: a long lookup window costs much more to
+// verify than the MTP's usual 3-4 tokens, and it is only worth that when enough of it is accepted. llama.cpp's
+// lookup decoding answers the same problem with confidence thresholds on its n-gram statistics; here the policy
+// learns both sides online and compares expected committed tokens per millisecond:
+//
+//   MTP     E = the mean tokens a window of that size has committed (EMA), at the measured cost of that size
+//   lookup  E(k) = 1 + q + q^2 + ... + q^k for k <= the proposal, q = the acceptance rate of lookup drafts whose
+//           match was about as long (4 buckets of match length, decayed counts), at the measured cost of k + 1
+//
+// and takes the lookup window only when its best E/cost beats the MTP's by `margin`. Costs are the measured round
+// times per window size (EMA; sizes not seen yet are scaled from seen ones by a prior shape), so the policy adapts
+// to the machine and the context length. It only chooses which drafts to verify: the output is unchanged.
+#pragma once
+
+#include <array>
+
+namespace strata::spec {
+
+class DraftPolicy {
+public:
+    static constexpr int kMaxT = 8;
+    static constexpr int kBuckets = 4;
+
+    explicit DraftPolicy(int max_t, double margin = 0.03);
+
+    struct Pick {
+        bool lookup = false;
+        int t = 1;                      // window size (1 + drafts)
+    };
+    /// `t_mtp`: the MTP's window; `lookup_k`: the lookup proposal's length (0 = none); `match`: its match length.
+    Pick choose(int t_mtp, int lookup_k, int match) const;
+    /// After the round: the window it used, the drafts accepted, and the round's time (verify + commit + draft).
+    void observe(bool lookup, int t, int accepted, int match, double round_ms);
+
+    double lookup_rate(int match) const;   // current q for a match length
+    double cost_ms(int t) const;           // measured or scaled round time of a window of t tokens
+
+private:
+    static int bucket(int match);
+    double mtp_tokens(int t) const;
+
+    int max_t_;
+    double margin_;
+    std::array<double, kMaxT + 1> cost_{}, cost_n_{};      // round ms by window size
+    std::array<double, kMaxT + 1> mtp_tok_{}, mtp_n_{};    // tokens committed by MTP windows of that size
+    std::array<double, kBuckets> ok_{}, bad_{};            // lookup drafts accepted / windows cut short, decayed
+};
+
+}  // namespace strata::spec

serve/server.py

@@ -10,6 +10,7 @@ data, http(s) URLs or local file paths) go through `strata-vision` (the model's
 embeddings (`GENI`).  JPEG/PNG/BMP/GIF go straight in; WebP, TIFF, AVIF, ... (agents like omp send WebP) are
 converted to PNG first with Pillow.
 Requests whose prompt plus max tokens exceed the engine's context are REJECTED with 400, never truncated.
+An unset (or 0, or -1) max tokens means "unlimited": whatever the prompt leaves of the context.
 
 The engine boundary is `Engine.generate(prompt_ids, max_new, sampling, cancel) -> iterator of token ids`.
 `StrataEngine` keeps one `strata --serve` process resident (weights, expert arena and VRAM tier load once) and
@@ -42,6 +43,7 @@ from serve.frontend import (ChatTemplate, Event, OutputParser, anthropic_to_mess
 
 IM_END = "<|im_end|>"
 IMAGE_PAD = "<|image_pad|>"
+CTX_SLACK = 8               # `strata --serve` rejects prompt + max_new + 8 > context: keep the same margin here
 
 
 # ------------------------------------------------------------------------------------------------ engines
@@ -71,7 +73,12 @@ class MockEngine:
 
 class StrataEngine:
     """The resident engine: `strata --serve` reads `GEN <max_new> <ids>` lines and streams `T <id>` lines, then
-    `DONE ...`.  Requests are serialized by the service's FIFO, so one pipe is enough."""
+    `DONE ...`.  Requests are serialized by the service's FIFO, so one pipe is enough.
+
+    Per-request sampling rides the same line as engine-side keys between max_new and the ids
+    (`temperature=F top_p=F top_k=N seed=N`, the engine's own spelling).  An absent temperature keeps the
+    engine's default, which is greedy; `temperature=0` means the same thing, so it is not forwarded.
+    """
 
     def __init__(self, exe: str, args: list[str], cwd: str | None = None, log: str | None = None,
                  env: dict | None = None):
@@ -103,11 +110,51 @@ class StrataEngine:
         self.last = {"generated": int(f[1]), "prompt_tokens": int(f[2]), "prompt_ms": float(f[3]),
                      "decode_ms": float(f[4]), "finish": f[5]}
 
+    @staticmethod
+    def sampling_keys(sampling: dict) -> str:
+        keys = ""
+        t = sampling.get("temperature")
+        if isinstance(t, (int, float)) and float(t) > 0.0:
+            keys += f" temperature={float(t)!r}"
+        tp = sampling.get("top_p")
+        if isinstance(tp, (int, float)) and float(tp) < 1.0:
+            keys += f" top_p={float(tp)!r}"
+        tk = sampling.get("top_k")
+        if isinstance(tk, int) and 1 <= tk <= 64:
+            keys += f" top_k={tk}"        # the engine's sampled path takes 1..64; outside it keeps its 20
+        mp = sampling.get("min_p")
+        if isinstance(mp, (int, float)) and 0.0 < float(mp) <= 1.0:
+            keys += f" min_p={float(mp)!r}"
+        rp = sampling.get("repetition_penalty")
+        rp_on = isinstance(rp, (int, float)) and float(rp) != 1.0
+        pf = sampling.get("frequency_penalty")
+        pf_on = isinstance(pf, (int, float)) and float(pf) != 0.0
+        pp = sampling.get("presence_penalty")
+        pp_on = isinstance(pp, (int, float)) and float(pp) != 0.0
+        if rp_on:
+            keys += f" penalty_repeat={float(rp)!r}"
+        if pf_on:
+            keys += f" penalty_freq={float(pf)!r}"
+        if pp_on:
+            keys += f" penalty_present={float(pp)!r}"
+        if rp_on or pf_on or pp_on:
+            # a penalty without a window counts over nothing: the engine's default is the last 64 tokens
+            pln = sampling.get("penalty_last_n")
+            if isinstance(pln, int) and not isinstance(pln, bool) and pln > 0:
+                keys += f" penalty_last_n={pln}"
+            else:
+                keys += " penalty_last_n=64"
+        seed = sampling.get("seed")
+        if isinstance(seed, int) and seed > 0:
+            keys += f" seed={seed}"
+        return keys
+
     def generate(self, ids, max_new, sampling, cancel, embeddings=None):
         """Yields token ids, and None as a heartbeat every 10 s while the engine is quiet (reading a long prompt):
         the HTTP layer turns it into an SSE comment, which keeps clients' watchdogs calm and notices a client that
         has gone.  A consumer that stops early (or `cancel`) makes the engine STOP, so it does not run to max_new."""
-        head = f"GENI {int(max_new)} {embeddings}" if embeddings else f"GEN {int(max_new)}"
+        head = f"GENI {int(max_new)} {embeddings}" if embeddings else \
+            f"GEN {int(max_new)}{self.sampling_keys(sampling or {}) if not embeddings else ''}"
         self.proc.stdin.write(f"{head} {','.join(str(int(t)) for t in ids)}\n")
         self.proc.stdin.flush()
         done = False
@@ -127,6 +174,10 @@ class StrataEngine:
                     if cancel.is_set():
                         return
                     yield int(line[2:])
+                elif line.startswith("PP "):             # prompt progress, one per chunk: also a heartbeat (the
+                    if cancel.is_set():                   # lines reset the 10 s wait, so without this a long prompt
+                        return                            # would send no keep-alives at all)
+                    yield None
                 elif line.startswith("DONE"):
                     self._parse_done(line)
                     done = True
@@ -307,8 +358,9 @@ class Detokenizer:
 
 class Service:
     def __init__(self, engine: Engine, tokenizer, template: ChatTemplate, model_name: str = "qwen3.8-flash-next",
-                 vision: Vision | None = None):
+                 vision: Vision | None = None, sampling_defaults: dict | None = None):
         self.engine, self.tok, self.template, self.model, self.vision = engine, tokenizer, template, model_name, vision
+        self.sampling_defaults = dict(sampling_defaults or {})   # the run config's `sampling` block
         self.fifo = threading.Lock()
         self.embeddings = threading.local()           # the current request's image embeddings file (GENI)
         self.api_key = ""                              # when set, /v1/* needs it (Bearer or x-api-key)
@@ -317,7 +369,9 @@ class Service:
         self.stop_ids = set(tokenizer.encode(IM_END, parse_special=True) +
                             tokenizer.encode("<|endoftext|>", parse_special=True))
 
-    def prepare(self, messages, tools, kwargs, max_new):
+    def prepare(self, messages, tools, kwargs, max_new=None):
+        """-> (ids, thinking, max_new). An unset or non-positive max_new (some clients send -1) means "unlimited":
+        the rest of the context."""
         prompt = self.template.render(messages, tools=tools, **kwargs)
         ids = self.tok.encode(prompt, parse_special=True)
         self.embeddings.path = None
@@ -348,10 +402,16 @@ class Service:
                 for path, _ in encoded:
                     f.write(path.read_bytes())
             self.embeddings.path = combined
-        if len(ids) + max_new > self.engine.max_context:
+        room = self.engine.max_context - CTX_SLACK - len(ids)
+        if max_new is None or max_new <= 0:
+            if room < 1:
+                raise ValueError(f"prompt ({len(ids)} tokens) leaves no room to answer in the context "
+                                 f"({self.engine.max_context}); requests are never truncated")
+            max_new = room
+        elif max_new > room:
             raise ValueError(f"prompt ({len(ids)} tokens) + max tokens ({max_new}) exceeds the context "
                              f"({self.engine.max_context}); requests are never truncated")
-        return ids, kwargs.get("enable_thinking", True) is not False
+        return ids, kwargs.get("enable_thinking", True) is not False, max_new
 
     def _note(self, n, evs):
         with self.status_lock:
@@ -370,7 +430,7 @@ class Service:
                     s["phase"] = "tool call complete"
                 s["tail"] = (s["tail"] + (ev.text or ""))[-600:]
 
-    def _progress(self, last_print, every=15.0):
+    def _progress(self, last_print, every=1.0):
         """A progress line in the server window every `every` seconds while a request runs."""
         now = time.time()
         if now - last_print < every:
@@ -388,8 +448,12 @@ class Service:
 
     def run(self, ids, thinking, tools, max_new, sampling, cancel) -> Iterator[tuple[str, object]]:
         """Yields ("event", Event) as text arrives, then ("done", {"finish": .., "completion_tokens": ..})."""
+        if self.sampling_defaults:     # config defaults; the request's own fields win (explicit 0 stays greedy)
+            req_values = {k: v for k, v in (sampling or {}).items() if v is not None}
+            sampling = {**self.sampling_defaults, **req_values}
         parser = OutputParser(thinking=thinking, tools=tools, stream_tools=True)
         detok, n, finish = Detokenizer(self.tok), 0, "length"
+        raw_ids = []                                    # every generated id (STRATA_DEBUG: dump raw model text)
         emb = getattr(self.embeddings, "path", None)
         with self.status_lock:
             self.status["queued"] += 1
@@ -402,35 +466,66 @@ class Service:
                 last_print = time.time()
                 gen = self.engine.generate(ids, max_new, sampling, cancel, embeddings=emb) if emb else \
                     self.engine.generate(ids, max_new, sampling, cancel)
-                for t in gen:
-                    if t is None:                       # heartbeat while the engine is quiet
+                try:
+                    for t in gen:
+                        if t is None:                   # heartbeat while the engine is quiet
+                            last_print = self._progress(last_print)
+                            yield "ping", None
+                            continue
+                        n += 1
+                        if t in self.stop_ids:
+                            finish = "stop"
+                            raw_ids.append(t)
+                            break
+                        raw_ids.append(t)
+                        evs = parser.feed(detok.push(t))
+                        self._note(n, evs)
                         last_print = self._progress(last_print)
-                        yield "ping", None
-                        continue
-                    n += 1
-                    if t in self.stop_ids:
-                        finish = "stop"
-                        break
-                    evs = parser.feed(detok.push(t))
-                    self._note(n, evs)
-                    last_print = self._progress(last_print)
-                    for ev in evs:
-                        yield "event", ev
-                if cancel.is_set():
-                    finish = "cancel"
+                        for ev in evs:
+                            yield "event", ev
+                    if cancel.is_set():
+                        finish = "cancel"
+                finally:
+                    gen.close()                         # STOP+drain to THIS request's DONE while still holding the
+                    #                                     fifo, so a stop-token break can't leave the shared engine
+                    #                                     queue mid-drain for the next request to read as its own DONE
+        except GeneratorExit:                           # the client disconnected mid-stream
+            finish = "disconnect"
+            raise
         finally:
             if emb:
                 Path(emb).unlink(missing_ok=True)
             with self.status_lock:
                 if self.status.get("busy"):
-                    el = time.time() - self.status.get("started", time.time())
-                    print(f"[strata] done: {n} tokens in {el:.0f} s ({finish})", flush=True)
+                    now = time.time()
+                    el = now - self.status.get("started", now)
+                    ft = self.status.get("first_token")
+                    rate = n / max(1e-6, now - ft) if ft else 0.0
+                    print(f"[strata] done: {n} tokens in {el:.0f} s ({rate:.1f} tok/s) "
+                          f"({finish}, cancel={cancel.is_set()})", flush=True)
+                    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
         for ev in parser.finish():
             yield "event", ev
         yield "done", {"finish": finish, "completion_tokens": n}
 
 
+def _debug_req(api, req, messages, tools, max_new, thinking, prompt_tokens):
+    """One compact line per request while diagnosing blank/empty turns. Set STRATA_DEBUG=1 to enable."""
+    if not os.environ.get("STRATA_DEBUG"):
+        return
+    last = messages[-1] if messages else {}
+    body = last.get("content")
+    if isinstance(body, list):
+        body = " ".join(p.get("text", "") for p in body if isinstance(p, dict))
+    preview = (str(body or "")[:80]).replace("\n", " ")
+    print(f"[strata] req {api}: msgs={len(messages)} tools={len(tools or [])} "
+          f"max_tokens_raw={req.get('max_tokens')!r}/{req.get('max_completion_tokens')!r} "
+          f"max_new={max_new} thinking={thinking} stream={bool(req.get('stream'))} "
+          f"prompt_tokens={prompt_tokens} last={last.get('role')!r}:{preview!r}", flush=True)
+
+
 # ------------------------------------------------------------------------------------------------ OpenAI
 def openai_chunks(svc: Service, req: dict, ids, thinking, tools, max_new, cancel):
     cid, created = "chatcmpl-" + uuid.uuid4().hex[:24], int(time.time())
@@ -667,8 +762,9 @@ def make_handler(svc: Service):
 
         def _openai(self, req):
             messages, tools, kw = openai_to_messages(req)
-            max_new = int(req.get("max_completion_tokens") or req.get("max_tokens") or 1024)
-            ids, thinking = svc.prepare(messages, tools, kw, max_new)
+            max_new = int(req.get("max_completion_tokens") or req.get("max_tokens") or 0)   # 0/-1: the rest of the context
+            ids, thinking, max_new = svc.prepare(messages, tools, kw, max_new)
+            _debug_req("openai", req, messages, tools, max_new, thinking, len(ids))
             cancel = threading.Event()
             chunks = openai_chunks(svc, req, ids, thinking, tools, max_new, cancel)
             if not req.get("stream"):
@@ -688,8 +784,9 @@ def make_handler(svc: Service):
 
         def _anthropic(self, req):
             messages, tools, kw = anthropic_to_messages(req)
-            max_new = int(req.get("max_tokens") or 1024)
-            ids, thinking = svc.prepare(messages, tools, kw, max_new)
+            max_new = int(req.get("max_tokens") or 0)                  # 0/-1: the rest of the context
+            ids, thinking, max_new = svc.prepare(messages, tools, kw, max_new)
+            _debug_req("anthropic", req, messages, tools, max_new, thinking, len(ids))
             cancel = threading.Event()
             events = anthropic_events(svc, req, ids, thinking, tools, max_new, cancel)
             if not req.get("stream"):
@@ -723,6 +820,59 @@ def serve(svc: Service, host="127.0.0.1", port=8095) -> ThreadingHTTPServer:
     return httpd
 
 
+def sampling_defaults_from_config(cfg: dict) -> dict:
+    """The run config's optional `sampling` block: defaults for the sampling fields a request leaves out, so
+    a plain client gets configured sampling instead of greedy.  Supported: temperature, top_p, top_k, min_p,
+    presence_penalty, repetition_penalty, frequency_penalty, penalty_last_n, seed.  The request's own fields
+    always win - an explicit temperature=0 still means greedy, a field set to null falls back to the default.
+    A bad value refuses to start the server (a typo'd config should not quietly change sampling); unknown keys
+    are named at startup and ignored."""
+    out = {}
+    for key, value in (cfg.get("sampling") or {}).items():
+        if value is None:
+            continue
+        number = isinstance(value, (int, float)) and not isinstance(value, bool)
+        if key == "temperature":
+            if not number or value < 0:
+                raise SystemExit(f"[strata] config sampling.temperature={value!r}: expected a number >= 0 (0 = greedy)")
+            out[key] = float(value)
+        elif key == "top_p":
+            if not number or not 0 < value <= 1:
+                raise SystemExit(f"[strata] config sampling.top_p={value!r}: expected 0 < top_p <= 1")
+            out[key] = float(value)
+        elif key == "min_p":
+            if not number or not 0 <= value <= 1:
+                raise SystemExit(f"[strata] config sampling.min_p={value!r}: expected 0 <= min_p <= 1")
+            out[key] = float(value)
+        elif key == "top_k":
+            if not number or value != int(value) or not 1 <= value <= 64:
+                raise SystemExit(f"[strata] config sampling.top_k={value!r}: the sampled path takes an integer 1..64")
+            out[key] = int(value)
+        elif key == "presence_penalty":
+            if not number or value < 0:
+                raise SystemExit(f"[strata] config sampling.presence_penalty={value!r}: expected a number >= 0")
+            out[key] = float(value)
+        elif key == "frequency_penalty":
+            if not number or value < 0:
+                raise SystemExit(f"[strata] config sampling.frequency_penalty={value!r}: expected a number >= 0")
+            out[key] = float(value)
+        elif key == "repetition_penalty":
+            if not number or value <= 0:
+                raise SystemExit(f"[strata] config sampling.repetition_penalty={value!r}: expected a number > 0 (1 = off)")
+            out[key] = float(value)
+        elif key == "penalty_last_n":
+            if not number or value != int(value) or value < 0:
+                raise SystemExit(f"[strata] config sampling.penalty_last_n={value!r}: expected a non-negative integer")
+            out[key] = int(value)
+        elif key == "seed":
+            if not number or value != int(value) or value <= 0:
+                raise SystemExit(f"[strata] config sampling.seed={value!r}: expected a positive integer")
+            out[key] = int(value)
+        else:
+            print(f"[strata] config sampling.{key}={value!r}: unknown key, ignored", flush=True)
+    return out
+
+
 def main() -> int:
     ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
     ap.add_argument("--engine", choices=["mock", "strata"], default="mock")
@@ -763,6 +913,10 @@ def main() -> int:
             ap.error("--engine strata needs --config")
         vision = None
         env = child_env(cfg)
+        sampling_defaults = sampling_defaults_from_config(cfg)
+        if sampling_defaults:
+            pretty = ", ".join(f"{k}={v}" for k, v in sampling_defaults.items())
+            print(f"[strata] sampling defaults from the config: {pretty}", flush=True)
         if cfg.get("vision"):
             print("loading the vision encoder ...", flush=True)
             vision = Vision(cfg["vision"], log=open(cfg["log"], "a", encoding="utf-8") if cfg.get("log") else None,
@@ -770,11 +924,12 @@ def main() -> int:
         print("loading the model (the first start takes a minute or two) ...", flush=True)
         engine = StrataEngine(cfg["exe"], cfg["args"], cwd=cfg.get("cwd"), log=cfg.get("log"), env=env)
     else:
-        engine, vision = MockEngine(tok, a.script), None
+        engine, vision, sampling_defaults = MockEngine(tok, a.script), None, {}
     # the model's own chat template (exported with its tokenizer), else the original model's
     tpl = tpath / "chat_template.jinja"
     svc = Service(engine, tok, ChatTemplate(tpl if tpl.exists() else ROOT / "serve/chat_template.jinja"),
-                  model_name=cfg.get("model_name", "qwen3.8-flash-next"), vision=vision)
+                  model_name=cfg.get("model_name", "qwen3.8-flash-next"), vision=vision,
+                  sampling_defaults=sampling_defaults)
     svc.api_key = a.api_key or cfg.get("api_key", "")
     httpd = serve(svc, host=a.host, port=a.port)
     print(f"ready: http://{a.host}:{a.port}/v1  (OpenAI: /v1/chat/completions, Anthropic: /v1/messages, "

serve/test_server.py

@@ -0,0 +1,130 @@
+"""serve/test_server.py - the max tokens budget over both APIs, against the mock engine (no GPU, no pack).
+
+    python -m unittest serve.test_server -v
+"""
+from __future__ import annotations
+
+import json
+import os
+import sys
+import unittest
+import urllib.error
+import urllib.request
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1]))
+from serve.frontend import ChatTemplate  # noqa: E402
+from serve.server import CTX_SLACK, ByteTokenizer, MockEngine, Service, serve  # noqa: E402
+
+ROOT = Path(__file__).resolve().parents[1]
+CTX = 4096
+ANSWER = "x" * 2000                              # longer than the old 1024 fallback: one token per byte
+
+
+class RecordingEngine(MockEngine):
+    def generate(self, ids, max_new, sampling, cancel, embeddings=None):
+        self.last_max_new = max_new
+        yield from super().generate(ids, max_new, sampling, cancel, embeddings)
+
+
+class MaxTokens(unittest.TestCase):
+    @classmethod
+    def setUpClass(cls):
+        tok = ByteTokenizer()
+        cls.engine = RecordingEngine(tok, "</think>\n\n" + ANSWER, max_context=CTX)
+        cls.svc = Service(cls.engine, tok, ChatTemplate(ROOT / "serve/chat_template.jinja"))
+        cls.httpd = serve(cls.svc, port=0)
+        cls.base = f"http://127.0.0.1:{cls.httpd.server_address[1]}"
+
+    @classmethod
+    def tearDownClass(cls):
+        cls.httpd.shutdown()
+        cls.httpd.server_close()
+
+    def post(self, path, body):
+        req = urllib.request.Request(self.base + path, data=json.dumps(body).encode(),
+                                     headers={"Content-Type": "application/json"})
+        try:
+            with urllib.request.urlopen(req, timeout=30) as r:
+                return r.status, json.loads(r.read())
+        except urllib.error.HTTPError as e:
+            with e:
+                return e.code, json.loads(e.read())
+
+    def call(self, api, text="hi", **budget):
+        """-> (status, body, prompt tokens, completion tokens); `budget` is merged into the request as given."""
+        msgs = [{"role": "user", "content": text}]
+        if api == "openai":
+            s, b = self.post("/v1/chat/completions", {"model": "m", "messages": msgs, **budget})
+            u = b.get("usage", {})
+            return s, b, u.get("prompt_tokens"), u.get("completion_tokens")
+        s, b = self.post("/v1/messages", {"model": "m", "messages": msgs, **budget})
+        u = b.get("usage", {})
+        return s, b, u.get("input_tokens"), u.get("output_tokens")
+
+    def test_unset_budget_is_the_rest_of_the_context(self):
+        cases = {"openai": [{"max_tokens": -1}, {"max_tokens": 0}, {}, {"max_tokens": None},
+                            {"max_completion_tokens": -1}, {"max_completion_tokens": None, "max_tokens": None}],
+                 "anthropic": [{"max_tokens": -1}, {"max_tokens": 0}, {}, {"max_tokens": None}]}
+        for api, budgets in cases.items():
+            for budget in budgets:
+                with self.subTest(api=api, budget=budget):
+                    s, b, pt, ct = self.call(api, **budget)
+                    self.assertEqual(s, 200, b)
+                    self.assertEqual(self.engine.last_max_new, CTX - CTX_SLACK - pt)
+                    self.assertGreater(ct, 1024)          # the whole answer, not cut at the old 1024 fallback
+
+    def test_explicit_budget_is_honoured(self):
+        for api, budget in [("openai", {"max_tokens": 50}), ("openai", {"max_completion_tokens": 50}),
+                            ("openai", {"max_completion_tokens": 50, "max_tokens": 9}),
+                            ("anthropic", {"max_tokens": 50}), ("openai", {"max_tokens": 1500}),
+                            ("anthropic", {"max_tokens": 1500})]:
+            with self.subTest(api=api, budget=budget):
+                want = budget.get("max_completion_tokens") or budget["max_tokens"]
+                s, b, _, ct = self.call(api, **budget)
+                self.assertEqual(s, 200, b)
+                self.assertEqual(self.engine.last_max_new, want)
+                self.assertEqual(ct, want)
+
+    def test_explicit_budget_over_the_context_is_rejected(self):
+        for api in ("openai", "anthropic"):
+            with self.subTest(api=api):
+                s, b, _, _ = self.call(api, max_tokens=CTX)
+                self.assertEqual(s, 400)
+                self.assertIn("exceeds the context", b["error"]["message"])
+
+    def test_unset_budget_with_a_near_full_prompt(self):
+        _, _, pt0, _ = self.call("openai", max_tokens=1)
+        overhead = pt0 - len("hi")                  # the template's tokens around the user text
+        for api in ("openai", "anthropic"):
+            _, _, pa, _ = self.call(api, max_tokens=1)
+            over = pa - pt0                          # the Anthropic template may differ slightly
+            with self.subTest(api=api, room=5):     # a few tokens left: the budget is exactly those
+                text = "y" * (CTX - CTX_SLACK - overhead - over - 5)
+                s, b, pt, ct = self.call(api, text=text, max_tokens=-1)
+                self.assertEqual(s, 200, b)
+                self.assertEqual(self.engine.last_max_new, 5)
+                self.assertEqual(ct, 5)
+            with self.subTest(api=api, room=0):     # nothing left: rejected, not truncated
+                text = "y" * (CTX - CTX_SLACK - overhead - over)
+                s, b, _, _ = self.call(api, text=text)
+                self.assertEqual(s, 400, b)
+                self.assertIn("no room to answer", b["error"]["message"])
+
+    def test_debug_log_shows_the_resolved_budget(self):
+        import contextlib
+        import io
+        os.environ["STRATA_DEBUG"] = "1"
+        try:
+            for api in ("openai", "anthropic"):
+                with self.subTest(api=api):
+                    out = io.StringIO()
+                    with contextlib.redirect_stdout(out):
+                        _, _, pt, _ = self.call(api, max_tokens=-1)
+                    self.assertIn(f"max_new={CTX - CTX_SLACK - pt} ", out.getvalue())
+        finally:
+            del os.environ["STRATA_DEBUG"]
+
+
+if __name__ == "__main__":
+    unittest.main()

setup.py

@@ -19,7 +19,7 @@ What the first run does (each step is skipped when it is already done):
   6. prepares the model for Strata and fetches the MTP draft layer (~5 GB, from the original Qwen checkpoint)
   7. writes run-<model>.bat / run-<model>.sh and starts the model
 
-Options: --family qwen|swift, --model Q2_0|IQ2_XS|IQ3_XXS, --context 32768, --vision yes|no|gpu|cpu, --port 8080, --yes (recommended
+Options: --family qwen|swift, --model Q2_0|IQ2_XS|IQ3_XXS|IQ3_S, --context 32768, --vision yes|no|gpu|cpu, --port 8080, --yes (recommended
 answers, no questions), --setup (install another model / change settings instead of starting), --no-start,
 --models-dir DIR, --gguf-dir DIR (use GGUF files you already have), --build (compile instead of the ready-made
 engine), --check (only check this PC).
@@ -54,15 +54,19 @@ 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, 2)                 # split models (Swift 1.5), STOP, cache sized after the slots are written
+MIN_ENGINE = (0, 1, 7)                 # per-request sampling on the GEN line (the server sends it), v0.1.7
 PY_PACKAGES = ["numpy", "jinja2", "regex", "pyyaml", "tqdm", "requests", "cmake", "ninja", "pillow"]
 
 MODELS = {
     "Q2_0": {"about": "2-bit, the fastest", "download_gb": 66.4, "ram_gb": 48, "arena_gb": 34.0},
     "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, the best quality, slower (more CPU work per token)", "download_gb": 75.8,
+    "IQ3_XXS": {"about": "3-bit i-quant, better quality, slower (more CPU work per token)", "download_gb": 75.8,
                 "ram_gb": 60, "arena_gb": 42.9},
+    # the original model only (Swift 1.5 has no IQ3_S): matches the full BF16 model on the published benchmarks
+    "IQ3_S": {"about": "3.5-bit i-quant, the best quality (matches the full model), the slowest; needs a 64 GB PC "
+                       "with little else running", "download_gb": 83.6, "ram_gb": 62, "arena_gb": 50.3,
+              "families": ("qwen",)},
 }
 CONTEXTS = [8192, 32768, 65536, 131072, 262144]
 # The model families: the same architecture, weights in the same three GSQ-RCO sizes, different files.
@@ -456,6 +460,42 @@ def get_prebuilt(url_base, gpu, vision) -> Path | None:
     return eng
 
 
+def update_installed_engine(url_base) -> None:
+    """An installed ready-made engine older than MIN_ENGINE is replaced before the model starts, so a plain
+    START-HERE.bat on an existing install picks up a new release.  If that cannot happen (no internet, the model
+    still running, no ready-made engine for this GPU) the installed engine is kept and starts as before."""
+    eng = ROOT / "engine"
+    info = eng / "BUILD.json"
+    if not info.exists() or not (eng / EXE).exists():
+        return
+    meta_text = info.read_text()
+    meta = json.loads(meta_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:
+        return
+    try:                                               # a running engine cannot be replaced (Windows keeps it locked)
+        for x in (EXE, VEXE):
+            if (eng / x).exists():
+                with open(eng / x, "r+b"):
+                    pass
+    except OSError:
+        warn(f"engine {meta.get('version')} is in use: close the model window and run this again to update it")
+        return
+    gpu = gpu_info()
+    new = None
+    if gpu is not None:
+        try:
+            new = get_prebuilt(url_base, gpu, "gpu")
+        except Exception as e:                         # a failed download must not stop the model from starting
+            warn(f"updating the engine failed ({e})")
+    if new is None:
+        if not info.exists():
+            info.write_text(meta_text)                 # get_prebuilt drops it before downloading: put it back
+        warn(f"could not update the engine: starting the installed {meta.get('version')}")
+        return
+    pip_install(CUDA_WHEELS, "NVIDIA CUDA libraries (cuBLAS, CUDA runtime; ~0.4 GB)")
+
+
 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()
@@ -627,6 +667,8 @@ def main() -> int:
     # ---- 0. already installed: just start it
     have = installed_configs()
     if have and not (a.setup or a.model or a.family or a.check or a.no_start):
+        if not a.build:
+            update_installed_engine(a.prebuilt)
         if len(have) == 1:
             return start(have[0], None)
         say()
@@ -654,7 +696,14 @@ def main() -> int:
         warn("less than 12 GB of VRAM: Strata will run, but most experts stay on the CPU and it will be slow")
     ram = ram_gb()
     cpu, avx2, avx512 = cpu_info()
-    ok(f"RAM: {ram:.0f} GB")
+    need = min(d["ram_gb"] for d in MODELS.values())
+    if ram < need - 4 and not a.check:
+        # every model keeps ALL its experts in RAM (34+ GB); VRAM only holds a copy of the most-used ones, so a
+        # bigger GPU does not lower this
+        fail(f"RAM: {ram:.0f} GB - the smallest model (Q2_0 / IQ2_XS) needs about {need} GB",
+             "Strata keeps all of the model's experts in RAM (34-50 GB, whatever the GPU) and the GPU holds a copy "
+             "of the most-used ones: it needs 48 GB of RAM or more")
+    ok(f"RAM: {ram:.0f} GB" if ram >= need - 4 else f"RAM: {ram:.0f} GB (less than the {need} GB the smallest model needs)")
     ok(f"CPU: {cpu} ({'AVX-512' if avx512 else 'AVX2' if avx2 else 'no AVX2'})")
     if not avx2:
         fail("this CPU has no AVX2; Strata needs at least AVX2")
@@ -681,13 +730,15 @@ def main() -> int:
     if fam.get("license"):
         say(f"  Its license: {fam['license']}")
     say()
-    names = list(MODELS)
+    names = [m for m in MODELS if family in MODELS[m].get("families", FAMILIES)]
+    if a.model and a.model not in names:
+        fail(f"{fam['title']} has no {a.model} model file", "choose one of: " + ", ".join(names))
     for i, m in enumerate(names, 1):
         d = MODELS[m]
         fit = "" if ram >= d["ram_gb"] else f"   <- needs {d['ram_gb']} GB RAM, you have {ram:.0f}"
         say(f"  {i}) {m:8s} {d['about']}; download {d['download_gb']:.0f} GB, uses ~{d['arena_gb']:.0f} GB of RAM{fit}")
-    rec = "3" if ram >= 60 else "1"
-    model = a.model or names[int(ask("Which size?", ["1", "2", "3"], rec, a.yes)) - 1]
+    rec = str(names.index("IQ3_XXS") + 1) if ram >= 60 else "1"
+    model = a.model or names[int(ask("Which size?", [str(i) for i in range(1, len(names) + 1)], rec, a.yes)) - 1]
     if ram < MODELS[model]["ram_gb"] - 4:
         fail(f"{model} needs about {MODELS[model]['ram_gb']} GB of RAM; this PC has {ram:.0f} GB",
              "choose Q2_0 or IQ2_XS, or add RAM")
@@ -703,8 +754,9 @@ def main() -> int:
         for i, c in enumerate(CONTEXTS, 1):
             say(f"  {i}) {c // 1024}K tokens" + ("   (recommended for your GPU)" if c == rec_ctx else ""))
         ctx = CONTEXTS[int(ask("Context?", [str(i) for i in range(1, 6)], str(CONTEXTS.index(rec_ctx) + 1), a.yes)) - 1]
-    if model == "IQ3_XXS" and ram < 90 and ctx > 131072:
-        warn("IQ3_XXS with a 262K context needs more than 64 GB of RAM (43 GB of experts + the context): using 128K")
+    if model in ("IQ3_XXS", "IQ3_S") and ram < 90 and ctx > 131072:
+        warn(f"{model} with a 262K context needs more than 64 GB of RAM ({MODELS[model]['arena_gb']:.0f} GB of experts "
+             "+ the context): using 128K")
         ctx = 131072
     ok(f"context: {ctx} tokens")
     if a.vision:
@@ -818,6 +870,13 @@ def main() -> int:
             "--max-context", str(ctx)]
     if ctx > 8192:
         args += ["--kv", "int8"]
+    # 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 (1.7 GB at 128K, 3.4 GB at 262K), so only when that fits.
+    kv_ram_gb = ctx * 13728 / 1e9
+    if ctx >= 65536 and ram >= MODELS[model]["ram_gb"] + kv_ram_gb + 1:
+        args += ["--kv-resident", "32768"]
+        ok(f"KV streaming on: the context's KV cache lives in RAM ({kv_ram_gb:.1f} GB), more experts fit in VRAM")
     if vision != "none":
         args += ["--vision", "--vram-reserve-mib", str(VISION[vision]["reserve_mib"])]
     cfg = {"exe": str(eng / EXE), "args": args, "cwd": str(ROOT), "tokenizer": str(pack / "tokenizer"),

src/core/layer.cpp

@@ -27,6 +27,8 @@
 #include "strata/kernels/qsa_decode_attn.hpp"
 #include "strata/kernels/fused_gdn.hpp"
 #include "strata/kernels/qsa_select.hpp"
+#include <algorithm>
+#include <cstdlib>
 #include <exception>
 #include <cuda_runtime.h>
 #include <cmath>
@@ -475,39 +477,207 @@ struct Cursor {    uint8_t* p;    uint64_t used = 0;    template <typename T>
 // dispatch too and the two must not be able to drift.
 }
 // namespace
-uint64_t qsa_state_bytes(const ModelGeometry& g, int64_t max_cells, bool with_rope) {    const QsaShapes s = qsa_shapes(g);    const int64_t pages = (max_cells + s.page_size - 1) / s.page_size;    uint64_t n = 0;    n += g_kv_int8 ? (uint64_t) pages * s.page_size * strata::kernels::kv_q8_bytes_per_cell(s) + 64                    : (uint64_t) pages * s.n_head_kv * s.page_size * s.head_dim * 2 * 2;
-// k_pool + v_pool, fp16
-n += (uint64_t) pages * 4;
-// page_table
-n += (uint64_t) (s.idx_block - 1) * s.idx_dim * 4;
-// tail
-n += (uint64_t) s.idx_dim * 4;
-// dead
-n += (uint64_t) ((max_cells / s.idx_block) + 2) * s.idx_dim * 4;
-// pooled
-n += 16;
-// block_pos
-if (with_rope) n += (uint64_t) max_cells * (s.n_rot / 2) * 4 * 2;
-// cos + sin tables
-n += strata::kernels::qsa_step_bytes() + 16; // counts and aligned attention status
-// the per-token counts
-n += (uint64_t) s.n_head * 4;
-// pos_dev
-return align_up16(n) + 256;}
-uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, QsaState& st,                        const QsaState* share_rope) {    const QsaShapes s = qsa_shapes(g);    const int64_t pages = (max_cells + s.page_size - 1) / s.page_size;    Cursor c{(uint8_t*) base};    st.kv_int8 = g_kv_int8;    if (g_kv_int8) {        const uint64_t cells = (uint64_t) pages * s.n_head_kv * s.page_size;        st.k_q = c.take<int8_t>(cells * s.head_dim);        st.v_q = c.take<int8_t>(cells * s.head_dim);        st.k_scale = c.take<uint16_t>(cells * (s.head_dim / strata::kernels::KV_Q8_GROUP));        st.v_scale = c.take<uint16_t>(cells * (s.head_dim / strata::kernels::KV_Q8_GROUP));    } else {        st.k_pool = c.take<uint16_t>((uint64_t) pages * s.n_head_kv * s.page_size * s.head_dim);        st.v_pool = c.take<uint16_t>((uint64_t) pages * s.n_head_kv * s.page_size * s.head_dim);    }    st.page_table = c.take<int32_t>((uint64_t) pages);    st.n_pages = pages;    st.max_cells = max_cells;    st.idx_tail = c.take<float>((uint64_t) (s.idx_block - 1) * s.idx_dim);    st.idx_dead = c.take<float>((uint64_t) s.idx_dim);    st.idx_pooled = c.take<float>((uint64_t) ((max_cells / s.idx_block) + 2) * s.idx_dim);    st.idx_block_pos = c.take<int32_t>(1);    if (share_rope != nullptr) {        st.cos_tab = share_rope->cos_tab;        st.sin_tab = share_rope->sin_tab;    } else {        st.cos_tab = c.take<float>((uint64_t) max_cells * (s.n_rot / 2));        st.sin_tab = c.take<float>((uint64_t) max_cells * (s.n_rot / 2));    }    st.step = c.take<int32_t>((uint64_t) strata::kernels::kStepCount); st.attention_status = c.take<int32_t>(1);    st.pos_dev = c.take<int32_t>((uint64_t) s.n_head);
-// the pinned staging the uploads copy FROM - see the note on `host_step` in the header
-if (cudaHostAlloc((void**) &st.host_step, strata::kernels::qsa_step_bytes() + sizeof(int32_t), cudaHostAllocMapped | cudaHostAllocPortable) !=            cudaSuccess ||        cudaHostAlloc((void**) &st.host_pos, (size_t) s.n_head * 4, cudaHostAllocMapped | cudaHostAllocPortable) != cudaSuccess) {        return 0;
-// the caller sees a zero byte count; a half-built state is worse than none
+namespace {
+// KV streaming (docs/kv-streaming-design.md): 0 keeps every cell in VRAM.
+int64_t g_kv_resident = 0;
+uint64_t g_kv_host_bytes = 0;
+
+/// How one state holds its K/V: `mode` as in QsaState::kv_mode, `slots` VRAM pages of `pages` logical ones.
+struct KvPlan {
+    int mode = 0;
+    int64_t pages = 0, slots = 0, pooled_rows = 0;
+};
+KvPlan kv_plan(const QsaShapes& s, int64_t max_cells, int64_t ring_cells) {
+    KvPlan p;
+    p.pages = (max_cells + s.page_size - 1) / s.page_size;
+    p.slots = p.pages;
+    p.pooled_rows = max_cells / s.idx_block + 2;
+    if (g_kv_resident <= 0 || ring_cells < 0) return p;   // ring_cells < 0: always fully resident
+    // A/B only: STRATA_KV_RING_OFF keeps the drafter fully resident, STRATA_KV_MAIN_OFF the main layers
+    static const bool ring_off = std::getenv("STRATA_KV_RING_OFF") != nullptr;
+    static const bool main_off = std::getenv("STRATA_KV_MAIN_OFF") != nullptr;
+    if (ring_cells > 0 ? ring_off : main_off) return p;
+    if (ring_cells > 0) {
+        const int64_t r = (ring_cells + s.page_size - 1) / s.page_size;
+        if (r < p.pages) { p.mode = 2; p.slots = r; p.pooled_rows = 2; }   // the drafter has no indexer
+    } else {
+        const int64_t r = (std::max(g_kv_resident, qsa_kv_resident_min()) + s.page_size - 1) / s.page_size;
+        if (r < p.pages) { p.mode = 1; p.slots = r; }
+    }
+    return p;
 }
-st.host_step[strata::kernels::kStepCount] = 0;
-// THE ROPE TABLE IS BUILT ON THE HOST IN FLOAT64 and uploaded once, because the reference computes its
-// frequencies in float64 and reproducing that on device means double-precision `pow`/`cos` that need not
-// agree with the host's libm.  A table shorter than the sequence would have the rotation read past it.
-if (share_rope == nullptr) {    std::vector<float> hc((size_t) max_cells * (s.n_rot / 2)), hs((size_t) max_cells * (s.n_rot / 2));    strata::kernels::build_rope_table((int) s.n_rot, strata::kernels::qsa_freq_base(), (int) max_cells,                                      hc.data(), hs.data());    cudaMemcpy(st.cos_tab, hc.data(), hc.size() * 4, cudaMemcpyHostToDevice);    cudaMemcpy(st.sin_tab, hs.data(), hs.size() * 4, cudaMemcpyHostToDevice);    }
-// the page table starts as the IDENTITY, which is the simplest legal mapping and what a caller that does
-// not page at all wants.  A real allocator re-points it; nothing outside `kv_append`/`kv_gather` may care.
-std::vector<int32_t> tab((size_t) pages);    for (int64_t i = 0; i < pages; ++i) tab[(size_t) i] = (int32_t) i;    cudaMemcpy(st.page_table, tab.data(), tab.size() * 4, cudaMemcpyHostToDevice);    return c.used;}
-void qsa_state_zero(const QsaState& st, const ModelGeometry& g, void* stream) {    const QsaShapes s = qsa_shapes(g);    cudaStream_t cs = (cudaStream_t) stream;    if (st.kv_int8) {        const size_t cells = (size_t) st.n_pages * s.n_head_kv * s.page_size;        cudaMemsetAsync(st.k_q, 0, cells * s.head_dim, cs);        cudaMemsetAsync(st.v_q, 0, cells * s.head_dim, cs);        cudaMemsetAsync(st.k_scale, 0, cells * (s.head_dim / strata::kernels::KV_Q8_GROUP) * 2, cs);        cudaMemsetAsync(st.v_scale, 0, cells * (s.head_dim / strata::kernels::KV_Q8_GROUP) * 2, cs);    } else {    cudaMemsetAsync(st.k_pool, 0, (size_t) st.n_pages * s.n_head_kv * s.page_size * s.head_dim * 2, cs);    cudaMemsetAsync(st.v_pool, 0, (size_t) st.n_pages * s.n_head_kv * s.page_size * s.head_dim * 2, cs);    }    cudaMemsetAsync(st.idx_tail, 0, (size_t) (s.idx_block - 1) * s.idx_dim * 4, cs);    cudaMemsetAsync(st.idx_dead, 0, (size_t) s.idx_dim * 4, cs);    cudaMemsetAsync(st.idx_pooled, 0, (size_t) ((st.max_cells / s.idx_block) + 2) * s.idx_dim * 4, cs);}
+uint64_t kv_pool_bytes(const QsaShapes& s, int64_t pages) {
+    return g_kv_int8 ? (uint64_t) pages * s.page_size * strata::kernels::kv_q8_bytes_per_cell(s) + 64
+                     : (uint64_t) pages * s.n_head_kv * s.page_size * s.head_dim * 2 * 2;
+}
+}  // namespace
+
+void qsa_set_kv_resident(int64_t cells) { g_kv_resident = cells > 0 ? cells : 0; }
+int64_t qsa_kv_resident() { return g_kv_resident; }
+int64_t qsa_kv_resident_min() { return 20480; }
+uint64_t qsa_kv_host_bytes() { return g_kv_host_bytes; }
+
+uint64_t qsa_state_bytes(const ModelGeometry& g, int64_t max_cells, bool with_rope, int64_t ring_cells) {
+    const QsaShapes s = qsa_shapes(g);
+    const KvPlan p = kv_plan(s, max_cells, ring_cells);
+    uint64_t n = 0;
+    n += kv_pool_bytes(s, p.slots) + 4 * 16;                                   // K/V pools (the VRAM slots)
+    n += (uint64_t) p.pages * 4;                                               // page_table
+    if (p.mode == 1) n += strata::kernels::kv_stream_map_bytes(p.slots) + 6 * 16;   // the residency map
+    n += (uint64_t) (s.idx_block - 1) * s.idx_dim * 4;                         // tail
+    n += (uint64_t) s.idx_dim * 4;                                             // dead
+    n += (uint64_t) p.pooled_rows * s.idx_dim * 4;                             // pooled
+    n += 16;                                                                   // block_pos
+    if (with_rope) n += (uint64_t) max_cells * (s.n_rot / 2) * 4 * 2;          // cos + sin tables
+    n += strata::kernels::qsa_step_bytes() + 16;                               // counts and aligned attention status
+    n += (uint64_t) s.n_head * 4;                                              // pos_dev
+    return align_up16(n) + 256;
+}
+
+uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, QsaState& st,
+                        const QsaState* share_rope, int64_t ring_cells) {
+    const QsaShapes s = qsa_shapes(g);
+    const KvPlan p = kv_plan(s, max_cells, ring_cells);
+    const int64_t pages = p.pages;
+    Cursor c{(uint8_t*) base};
+    st.kv_int8 = g_kv_int8;
+    st.kv_mode = p.mode;
+    st.n_slots = p.slots;
+    const uint64_t rows = (uint64_t) p.slots * s.n_head_kv * s.page_size;   // VRAM rows: the slots
+    if (g_kv_int8) {
+        st.k_q = c.take<int8_t>(rows * s.head_dim);
+        st.v_q = c.take<int8_t>(rows * s.head_dim);
+        st.k_scale = c.take<uint16_t>(rows * (s.head_dim / strata::kernels::KV_Q8_GROUP));
+        st.v_scale = c.take<uint16_t>(rows * (s.head_dim / strata::kernels::KV_Q8_GROUP));
+    } else {
+        st.k_pool = c.take<uint16_t>(rows * s.head_dim);
+        st.v_pool = c.take<uint16_t>(rows * s.head_dim);
+    }
+    st.page_table = c.take<int32_t>((uint64_t) pages);
+    st.n_pages = pages;
+    st.max_cells = max_cells;
+    st.map = strata::kernels::KvStreamMap{};
+    st.map.page_table = st.page_table;
+    st.map.n_blocks = pages;
+    st.map.n_slots = p.slots;
+    if (p.mode == 1) {
+        st.map.slot_block = c.take<int32_t>((uint64_t) p.slots);
+        st.map.slot_stamp = c.take<int32_t>((uint64_t) p.slots);
+        st.map.slot_ref = c.take<int32_t>((uint64_t) p.slots);
+        st.map.miss_block = c.take<int32_t>((uint64_t) p.slots);
+        st.map.miss_slot = c.take<int32_t>((uint64_t) p.slots);
+        st.map.ctl = c.take<int32_t>((uint64_t) strata::kernels::kKvCtlInts);
+    }
+    st.idx_tail = c.take<float>((uint64_t) (s.idx_block - 1) * s.idx_dim);
+    st.idx_dead = c.take<float>((uint64_t) s.idx_dim);
+    st.idx_pooled_rows = p.pooled_rows;
+    st.idx_pooled = c.take<float>((uint64_t) p.pooled_rows * s.idx_dim);
+    st.idx_block_pos = c.take<int32_t>(1);
+    if (share_rope != nullptr) {
+        st.cos_tab = share_rope->cos_tab;
+        st.sin_tab = share_rope->sin_tab;
+    } else {
+        st.cos_tab = c.take<float>((uint64_t) max_cells * (s.n_rot / 2));
+        st.sin_tab = c.take<float>((uint64_t) max_cells * (s.n_rot / 2));
+    }
+    st.step = c.take<int32_t>((uint64_t) strata::kernels::kStepCount);
+    st.attention_status = c.take<int32_t>(1);
+    st.pos_dev = c.take<int32_t>((uint64_t) s.n_head);
+    // the pinned staging the uploads copy FROM - see the note on `host_step` in the header
+    if (cudaHostAlloc((void**) &st.host_step, strata::kernels::qsa_step_bytes() + sizeof(int32_t),
+                      cudaHostAllocMapped | cudaHostAllocPortable) != cudaSuccess ||
+        cudaHostAlloc((void**) &st.host_pos, (size_t) s.n_head * 4, cudaHostAllocMapped | cudaHostAllocPortable) !=
+            cudaSuccess) {
+        return 0;   // the caller sees a zero byte count; a half-built state is worse than none
+    }
+    st.host_step[strata::kernels::kStepCount] = 0;
+    // KV streaming: the authoritative K/V of every cell, pinned and device-mapped, in the identity layout
+    st.host = strata::kernels::KvHostPools{};
+    if (p.mode != 0) {
+        const uint64_t hrows = (uint64_t) pages * s.n_head_kv * s.page_size;
+        const uint64_t bytes = (uint64_t) pages * strata::kernels::kv_block_bytes(s, g_kv_int8) + 4 * 256;
+        uint8_t* h = nullptr;
+        uint8_t* d = nullptr;
+        if (cudaHostAlloc((void**) &h, bytes, cudaHostAllocMapped | cudaHostAllocPortable) != cudaSuccess ||
+            cudaHostGetDevicePointer((void**) &d, h, 0) != cudaSuccess) {
+            if (p.mode == 1) std::fprintf(stderr, "strata: KV streaming: cannot pin %.2f GiB of RAM for a layer's KV copy "
+                                 "(lower the context, or run without --kv-resident)\n", (double) bytes / 1073741824.0);
+            return 0;
+        }
+        g_kv_host_bytes += bytes;
+        Cursor hc{d};
+        if (g_kv_int8) {
+            st.host.k_q = hc.take<int8_t>(hrows * s.head_dim);
+            st.host.v_q = hc.take<int8_t>(hrows * s.head_dim);
+            st.host.k_scale = hc.take<uint16_t>(hrows * (s.head_dim / strata::kernels::KV_Q8_GROUP));
+            st.host.v_scale = hc.take<uint16_t>(hrows * (s.head_dim / strata::kernels::KV_Q8_GROUP));
+        } else {
+            st.host.k_pool = hc.take<uint16_t>(hrows * s.head_dim);
+            st.host.v_pool = hc.take<uint16_t>(hrows * s.head_dim);
+        }
+    }
+    // THE ROPE TABLE IS BUILT ON THE HOST IN FLOAT64 and uploaded once, because the reference computes its
+    // frequencies in float64 and reproducing that on device means double-precision `pow`/`cos` that need not
+    // agree with the host's libm.  A table shorter than the sequence would have the rotation read past it.
+    if (share_rope == nullptr) {
+        std::vector<float> hc((size_t) max_cells * (s.n_rot / 2)), hs((size_t) max_cells * (s.n_rot / 2));
+        strata::kernels::build_rope_table((int) s.n_rot, strata::kernels::qsa_freq_base(), (int) max_cells,
+                                          hc.data(), hs.data());
+        cudaMemcpy(st.cos_tab, hc.data(), hc.size() * 4, cudaMemcpyHostToDevice);
+        cudaMemcpy(st.sin_tab, hs.data(), hs.size() * 4, cudaMemcpyHostToDevice);
+    }
+    // the page table starts as the IDENTITY, which is the simplest legal mapping and what a caller that does
+    // not page at all wants; a streamed state starts with nothing resident, a ring at `block % n_slots`.
+    if (p.mode == 0) {
+        std::vector<int32_t> tab((size_t) pages);
+        for (int64_t i = 0; i < pages; ++i) tab[(size_t) i] = (int32_t) i;
+        cudaMemcpy(st.page_table, tab.data(), tab.size() * 4, cudaMemcpyHostToDevice);
+    } else if (p.mode == 1) {
+        strata::kernels::kv_stream_reset(st.map, nullptr);
+    } else {
+        strata::kernels::kv_ring_table(st.page_table, pages, p.slots, nullptr);
+    }
+    cudaDeviceSynchronize();
+    return c.used;
+}
+
+void qsa_state_zero(const QsaState& st, const ModelGeometry& g, void* stream) {
+    const QsaShapes s = qsa_shapes(g);
+    cudaStream_t cs = (cudaStream_t) stream;
+    const size_t rows = (size_t) st.n_slots * s.n_head_kv * s.page_size;
+    if (st.kv_int8) {
+        cudaMemsetAsync(st.k_q, 0, rows * s.head_dim, cs);
+        cudaMemsetAsync(st.v_q, 0, rows * s.head_dim, cs);
+        cudaMemsetAsync(st.k_scale, 0, rows * (s.head_dim / strata::kernels::KV_Q8_GROUP) * 2, cs);
+        cudaMemsetAsync(st.v_scale, 0, rows * (s.head_dim / strata::kernels::KV_Q8_GROUP) * 2, cs);
+    } else {
+        cudaMemsetAsync(st.k_pool, 0, rows * s.head_dim * 2, cs);
+        cudaMemsetAsync(st.v_pool, 0, rows * s.head_dim * 2, cs);
+    }
+    // A streamed state starts over with nothing resident. Its host copy is not cleared (GBs over PCIe per new
+    // conversation): no reader names a cell before this sequence has written it, and a block copied in whole
+    // carries the unwritten cells past the end, which nothing reads.
+    if (st.kv_mode == 1) strata::kernels::kv_stream_reset(st.map, stream);
+    cudaMemsetAsync(st.idx_tail, 0, (size_t) (s.idx_block - 1) * s.idx_dim * 4, cs);
+    cudaMemsetAsync(st.idx_dead, 0, (size_t) s.idx_dim * 4, cs);
+    cudaMemsetAsync(st.idx_pooled, 0, (size_t) st.idx_pooled_rows * s.idx_dim * 4, cs);
+}
+
+strata::kernels::QsaAttnPools qsa_attn_pools(const QsaState& st) {
+    strata::kernels::QsaAttnPools pools;
+    pools.page_table = st.page_table;
+    if (st.kv_int8) { pools.k_q = st.k_q; pools.v_q = st.v_q; pools.k_scale = st.k_scale; pools.v_scale = st.v_scale; }
+    else { pools.k_pool = st.k_pool; pools.v_pool = st.v_pool; }
+    return pools;
+}
+
+void qsa_kv_resolve(const QsaState& st, const ModelGeometry& g, const int32_t* ids, const int32_t* steps, int64_t n_q,
+                    int64_t cap, void* stream) {
+    if (st.kv_mode != 1) return;
+    strata::kernels::kv_stream_resolve(st.map, qsa_attn_pools(st), st.host, st.kv_int8, ids, steps, n_q, cap,
+                                       qsa_shapes(g), stream);
+}
+
 uint64_t qsa_buffers_bytes(const ModelGeometry& g, int64_t max_cells) {    const QsaShapes s = qsa_shapes(g);    const int64_t cap = strata::kernels::qsa_selection_width(strata::kernels::kTopkMaxCells, s);    uint64_t n = 0;    n += (uint64_t) q8k_bytes(g.n_embd);    n += (uint64_t) (g.n_embd / 32) * 34;
 // block_q8_0
 n += (uint64_t) g.n_embd * 2;    n += (uint64_t) g.n_head * 2 * g.head_dim * 4;    n += (uint64_t) g.n_head * g.head_dim * 4;    n += (uint64_t) g.n_head_kv * g.head_dim * 4 * 2;    n += (uint64_t) g.idx_key_dim * 4;    n += (uint64_t) g.idx_q_heads * g.idx_key_dim * 4;    n += (uint64_t) max_cells * 4;    n += (uint64_t) cap * 4;    n += (uint64_t) cap * g.n_head_kv * g.head_dim * 2 * 2;    n += (uint64_t) g.n_head * g.head_dim * 4;
@@ -676,7 +846,7 @@ if (!gemv_quantized(*w_attnk, p_k, f_k, b.x_q8_0, b.x_q8k, b.kcur, g.n_embd, g.n
 // EVERY ENTRY POINT FROM HERE ON IS THE CAPTURABLE ONE: the per-token counts come from `st.step` and every
 // launch is sized from a capacity in `st`/`b`, so this sequence can be captured and replayed.  The
 // host-scalar wrappers would be correct here today and silently wrong in a graph.
-if (st.kv_int8) kv_append_q8_step(st.k_q, st.v_q, st.k_scale, st.v_scale, st.page_table, st.step, b.kcur, b.vcur, s, stream);    else kv_append_step(st.k_pool, st.v_pool, st.page_table, st.step, b.kcur, b.vcur, s, stream);    {        const uint64_t nvk = (uint64_t) g.n_head_kv * g.head_dim;        const uint64_t base = (uint64_t) 2 * g.n_embd + 2 * g.hc + (uint64_t) g.n_head * g.head_dim + 2 * nvk + 8;        if (!st.kv_int8) dump_slot(dump, g, layer, (const float*) st.k_pool, base, nvk / 2, stream);        if (!st.kv_int8) dump_slot(dump, g, layer, (const float*) st.v_pool, base + nvk / 2, nvk / 2, stream);        dump_slot(dump, g, layer, b.vcur, base + nvk, nvk, stream);        dump_slot(dump, g, layer, b.kcur, base + 2 * nvk, nvk, stream);    }    {        const QsaIndexerBuffers ib{st.idx_tail, st.idx_dead, st.idx_pooled, st.idx_block_pos};        if (native_qsa_indexer_enabled()) {
+if (st.kv_int8) kv_append_q8_step(st.k_q, st.v_q, st.k_scale, st.v_scale, st.page_table, st.step, b.kcur, b.vcur, s, stream, &st.host);    else kv_append_step(st.k_pool, st.v_pool, st.page_table, st.step, b.kcur, b.vcur, s, stream, &st.host);    {        const uint64_t nvk = (uint64_t) g.n_head_kv * g.head_dim;        const uint64_t base = (uint64_t) 2 * g.n_embd + 2 * g.hc + (uint64_t) g.n_head * g.head_dim + 2 * nvk + 8;        if (!st.kv_int8) dump_slot(dump, g, layer, (const float*) st.k_pool, base, nvk / 2, stream);        if (!st.kv_int8) dump_slot(dump, g, layer, (const float*) st.v_pool, base + nvk / 2, nvk / 2, stream);        dump_slot(dump, g, layer, b.vcur, base + nvk, nvk, stream);        dump_slot(dump, g, layer, b.kcur, base + 2 * nvk, nvk, stream);    }    {        const QsaIndexerBuffers ib{st.idx_tail, st.idx_dead, st.idx_pooled, st.idx_block_pos};        if (native_qsa_indexer_enabled()) {
     try {
         native_qsa_indexer_append(b.idx_raw, st.step + kStepPos, pos_base,
             (const float*) w_ikn->data, RMS_EPS, ib, s, st.max_cells, (float) qsa_freq_base(), stream);
@@ -708,6 +878,8 @@ int64_t max_blocks = (st.max_cells / s.idx_block) + 2;
     } else {
     qsa_index_step(st.idx_pooled, b.q_idx, nullptr, s, st.step, max_blocks, b.cell_scores, stream);    topk_512_step(b.cell_scores, s, cap, st.step, b.ids, stream);
     }
+    // KV streaming: every block the selection names is made resident before anything reads it
+    qsa_kv_resolve(st, g, b.ids, st.step, 1, cap, stream);
     if (g_fast_attn && !native_flash_attn_short && dump == nullptr) {
         strata::kernels::QsaAttnPools pools;
         pools.k_pool = st.k_pool; pools.v_pool = st.v_pool; pools.k_q = st.k_q; pools.v_q = st.v_q;

src/core/mtp.cpp

@@ -157,9 +157,23 @@ bool MtpDrafter::load(const std::string& rt_dir, const ModelGeometry& g, Session
     // ---- the layer's own K/V (dense attention: no indexer state is read)
     const strata::kernels::QsaShapes s = shapes_of(g);
     const int64_t max_cells = ss.qsa_states[0].max_cells;
-    const uint64_t sb = qsa_state_bytes(g, max_cells, false);
+    // KV streaming: the drafter only reads its last `window` cells, so with streaming on its K/V is a ring of the
+    // window (plus the cells a round writes ahead of its queries) over a host copy, refilled on a resume. The host copy
+    // is pinned after the expert arena has pinned what it could: if it does not fit, the K/V stays whole in VRAM.
+    int64_t ring = (window > 0 && window < max_cells) ? window + 4 * (int64_t) max_t + 64 : 0;
+    uint64_t sb = qsa_state_bytes(g, max_cells, false, ring);
     if (cudaMalloc(&state_arena_, sb) != cudaSuccess) { err = "mtp: the K/V state does not fit"; return false; }
-    if (qsa_state_init(g, max_cells, state_arena_, st_, &ss.qsa_states[0]) == 0) { err = "mtp: state init failed"; return false; }
+    if (qsa_state_init(g, max_cells, state_arena_, st_, &ss.qsa_states[0], ring) == 0) {
+        if (st_.kv_mode == 0) { err = "mtp: state init failed"; return false; }
+        std::fprintf(stderr, "strata mtp: no pinned RAM left for the draft layer's K/V copy; keeping it in VRAM\n");
+        cudaGetLastError();
+        cudaFree(state_arena_);
+        st_ = QsaState{};
+        ring = -1;   // fully resident
+        sb = qsa_state_bytes(g, max_cells, false, ring);
+        if (cudaMalloc(&state_arena_, sb) != cudaSuccess) { err = "mtp: the K/V state does not fit"; return false; }
+        if (qsa_state_init(g, max_cells, state_arena_, st_, &ss.qsa_states[0], ring) == 0) { err = "mtp: state init failed"; return false; }
+    }
     qsa_state_zero(st_, g, nullptr);
     cudaDeviceSynchronize();
     vram_ += sb;
@@ -326,10 +340,10 @@ bool MtpDrafter::record_forward(int T, int step_row0, cudaStream_t cs, std::stri
             norm_rope(kcur_ + t * NKV * HD, f32("self_attn.k_norm.weight"), (int) NKV, (int) HD, pos + t * NH);
             if (st_.kv_int8)
                 kv_append_q8_step(st_.k_q, st_.v_q, st_.k_scale, st_.v_scale, st_.page_table, step + t * 4,
-                                  kcur_ + t * NKV * HD, vcur_ + t * NKV * HD, s, cs);
+                                  kcur_ + t * NKV * HD, vcur_ + t * NKV * HD, s, cs, &st_.host);
             else
                 kv_append_step(st_.k_pool, st_.v_pool, st_.page_table, step + t * 4, kcur_ + t * NKV * HD,
-                               vcur_ + t * NKV * HD, s, cs);
+                               vcur_ + t * NKV * HD, s, cs, &st_.host);
         }
         if (!full) return true;
         native_mmvq(GGML_Q8_0, q8("self_attn.q_proj.weight"), xq_, qfull_, (int) N, (int) (NH * 2 * HD), T, cs);
@@ -342,10 +356,7 @@ bool MtpDrafter::record_forward(int T, int step_row0, cudaStream_t cs, std::stri
             }
             norm_rope(qc, f32("self_attn.q_norm.weight"), (int) NH, (int) HD, pos + t * NH);
         }
-        QsaAttnPools pools;
-        pools.page_table = st_.page_table;
-        if (st_.kv_int8) { pools.k_q = st_.k_q; pools.v_q = st_.v_q; pools.k_scale = st_.k_scale; pools.v_scale = st_.v_scale; }
-        else { pools.k_pool = st_.k_pool; pools.v_pool = st_.v_pool; }
+        const QsaAttnPools pools = qsa_attn_pools(st_);
         if (window_ > 0) window_ids(const_cast<int32_t*>(step), T, (int) window_, ident_, cap_, cs);
         qsa_decode_attn_batch(qcur_, pools, ident_, step, cap_, s, attn_scratch_, attn_, T, cs);
         for (int t = 0; t < T; ++t)
@@ -491,6 +502,15 @@ bool MtpDrafter::capture_step(int j, std::string& err) {
     return finish_capture(cs_, ok, step_exec_[j], "step", err);
 }
 
+void MtpDrafter::kv_restore(int64_t upto) {
+    if (st_.kv_mode != 2 || upto <= 0) return;
+    // the ring's blocks below `upto`, from the host copy: a checkpoint resume may have left later cells in them
+    const strata::kernels::QsaShapes s = shapes_of(*g_);
+    const int64_t b1 = (upto + s.page_size - 1) / s.page_size, b0 = std::max<int64_t>(0, b1 - st_.n_slots);
+    strata::kernels::kv_ring_restore(qsa_attn_pools(st_), st_.host, st_.kv_int8, b0, b1, st_.n_slots, s, cs_);
+    cudaStreamSynchronize(cs_);
+}
+
 bool MtpDrafter::prefill(const float* R_rows, const int32_t* next_tokens, int64_t n, int64_t cell0, std::string& err) {
     const Clock::time_point t0 = Clock::now();
     const int64_t HCN = g_->hc * g_->n_embd;
@@ -551,7 +571,7 @@ bool MtpDrafter::draft(int T, const int32_t* tokens, int64_t p, int a, int32_t*
     if (probs) probs[0] = pj;
     int n = 1;
     // the chain continues while the last draft is likely enough to be verified
-    for (int j = 1; j < max_t_ - 1 && pj >= min_p; ++j) {
+    for (int j = 1; j < std::min(max_t_ - 1, max_drafts_) && pj >= min_p; ++j) {
         if (!capture_step(j, err)) return false;
         put(max_t_ + j - 1, p + a + j);
         std::atomic_thread_fence(std::memory_order_seq_cst);

src/core/verify.cpp

@@ -111,6 +111,8 @@ bool Verifier::init(const WeightTable& wt, const ModelGeometry& g, SessionState&
     hits_ = hits;
     head_ = head;
     max_t_ = max_t;
+    sampling_.greedy = true;      // a fresh verifier samples greedily until set_sampling says otherwise
+    sampling_.temperature = 0.0f;
     if (max_t < 2 || max_t > strata::kernels::kVerifyMaxT || max_t > strata::kernels::cpu::MAXT) {
         err = "verify: the window must hold 2.." + std::to_string(strata::kernels::kVerifyMaxT) + " tokens";
         return false;
@@ -424,10 +426,10 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                     const int32_t* step_t = step_ + t * kStepCount;
                     if (st.kv_int8)
                         kv_append_q8_step(st.k_q, st.v_q, st.k_scale, st.v_scale, st.page_table, step_t,
-                                          kcur_ + t * NKV * HD, vcur_ + t * NKV * HD, s, cs);
+                                          kcur_ + t * NKV * HD, vcur_ + t * NKV * HD, s, cs, &st.host);
                     else
                         kv_append_step(st.k_pool, st.v_pool, st.page_table, step_t, kcur_ + t * NKV * HD,
-                                       vcur_ + t * NKV * HD, s, cs);
+                                       vcur_ + t * NKV * HD, s, cs, &st.host);
                 }
                 const QsaIndexerBuffers ib{st.idx_tail, st.idx_dead, st.idx_pooled, st.idx_block_pos};
                 for (int t = tb; t < te; ++t)
@@ -454,10 +456,9 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                                  s, scores_ + (size_t) tb * max_blocks_, cs);
                 qsa_block_topk(scores_ + (size_t) tb * max_blocks_, step_ + tb * kStepCount, n, max_blocks_, cap_, s,
                                sel_ + (size_t) tb * cap_, cs);
-                QsaAttnPools pools;
-                pools.page_table = st.page_table;
-                if (st.kv_int8) { pools.k_q = st.k_q; pools.v_q = st.v_q; pools.k_scale = st.k_scale; pools.v_scale = st.v_scale; }
-                else { pools.k_pool = st.k_pool; pools.v_pool = st.v_pool; }
+                // KV streaming: the n selections' blocks resident (device-side, inside the graph)
+                qsa_kv_resolve(st, *g_, sel_ + (size_t) tb * cap_, step_ + tb * kStepCount, n, cap_, cs);
+                const QsaAttnPools pools = qsa_attn_pools(st);
                 qsa_decode_attn_batch(qcur_ + tb * NH * HD, pools, sel_ + (size_t) tb * cap_, step_ + tb * kStepCount, cap_,
                                       s, attn_scratch_ + (size_t) tb * attn_scratch_floats_, attn_ + tb * NH * HD, n, cs);
                 for (int t = tb; t < te; ++t) {
@@ -595,6 +596,9 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 return false;
             }
         }
+        // Greedy, the default, is recorded here as before (no extra launch or sync per window). A request that
+        // samples or penalizes is sampled again host-side after the replay (run()) with its own parameters and a
+        // fresh draw counter: a captured sampler would bake them in and replay the same draws forever.
         SamplerParams sp;
         sp.greedy = true;
         sp.temperature = 0.0f;
@@ -794,6 +798,20 @@ bool Verifier::run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool,
     const cudaError_t se = cudaStreamSynchronize(cs_);
     if (se != cudaSuccess) { err = std::string("verify: ") + cudaGetErrorString(se); return false; }
     cudaStreamSynchronize(copy_);   // no host function of this window may raise flag B in the next one
+    // ---- a sampled or penalized request: the head's sampling again, host-side so its parameters are this call's
+    // own (a captured kernel would replay the same draws forever).  Row t's draw is Philox(seed, pos0 + t): tied to
+    // the POSITION it samples, not to how the text was cut into windows, so a seed replays the same text whatever
+    // the drafts were. Exact: a rejected row's draw is discarded, and no kept decision depends on a reused draw.
+    const bool sampled = !sampling_.greedy && sampling_.temperature > 0.0f;
+    if (sampled || hist_d_ != nullptr) {
+        SamplerParams sp = sampling_;
+        sp.counter = (uint64_t) pos0;
+        sample_tokens(head_logits_, T, (int) n_vocab_, hist_d_, hist_len_, sp, m_out_, cs_);
+        if (cudaStreamSynchronize(cs_) != cudaSuccess) {   // m_out_ is the mapped h_out_: synced, it is readable
+            err = "verify: the head sampling failed";
+            return false;
+        }
+    }
     for (int t = 0; t < T; ++t) out[t] = ((volatile int32_t*) h_out_)[t];
     VDBG("window done\n");
     ++windows;

src/kernels/cpu/pool.cpp

@@ -49,11 +49,12 @@ std::vector<int> physical_cores(bool skip_first) {
 #else
     cpu_set_t set;
     CPU_ZERO(&set);
-    if (sched_getaffinity(0, sizeof set, &set) == 0)
+    if (sched_getaffinity(0, sizeof set, &set) == 0) {
         for (int i = 0; i < CPU_SETSIZE; ++i)
             if (CPU_ISSET(i, &set)) cores.push_back(i);
-    else
+    } else {
         for (unsigned i = 0; i < std::thread::hardware_concurrency(); ++i) cores.push_back((int) i);
+    }
 #endif
     if (skip_first && !cores.empty()) cores.erase(cores.begin());
     return cores;
@@ -127,10 +128,23 @@ ExpertPool::ExpertPool(int n_workers, bool pin, bool host_works) : host_works_(h
 ExpertPool::~ExpertPool() {
     stop_.store(true, std::memory_order_release);
     // Bump the epoch so a PARKED worker notices the stop flag rather than sleeping through it.
-    epoch_.fetch_add(1, std::memory_order_release);
+    publish();
     for (auto& t : threads_) t.join();
 }
 
+void ExpertPool::publish() {
+    // Both sides are seq_cst, and that is the whole lost-wakeup argument: a worker going to sleep does
+    // `sleepers_++` and then reads `epoch_`, the host does `epoch_++` and then reads `sleepers_`.  In one total
+    // order at least one of them sees the other's write - the worker sees the new epoch and does not sleep, or
+    // the host sees the sleeper and notifies under the mutex the worker holds until it is inside `wait`.
+    // On x86 the fetch_add is a locked xadd either way, so this costs the token path nothing.
+    epoch_.fetch_add(1, std::memory_order_seq_cst);
+    if (sleepers_.load(std::memory_order_seq_cst) != 0) {
+        std::lock_guard<std::mutex> lk(sleep_mu_);
+        sleep_cv_.notify_all();
+    }
+}
+
 void ExpertPool::worker(int id) {
     uint32_t seen = 0;
     // ARRIVE at the park before the first wait, so `parked_ == n_` is true from construction.  Counting only
@@ -147,9 +161,22 @@ void ExpertPool::worker(int id) {
         // - a locked read-modify-write, five workers against one cache line - so the workers spent their wait
         // invalidating each other's caches and the very line the host writes to publish work.  The counter was
         // diagnostic and nothing branched on it.  See the note on the atomics in pool.hpp.
+        //
+        // After `kSpinBeforeSleep` with no work the worker sleeps instead (issue #4).  The clock is read once
+        // every 1024 pauses, so the spin itself is unchanged.
+        const auto parked_at = std::chrono::steady_clock::now();
+        uint32_t spins = 0;
         while (epoch_.load(std::memory_order_acquire) == seen) {
             if (stop_.load(std::memory_order_relaxed)) return;
             _mm_pause();
+            if ((++spins & 1023u) != 0) continue;
+            if (std::chrono::steady_clock::now() - parked_at < kSpinBeforeSleep) continue;
+            std::unique_lock<std::mutex> lk(sleep_mu_);
+            sleepers_.fetch_add(1, std::memory_order_seq_cst);
+            sleep_cv_.wait(lk, [&] {
+                return epoch_.load(std::memory_order_seq_cst) != seen || stop_.load(std::memory_order_relaxed);
+            });
+            sleepers_.fetch_sub(1, std::memory_order_relaxed);
         }
         if (stop_.load(std::memory_order_acquire)) return;
         seen = epoch_.load(std::memory_order_relaxed);
@@ -245,7 +272,7 @@ void ExpertPool::run_phase(int mode, int n_tasks) {
     njobs_ = n_tasks;
     head_.store(0, std::memory_order_relaxed);
     done_.store(0, std::memory_order_relaxed);
-    epoch_.fetch_add(1, std::memory_order_release);
+    publish();
     if (host_works_) drain(-1, host_scratch_);
     while (done_.load(std::memory_order_acquire) != (uint32_t) n_tasks) _mm_pause();
     while (parked_.load(std::memory_order_acquire) != (uint32_t) n_) _mm_pause();
@@ -356,7 +383,7 @@ void ExpertPool::run(ExpertJob* jobs, int n) {
     mode_ = 0;
     head_.store(0, std::memory_order_relaxed);
     done_.store(0, std::memory_order_relaxed);
-    epoch_.fetch_add(1, std::memory_order_release);   // release: jobs_/njobs_ are visible before the bump
+    publish();   // a release (seq_cst): jobs_/njobs_ are visible before the bump
 
     // ---- **THE HOST DRAINS TOO (R2.2), INSTEAD OF SPINNING ON `done_`.**
     //

src/kernels/cuda/iq_kernels.cu

@@ -268,6 +268,28 @@ __device__ __forceinline__ float vec_dot_iq4_nl_q8_1(const void* __restrict__ vb
     return d * sumi;
 }
 
+// IQ4_XS: 256 values as 8 sub-blocks of 32 (6-bit scale each); one call covers one sub-block (iqs = 4 * sub-block),
+// and `bq8_1` is the super-block's first q8_1 block, so the call's activation is bq8_1[iqs / 4].  The GSQ-RCO IQ3_S
+// file keeps one layer's routed gate/up experts in this format.
+__device__ __forceinline__ float vec_dot_iq4_xs_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                     const int& kbx, const int& iqs) {
+    const block_iq4_xs* bq4 = (const block_iq4_xs*) vbq + kbx;
+    int sumi = 0;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int aux_q4 = get_int_b4(bq4->qs, iqs + j);
+        const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl);
+        const int u0 = get_int_b4(bq8_1[iqs / 4].qs, j + 0);
+        const int u1 = get_int_b4(bq8_1[iqs / 4].qs, j + 4);
+        sumi = ggml_cuda_dp4a(v.x, u0, sumi);
+        sumi = ggml_cuda_dp4a(v.y, u1, sumi);
+    }
+    const int ls = ((bq4->scales_l[iqs / 8] >> (iqs & 0x04)) & 0x0F) | (((bq4->scales_h >> (iqs / 2)) & 0x03) << 4);
+    sumi *= ls - 32;
+    const float d = __half2float(bq4->d) * __low2float(bq8_1[iqs / 4].ds);
+    return d * sumi;
+}
+
 // ---------------------------------------------------------------- the formats
 // qk = values per block, ipb = dot calls per block (qi / vdr), step = the iqs stride between calls.
 template<int TY> struct Fmt;
@@ -281,6 +303,8 @@ template<> struct Fmt<20> { static constexpr int qk = 32, ipb = 2, step = 2;
     __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq4_nl_q8_1(v, y, kbx, iqs); } };
 template<> struct Fmt<21> { static constexpr int qk = 256, ipb = 8, step = 2;
     __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq3_s_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<23> { static constexpr int qk = 256, ipb = 8, step = 4;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq4_xs_q8_1(v, y, kbx, iqs); } };
 template<> struct Fmt<22> { static constexpr int qk = 256, ipb = 8, step = 2;
     __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq2_s_q8_1(v, y, kbx, iqs); } };
 template<> struct Fmt<29> { static constexpr int qk = 256, ipb = 8, step = 1;
@@ -618,6 +642,7 @@ void iq_mmvq(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n
         case 20: mmvq_kernel<20><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
         case 21: mmvq_kernel<21><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
         case 22: mmvq_kernel<22><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
+        case 23: mmvq_kernel<23><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
         case 29: mmvq_kernel<29><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
         case 42: mmvq_kernel<42><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
         default: std::fprintf(stderr, "iq_mmvq: type %d is not supported\n", t); std::exit(1);
@@ -700,6 +725,7 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
         case 18: native_gu_kernel<18><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 21: native_gu_kernel<21><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 22: native_gu_kernel<22><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
+        case 23: native_gu_kernel<23><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 29: native_gu_kernel<29><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 42: native_gu_kernel<42><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         default: std::fprintf(stderr, "native_expert_grouped: gate/up type %d\n", L.gu_type); std::exit(1);
@@ -711,6 +737,7 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
     const dim3 gd((unsigned) ((L.n_embd + 7) / 8), (unsigned) cap_groups);
     switch (L.d_type) {
         case 20: native_down_kernel<20><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
+        case 23: native_down_kernel<23><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
         case 42: native_down_kernel<42><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
         default: std::fprintf(stderr, "native_expert_grouped: down type %d\n", L.d_type); std::exit(1);
     }

src/kernels/cuda/kv_q8.cu

@@ -31,7 +31,7 @@ __global__ void kv_append_q8_kernel(int8_t* __restrict__ k_q, int8_t* __restrict
                                     uint16_t* __restrict__ k_scale, uint16_t* __restrict__ v_scale,
                                     const int32_t* __restrict__ table, const int32_t* __restrict__ step,
                                     const float* __restrict__ kcur, const float* __restrict__ vcur, int kv_heads,
-                                    int head_dim, int page_size) {
+                                    int head_dim, int page_size, KvHostPools host) {
     const long long pos = (long long) __ldg(step + kStepPos);
     const int h = blockIdx.x, g = blockIdx.y, t = threadIdx.x;
     const bool is_v = blockIdx.z == 1;
@@ -51,10 +51,18 @@ __global__ void kv_append_q8_kernel(int8_t* __restrict__ k_q, int8_t* __restrict
         q = __float2int_rn(x / sf);
         q = q < -127 ? -127 : (q > 127 ? 127 : q);
     }
+    // KV streaming: the host copy (identity layout) always, the VRAM page only if the block is resident
     const long long page = (long long) table[pos / page_size];
-    const long long row = (page * kv_heads + h) * page_size + (pos % page_size);
-    (is_v ? v_q : k_q)[row * head_dim + g * KV_Q8_GROUP + t] = (int8_t) q;
-    if (t == 0) (is_v ? v_scale : k_scale)[row * groups + g] = sbits;
+    if (page >= 0) {
+        const long long row = (page * kv_heads + h) * page_size + (pos % page_size);
+        (is_v ? v_q : k_q)[row * head_dim + g * KV_Q8_GROUP + t] = (int8_t) q;
+        if (t == 0) (is_v ? v_scale : k_scale)[row * groups + g] = sbits;
+    }
+    if (host.k_q != nullptr) {
+        const long long row = ((pos / page_size) * kv_heads + h) * page_size + (pos % page_size);
+        (is_v ? host.v_q : host.k_q)[row * head_dim + g * KV_Q8_GROUP + t] = (int8_t) q;
+        if (t == 0) (is_v ? host.v_scale : host.k_scale)[row * groups + g] = sbits;
+    }
 }
 
 // One thread = 4 consecutive values of one cell and head (as the FP16 gather does with uint2).
@@ -93,12 +101,13 @@ __global__ void kv_gather_q8_kernel(const int8_t* __restrict__ k_q, const int8_t
 }  // namespace
 
 void kv_append_q8_step(int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale, const int32_t* page_table,
-                       const int32_t* step, const float* kcur, const float* vcur, const QsaShapes& s, void* stream) {
+                       const int32_t* step, const float* kcur, const float* vcur, const QsaShapes& s, void* stream,
+                       const KvHostPools* host) {
     validate(s, "kv_append_q8");
     const dim3 grid((unsigned) s.n_head_kv, (unsigned) (s.head_dim / KV_Q8_GROUP), 2);
     kv_append_q8_kernel<<<grid, KV_Q8_GROUP, 0, (cudaStream_t) stream>>>(
         k_q, v_q, k_scale, v_scale, page_table, step, kcur, vcur, (int) s.n_head_kv, (int) s.head_dim,
-        (int) s.page_size);
+        (int) s.page_size, host ? *host : KvHostPools{});
     check("kv_append_q8 launch");
 }
 

src/kernels/cuda/kv_stream.cu

@@ -0,0 +1,251 @@
+// src/kernels/cuda/kv_stream.cu - see include/strata/kernels/kv_stream.hpp.
+#include "strata/kernels/kv_stream.hpp"
+#include "strata/kernels/kv_q8.hpp"
+
+#include <cuda_runtime.h>
+
+#include <algorithm>
+#include <cstdio>
+#include <cstdlib>
+
+namespace strata::kernels {
+namespace {
+
+void check(const char* what) {
+    const cudaError_t e = cudaGetLastError();
+    if (e != cudaSuccess) {
+        std::fprintf(stderr, "kv_stream: %s: %s\n", what, cudaGetErrorString(e));
+        std::exit(1);
+    }
+}
+
+constexpr int RT = 1024;   // the resolve block
+
+// The per-block byte runs of the (up to four) pool arrays: block b of array i is bytes [b * len, (b + 1) * len).
+struct Runs {
+    const uint8_t* src[4];
+    uint8_t* dst[4];
+    int len[4];
+    int n;
+};
+
+Runs runs_of(const QsaAttnPools& slots, const KvHostPools& host, bool int8, const QsaShapes& s) {
+    const int rows = (int) (s.n_head_kv * s.page_size);
+    Runs r{};
+    if (int8) {
+        const int codes = rows * (int) s.head_dim, scales = rows * (int) (s.head_dim / KV_Q8_GROUP) * 2;
+        r.src[0] = (const uint8_t*) host.k_q;     r.dst[0] = (uint8_t*) slots.k_q;     r.len[0] = codes;
+        r.src[1] = (const uint8_t*) host.v_q;     r.dst[1] = (uint8_t*) slots.v_q;     r.len[1] = codes;
+        r.src[2] = (const uint8_t*) host.k_scale; r.dst[2] = (uint8_t*) slots.k_scale; r.len[2] = scales;
+        r.src[3] = (const uint8_t*) host.v_scale; r.dst[3] = (uint8_t*) slots.v_scale; r.len[3] = scales;
+        r.n = 4;
+    } else {
+        const int bytes = rows * (int) s.head_dim * 2;
+        r.src[0] = (const uint8_t*) host.k_pool; r.dst[0] = (uint8_t*) slots.k_pool; r.len[0] = bytes;
+        r.src[1] = (const uint8_t*) host.v_pool; r.dst[1] = (uint8_t*) slots.v_pool; r.len[1] = bytes;
+        r.n = 2;
+    }
+    return r;
+}
+
+// Block-wide exclusive prefix sum of one int per thread (RT threads); `total` is the sum over the block.
+__device__ int block_scan(int v, int* warp_sums, int& total) {
+    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
+    int x = v;
+    for (int o = 1; o < 32; o <<= 1) {
+        const int y = __shfl_up_sync(0xffffffffu, x, o);
+        if (lane >= o) x += y;
+    }
+    if (lane == 31) warp_sums[w] = x;
+    __syncthreads();
+    if (w == 0) {
+        int t = warp_sums[lane];
+        for (int o = 1; o < 32; o <<= 1) {
+            const int y = __shfl_up_sync(0xffffffffu, t, o);
+            if (lane >= o) t += y;
+        }
+        warp_sums[lane] = t;
+    }
+    __syncthreads();
+    total = warp_sums[31];
+    const int excl = x - v + (w > 0 ? warp_sums[w - 1] : 0);
+    __syncthreads();
+    return excl;
+}
+
+__global__ void __launch_bounds__(RT) resolve_kernel(KvStreamMap m, const int32_t* __restrict__ ids,
+                                                     const int32_t* __restrict__ steps, int n_q, int cap, int page_size) {
+    __shared__ int s_nmiss, s_lookups, s_cut;
+    __shared__ int warp_sums[32];
+    const int epoch = m.ctl[0] + 1;
+    if (threadIdx.x == 0) { s_nmiss = 0; s_lookups = 0; }
+    __syncthreads();
+    // 1. hits take this epoch and their reference bit; a missing block is claimed exactly once (-1 -> -2)
+    int lookups = 0;
+    for (int q = 0; q < n_q; ++q) {
+        const int width = steps[q * kStepCount + kStepWidth];
+        const int32_t* qi = ids + (long long) q * cap;
+        for (int i = threadIdx.x; i < width; i += RT) {
+            const int b = qi[i] / page_size;
+            if (i > 0 && qi[i - 1] / page_size == b) continue;   // ids are ascending: one lookup per block
+            ++lookups;
+            const int sl = m.page_table[b];
+            if (sl >= 0) {
+                m.slot_stamp[sl] = epoch;
+                m.slot_ref[sl] = 1;
+            } else if (sl == -1 && atomicCAS(&m.page_table[b], -1, -2) == -1) {
+                m.miss_block[atomicAdd(&s_nmiss, 1)] = b;
+            }
+        }
+    }
+    atomicAdd(&s_lookups, lookups);
+    __syncthreads();
+    // 2. one victim per miss: a clock sweep from the hand. A slot this call uses (stamp == epoch) is never taken;
+    //    a referenced one loses its bit as the hand passes it and is taken on the next pass.
+    const int need = s_nmiss, n = (int) m.n_slots;
+    int hand = m.ctl[1], got = 0;
+    for (int scanned = 0; got < need && scanned < 3 * n; scanned += RT) {
+        const int j = (int) (((long long) hand + threadIdx.x) % n);
+        const bool mine = m.slot_stamp[j] == epoch;
+        const bool cand = !mine && (m.slot_block[j] < 0 || m.slot_ref[j] == 0);
+        int total = 0;
+        const int rank = block_scan(cand ? 1 : 0, warp_sums, total);
+        const int want = need - got;
+        if (threadIdx.x == 0) s_cut = RT;
+        __syncthreads();
+        if (cand && rank == want - 1) s_cut = threadIdx.x + 1;   // the hand stops just past the last slot taken
+        __syncthreads();
+        const int cut = s_cut;
+        if (cand && rank < want) {
+            m.miss_slot[got + rank] = j;
+            m.slot_stamp[j] = epoch;   // taken: a sweep that wraps around must not take it twice
+        } else if (threadIdx.x < cut && !mine) {
+            m.slot_ref[j] = 0;
+        }
+        got += total < want ? total : want;
+        hand = (int) (((long long) hand + cut) % n);
+        __syncthreads();
+    }
+    // 3. re-point the table; the copy kernel fills the slots
+    const int placed = got < need ? got : need;
+    for (int k = threadIdx.x; k < need; k += RT) {
+        const int b = m.miss_block[k];
+        if (k >= placed) { m.page_table[b] = -1; continue; }   // overflow: never happens with a legal n_slots
+        const int sl = m.miss_slot[k];
+        const int old = m.slot_block[sl];
+        if (old >= 0) m.page_table[old] = -1;
+        m.slot_block[sl] = b;
+        m.slot_stamp[sl] = epoch;
+        m.slot_ref[sl] = 1;
+        m.page_table[b] = sl;
+    }
+    if (threadIdx.x == 0) {
+        m.ctl[0] = epoch;
+        m.ctl[1] = hand;
+        m.ctl[2] = placed;
+        if (placed < need) m.ctl[3] = 1;
+        unsigned long long* c = reinterpret_cast<unsigned long long*>(m.ctl + 4);
+        c[0] += (unsigned long long) placed;
+        c[1] += (unsigned long long) s_lookups;
+        c[2] += 1ull;
+    }
+}
+
+// One block per missed block (grid-stride): copy its runs from the host copy into its slot, 16 B per thread.
+__global__ void copy_kernel(KvStreamMap m, Runs r) {
+    const int need = m.ctl[2];
+    for (int k = blockIdx.x; k < need; k += gridDim.x) {
+        const long long b = m.miss_block[k], sl = m.miss_slot[k];
+        for (int a = 0; a < r.n; ++a) {
+            const uint4* src = reinterpret_cast<const uint4*>(r.src[a] + b * r.len[a]);
+            uint4* dst = reinterpret_cast<uint4*>(r.dst[a] + sl * r.len[a]);
+            for (int i = threadIdx.x; i < r.len[a] / 16; i += blockDim.x) dst[i] = src[i];
+        }
+    }
+}
+
+__global__ void reset_kernel(KvStreamMap m) {
+    const long long i0 = (long long) blockIdx.x * blockDim.x + threadIdx.x, st = (long long) gridDim.x * blockDim.x;
+    for (long long i = i0; i < m.n_blocks; i += st) m.page_table[i] = -1;
+    for (long long i = i0; i < m.n_slots; i += st) {
+        m.slot_block[i] = -1;
+        m.slot_stamp[i] = -1;
+        m.slot_ref[i] = 0;
+    }
+    if (i0 < kKvCtlInts) m.ctl[i0] = 0;
+}
+
+__global__ void ring_kernel(int32_t* table, long long n_blocks, long long n_slots) {
+    for (long long i = (long long) blockIdx.x * blockDim.x + threadIdx.x; i < n_blocks;
+         i += (long long) gridDim.x * blockDim.x)
+        table[i] = (int32_t) (i % n_slots);
+}
+
+}  // namespace
+
+uint64_t kv_block_bytes(const QsaShapes& s, bool int8) {
+    const uint64_t rows = (uint64_t) (s.n_head_kv * s.page_size);
+    return int8 ? rows * (uint64_t) s.head_dim * 2 + rows * (uint64_t) (s.head_dim / KV_Q8_GROUP) * 2 * 2
+                : rows * (uint64_t) s.head_dim * 2 * 2;
+}
+
+void kv_stream_reset(const KvStreamMap& m, void* stream) {
+    reset_kernel<<<128, 256, 0, (cudaStream_t) stream>>>(m);
+    check("reset");
+}
+
+void kv_stream_resolve(const KvStreamMap& m, const QsaAttnPools& slots, const KvHostPools& host, bool int8,
+                       const int32_t* ids, const int32_t* steps, int64_t n_q, int64_t cap, const QsaShapes& s,
+                       void* stream) {
+    if (n_q <= 0) return;
+    if (s.n_head_kv * s.page_size * (s.head_dim / KV_Q8_GROUP) * 2 % 16 != 0) {
+        std::fprintf(stderr, "kv_stream: a block's scale run must be a multiple of 16 bytes\n");
+        std::exit(1);
+    }
+    resolve_kernel<<<1, RT, 0, (cudaStream_t) stream>>>(m, ids, steps, (int) n_q, (int) cap, (int) s.page_size);
+    check("resolve");
+    copy_kernel<<<96, 128, 0, (cudaStream_t) stream>>>(m, runs_of(slots, host, int8, s));
+    check("copy");
+}
+
+void kv_ring_table(int32_t* page_table, int64_t n_blocks, int64_t n_slots, void* stream) {
+    ring_kernel<<<64, 256, 0, (cudaStream_t) stream>>>(page_table, n_blocks, n_slots);
+    check("ring table");
+}
+
+void kv_ring_restore(const QsaAttnPools& slots, const KvHostPools& host, bool int8, int64_t b0, int64_t b1,
+                     int64_t n_slots, const QsaShapes& s, void* stream) {
+    const Runs r = runs_of(slots, host, int8, s);
+    for (int64_t b = b0; b < b1;) {
+        const int64_t sl = b % n_slots, run = std::min<int64_t>(b1 - b, n_slots - sl);   // up to the ring's end
+        for (int a = 0; a < r.n; ++a)
+            if (cudaMemcpyAsync(r.dst[a] + sl * r.len[a], r.src[a] + b * r.len[a], (size_t) (run * r.len[a]),
+                                cudaMemcpyDefault, (cudaStream_t) stream) != cudaSuccess)
+                check("ring restore");
+        b += run;
+    }
+}
+
+void kv_stage_from_host(const QsaAttnPools& stage, const KvHostPools& host, bool int8, int64_t n_blocks,
+                        const QsaShapes& s, void* stream) {
+    if (n_blocks <= 0) return;
+    const Runs r = runs_of(stage, host, int8, s);
+    for (int a = 0; a < r.n; ++a)
+        if (cudaMemcpyAsync(r.dst[a], r.src[a], (size_t) (n_blocks * r.len[a]), cudaMemcpyDefault,
+                            (cudaStream_t) stream) != cudaSuccess)
+            check("stage");
+}
+
+KvStreamCounters kv_stream_counters(const KvStreamMap& m) {
+    int32_t c[kKvCtlInts] = {};
+    KvStreamCounters r;
+    if (m.ctl == nullptr || cudaMemcpy(c, m.ctl, sizeof(c), cudaMemcpyDeviceToHost) != cudaSuccess) return r;
+    const unsigned long long* u = reinterpret_cast<const unsigned long long*>(c + 4);
+    r.misses = u[0];
+    r.lookups = u[1];
+    r.calls = u[2];
+    r.overflow = c[3] != 0;
+    return r;
+}
+
+}  // namespace strata::kernels

src/kernels/cuda/qsa.cu

@@ -43,6 +43,7 @@
 #include "strata/kernels/qsa.hpp"
 
 #include "strata/kernels/f16_bits.hpp"
+#include "strata/kernels/kv_stream.hpp"
 #include "strata/kernels/rope.hpp"
 #include "strata/kernels/mrope.hpp"
 
@@ -120,16 +121,24 @@ __device__ __forceinline__ float h2f(uint16_t bits) { return f32_from_f16(bits);
 __global__ void kv_append_kernel(uint16_t* __restrict__ k_pool, uint16_t* __restrict__ v_pool,
                                  const int32_t* __restrict__ table, const int32_t* __restrict__ step,
                                  const float* __restrict__ kcur, const float* __restrict__ vcur, int kv_heads,
-                                 int head_dim, int page_size) {
+                                 int head_dim, int page_size, KvHostPools host) {
     const long long pos = (long long) __ldg(step + kStepPos);
     const int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= kv_heads * head_dim) return;
     const int h = i / head_dim, d = i - h * head_dim;
-    // `[page][kv_head][page_size][head_dim]`: one head's consecutive cells are contiguous inside a page
+    // `[page][kv_head][page_size][head_dim]`: one head's consecutive cells are contiguous inside a page.
+    // KV streaming: the host copy (identity layout) always, the VRAM page only if the block is resident.
     const long long page = (long long) table[pos / page_size];
-    const long long row = (page * kv_heads + h) * page_size + (pos % page_size);
-    k_pool[row * head_dim + d] = f16_from_f32(kcur[i]);
-    v_pool[row * head_dim + d] = f16_from_f32(vcur[i]);
+    if (page >= 0) {
+        const long long row = (page * kv_heads + h) * page_size + (pos % page_size);
+        k_pool[row * head_dim + d] = f16_from_f32(kcur[i]);
+        v_pool[row * head_dim + d] = f16_from_f32(vcur[i]);
+    }
+    if (host.k_pool != nullptr) {
+        const long long row = ((pos / page_size) * kv_heads + h) * page_size + (pos % page_size);
+        host.k_pool[row * head_dim + d] = f16_from_f32(kcur[i]);
+        host.v_pool[row * head_dim + d] = f16_from_f32(vcur[i]);
+    }
 }
 
 // ================= 2. indexer_key_append =================
@@ -575,13 +584,14 @@ const int32_t* step_upload_width(int64_t width, const QsaShapes& s) {
 // buffer and forward this token's counts as the capacities.
 
 void kv_append_step(uint16_t* k_pool, uint16_t* v_pool, const int32_t* page_table, const int32_t* step,
-                    const float* kcur, const float* vcur, const QsaShapes& s, void* stream) {
+                    const float* kcur, const float* vcur, const QsaShapes& s, void* stream, const KvHostPools* host) {
     validate(s, "kv_append");
     if (step == nullptr) fail("kv_append: step is null");
     // The grid is the head x dim count, which is a CONSTANT - `kv_append` writes one cell.
     const long long n = s.n_head_kv * s.head_dim;
     kv_append_kernel<<<grid_for(n, THREADS), THREADS, 0, (cudaStream_t) stream>>>(
-        k_pool, v_pool, page_table, step, kcur, vcur, (int) s.n_head_kv, (int) s.head_dim, (int) s.page_size);
+        k_pool, v_pool, page_table, step, kcur, vcur, (int) s.n_head_kv, (int) s.head_dim, (int) s.page_size,
+        host ? *host : KvHostPools{});
     check_launch("kv_append");
     if (stream == nullptr) check_sync("kv_append");
 }

src/kernels/cuda/sampler.cu

@@ -1,6 +1,6 @@
 // src/kernels/cuda/sampler.cu - P2.S2: the sampler chain, in the order docs/sampling.md settles.
 //
-//     penalties  ->  top_k  ->  top_p  ->  temperature  ->  pick
+//     penalties -> top_k -> min_p -> top_p -> temperature -> penalties -> pick
 //
 // THE ORDER IS THE WHOLE CONTENT OF THIS FILE.  `docs/sampling.md` transcribes it from llama.cpp's own chain
 // (`common/sampling.cpp` L357/360/375/381/399) and the two facts that are easy to get backwards are that
@@ -9,10 +9,9 @@
 // valid token, so only a comparison at the distribution level can tell them apart; the parity test does that
 // explicitly by running the wrong order and requiring it to differ.
 //
-// Everything happens on the logits of ONE token in one thread.  The vocab is 248,320, which is far too large
-// to sort per token on a naive path, so `top_k` uses `k` passes of a maximum scan - 20 x 248,320 = 5.0M
-// comparisons per token.  Phase 3's note ("sort-free sampler: top-k = 20 makes this easy") is the optimisation;
-// this is the correctness version.
+// Both kernels put ONE BLOCK per token over the vocabulary: `sampler_greedy_kernel` is the plain argmax,
+// `sampler_kernel` runs the sampled chain as `top_k` block-argmax rounds followed by the top_p / temperature /
+// draw chain (its header says why the selection must be parallel and why the tie rule keeps the semantics).
 #include "strata/kernels/sampler.hpp"
 
 #include <cuda_runtime.h>
@@ -102,12 +101,32 @@ __global__ void sampler_greedy_kernel(const float* __restrict__ logits, int n_vo
         hrow += history_len - hlen;          // the window is the TAIL
     }
 
+    // PENALTY MEMBERSHIP AS A BITMAP.  The history touches at most `hlen` tokens of a quarter-million
+    // vocabulary, but the naive `history_count` per candidate per argmax round costs O(k x n_vocab x hlen)
+    // integer compares (~318 M per token at k=20, hlen=64 - measured 45 -> 31 tok/s on a real workload).
+    // A shared bitmap gives an O(1) membership test, and only the (at most hlen) hits pay the count scan;
+    // the counts - and therefore every sampled value - are exactly what the per-candidate scan produced.
+    extern __shared__ unsigned int penal_bits[];
+    const int bits_words = (int) ((n_vocab + 31) / 32);
+    const bool use_bits = hrow != nullptr && bits_words > 0;
+    if (use_bits) {
+        for (int w = threadIdx.x; w < bits_words; w += blockDim.x) penal_bits[w] = 0u;
+        __syncthreads();
+        for (int i = threadIdx.x; i < hlen; i += blockDim.x)
+            if (hrow[i] >= 0) atomicOr(&penal_bits[hrow[i] >> 5], 1u << (hrow[i] & 31));
+        __syncthreads();
+    }
+    auto hit_count = [&](int v) -> int {
+        if (!use_bits || !(penal_bits[v >> 5] & (1u << (v & 31)))) return 0;
+        return history_count(hrow, hlen, v);
+    };
+
     // `n_vocab` is the "no candidate" index: it loses every comparison to a real one, so a thread with no
     // elements contributes nothing rather than contributing a bogus zero.
     float bv = __int_as_float(0xff800000);   // -inf
     int best = n_vocab;
     for (int v = threadIdx.x; v < n_vocab; v += blockDim.x) {
-        const float s = apply_penalties(l[v], hrow ? history_count(hrow, hlen, v) : 0, p);
+        const float s = apply_penalties(l[v], hit_count(v), p);
         if (s > bv) { bv = s; best = v; }
     }
     for (int off = 16; off > 0; off >>= 1) {
@@ -134,47 +153,32 @@ __global__ void sampler_greedy_kernel(const float* __restrict__ logits, int n_vo
     }
 }
 
+/// **THE SAMPLED PATH, ONE BLOCK PER TOKEN.**  The kernel below replaced a version that ran the whole chain
+/// in ONE THREAD per token (`<<<ceil(T/64), 64>>>`, so a 4-token window fielded four threads): `top_k` alone
+/// was `k` sequential scans of the vocabulary with an inner sweep over the already-taken list - 20 x 248,320
+/// iterations of dependent work on one SM - and a verify window measured **1.6 s in the sampler**, which made
+/// every temperature-bearing request ~30x slower than a greedy one.  The selection is `k` argmax rounds, and
+/// an argmax over the vocabulary parallelises exactly like `sampler_greedy_kernel` (block over the vocab), so
+/// the rounds run back to back inside a block-per-token launch: the per-token cost falls to
+/// `k x n_vocab / 1024` plus `k` block reductions.
+///
+/// THE SEMANTICS ARE THE SERIAL ONES, EXACTLY.  Each round's argmax resolves ties to the LOWEST index (the
+/// serial scan's strict `>` keeps the first maximum it meets), so the kept sequence - both its set and its
+/// order - is unchanged; `top_p`'s cut reads that order in double arithmetic as before; temperature and the
+/// Philox draw apply after the cut.  `sampler_parity` pins all of it against the host reference.
 __global__ void sampler_kernel(const float* __restrict__ logits, int n_vocab, int n_tokens,
                                const int* __restrict__ history, int history_len, const SamplerParams p,
                                int* __restrict__ out) {
-    const int t = blockIdx.x * blockDim.x + threadIdx.x;
+    const int t = blockIdx.x;
     if (t >= n_tokens) return;
-
-    // Work on a private copy: the chain mutates the distribution and the caller's logits are shared with the
-    // next stage.  A 248,320-float allocation per thread is far too much for shared memory, so this uses a
-    // two-pass approach instead - first find the keep set, then normalise over it - and never materialises a
-    // second full-size buffer.
     const float* l = logits + (size_t) t * n_vocab;
 
-    // Temperature is needed by BOTH paths below, so it is computed here; the chain still APPLIES it after
+    // Temperature is needed by BOTH stages below, so it is computed here; the chain still APPLIES it after
     // the truncation filters - the survivors are chosen on the raw logits and only then scaled.
     const float inv_t = p.temperature > 0.0f ? 1.0f / p.temperature : 0.0f;
 
-    // ---- penalties: NOT IMPLEMENTED, and the ORDER they must take is transcribed here so the next person
-    // does not have to guess it.  `llama_sampler_penalties_apply` (src/llama-sampler.cpp) runs AFTER
-    // temperature in the chain (L381 against L375), and for every token seen in the last `penalty_last_n`:
-    //
-    //     if (logit <= 0) logit *= penalty_repeat;   else logit /= penalty_repeat;
-    //     logit -= float(count) * penalty_freq + float(count > 0) * penalty_present;
-    //
-    // TWO details there are counter-intuitive.  The repeat penalty MULTIPLIES for non-positive logits and
-    // DIVIDES for positive ones - the source notes the paper only ever divided, "but that would cause tokens
-    // with negative logits to become more likely, which is obviously wrong".  And the presence penalty rides
-    // on `float(count > 0)`, a BOOLEAN cast to float, not the count - so it applies once however many times
-    // the token appeared.  Both are the kind of thing a plausible implementation gets wrong while still
-    // producing valid tokens.
-    //
-    // It needs the token history, which is not in this signature, so a caller passing a non-zero
-    // `penalty_last_n` gets a loud failure rather than a silently unpenalised sample.
-    // ---- top_k: k passes of a maximum scan, skipping the already-taken
-    // ---- top_p: cumulative mass over the survivors in descending order
-    // The two are done together in one selection pass so the survivors and their mass are consistent.
-    // GREEDY IS THE GLOBAL ARGMAX, WHATEVER THE FILTERS SAY - so it is computed directly and the filter chain
-    // is skipped entirely.  No filter can remove the maximum: top_k keeps the k largest, top_p keeps a prefix
-    // of the descending order that always contains the largest (min_keep >= 1), and temperature is monotonic so
-    // it cannot reorder anything.  That also makes greedy INDEPENDENT OF THE ORDER, which is why the ordering
-    // test cannot use it - see sampler_parity.cpp.
-    // The penalty window is the last `penalty_last_n` entries of this row's history.
+    // The penalty window is the last `penalty_last_n` entries of this row's history (disabled at this
+    // launch: `sample_tokens` refuses a non-zero `penalty_last_n` without a history buffer).
     const int* hrow = history ? history + (size_t) t * history_len : nullptr;
     int hlen = 0;
     if (hrow) {
@@ -183,109 +187,117 @@ __global__ void sampler_kernel(const float* __restrict__ logits, int n_vocab, in
         hrow += history_len - hlen;          // the window is the TAIL
     }
 
-    if (p.greedy || p.temperature <= 0.0f) {
-        // Greedy is the global argmax ONLY when no penalty is active: a penalty can lower a token enough to
-        // change the maximum, so the argmax has to be taken over the penalised logits, not the raw ones.
-        //
-        // *** TEMPERATURE MUST NOT BE APPLIED HERE.  THIS WAS A REAL BUG. ***
-        // This branch used to read `apply_penalties(l[v] * inv_t, ...)`.  `inv_t` is 0.0f whenever
-        // temperature <= 0 (see above), so at temperature 0 EVERY logit became 0.0f and the argmax returned
-        // index 0 - the sampler emitted token 0 forever, whatever the model predicted.  OpenAI clients send
-        // `temperature: 0` for greedy decoding, so this was reachable from any ordinary client.
-        // It is also pointless even when inv_t is non-zero: scaling by a positive constant is monotonic and
-        // cannot reorder the argmax, which is exactly what the note above this block already says.
-        int best = 0;
-        float bv = apply_penalties(l[0], hrow ? history_count(hrow, hlen, 0) : 0, p);
-        for (int v = 1; v < n_vocab; ++v) {
-            const float s = apply_penalties(l[v], hrow ? history_count(hrow, hlen, v) : 0, p);
-            if (s > bv) { bv = s; best = v; }
-        }
-        out[t] = best;
-        return;
+    // the membership bitmap, as in `sampler_greedy_kernel` - see the cost note there
+    extern __shared__ unsigned int penal_bits[];
+    const int bits_words = (int) ((n_vocab + 31) / 32);
+    const bool use_bits = hrow != nullptr && bits_words > 0;
+    if (use_bits) {
+        for (int w = threadIdx.x; w < bits_words; w += blockDim.x) penal_bits[w] = 0u;
+        __syncthreads();
+        for (int i = threadIdx.x; i < hlen; i += blockDim.x)
+            if (hrow[i] >= 0) atomicOr(&penal_bits[hrow[i] >> 5], 1u << (hrow[i] & 31));
+        __syncthreads();
     }
+    auto hit_count = [&](int v) -> int {
+        if (!use_bits || !(penal_bits[v >> 5] & (1u << (v & 31)))) return 0;
+        return history_count(hrow, hlen, v);
+    };
 
     const int KMAX = 64;
-    int keep_ids[KMAX];
-    float keep_logit[KMAX];
-    int k = p.top_k > 0 ? (p.top_k < KMAX ? p.top_k : KMAX) : KMAX;
-    if (p.top_k <= 0) k = 0;                      // 0 means "no top-k filter" -> handled by the mask below
-
+    int k = p.top_k > 0 ? (p.top_k < KMAX ? p.top_k : KMAX) : 0;
     if (k <= 0) {
-        std::printf("sampler: the sampled path needs top_k in 1..%d (got %d); greedy needs no filters\n", KMAX, p.top_k);
+        if (threadIdx.x == 0)
+            std::printf("sampler: the sampled path needs top_k in 1..%d (got %d); greedy needs no filters\n",
+                        KMAX, p.top_k);
         return;   // leave out[t] unwritten rather than returning an uninitialised token
     }
-    int n_keep = 0;
-    {
-        for (int i = 0; i < k; ++i) {
-            int best = -1;
-            float bv = 0.0f;
-            for (int v = 0; v < n_vocab; ++v) {
-                bool taken = false;
-                for (int j = 0; j < i; ++j) if (keep_ids[j] == v) { taken = true; break; }
-                if (taken) continue;
-                if (best < 0 || l[v] > bv) { best = v; bv = l[v]; }
+
+    // ---- top_k: k rounds of a block argmax over the not-yet-taken.  `sel_*` holds the kept ids and their
+    // raw logits in selection order: descending by value, ties to the lower index, which is the order the
+    // top_p cut below is defined over.
+    __shared__ int sel_ids[KMAX];
+    __shared__ float sel_logit[KMAX];
+    __shared__ float sv[32];
+    __shared__ int si[32];
+    for (int i = 0; i < k; ++i) {
+        // `n_vocab` is the "no candidate" index: it loses every comparison to a real one (same convention as
+        // the greedy kernel, whose tie rule this reduction shares).
+        float bv = __int_as_float(0xff800000);   // -inf
+        int best = n_vocab;
+        for (int v = threadIdx.x; v < n_vocab; v += blockDim.x) {
+            bool taken = false;
+            for (int j = 0; j < i; ++j) if (sel_ids[j] == v) { taken = true; break; }
+            if (taken) continue;
+            const float s = apply_penalties(l[v], hit_count(v), p);
+            if (s > bv) { bv = s; best = v; }
+        }
+        for (int off = 16; off > 0; off >>= 1) {
+            const float ov = __shfl_down_sync(0xFFFFFFFFu, bv, off);
+            const int oi = __shfl_down_sync(0xFFFFFFFFu, best, off);
+            if (ov > bv || (ov == bv && oi < best)) { bv = ov; best = oi; }
+        }
+        const int warp = (int) (threadIdx.x >> 5), lane = (int) (threadIdx.x & 31);
+        if (lane == 0) { sv[warp] = bv; si[warp] = best; }
+        __syncthreads();
+        if (warp == 0) {
+            const int nw = (int) ((blockDim.x + 31) >> 5);
+            float wv = lane < nw ? sv[lane] : __int_as_float(0xff800000);
+            int wi = lane < nw ? si[lane] : n_vocab;
+            for (int off = 16; off > 0; off >>= 1) {
+                const float ov = __shfl_down_sync(0xFFFFFFFFu, wv, off);
+                const int oi = __shfl_down_sync(0xFFFFFFFFu, wi, off);
+                if (ov > wv || (ov == wv && oi < wi)) { wv = ov; wi = oi; }
             }
-            keep_ids[i] = best;
-            keep_logit[i] = bv;
-        }
-        n_keep = k;
-
-        // ---- top_p over the survivors, in descending order (which the selection above already produced)
-        if (p.top_p < 1.0f) {
-            // softmax over the keep set for the cumulative mass
-            float mx = keep_logit[0];
-            for (int i = 1; i < n_keep; ++i) mx = fmaxf(mx, keep_logit[i]);
-            double sum = 0.0;
-            for (int i = 0; i < n_keep; ++i) sum += exp((double) keep_logit[i] - (double) mx);
-            double cum = 0.0;
-            int cut = n_keep;
-            for (int i = 0; i < n_keep; ++i) {
-                cum += exp((double) keep_logit[i] - (double) mx) / sum;
-                if (cum >= (double) p.top_p) { cut = i + 1; break; }
-            }
-            if (cut < p.min_keep) cut = p.min_keep < n_keep ? p.min_keep : n_keep;
-            n_keep = cut;
+            if (lane == 0) { sel_ids[i] = (wi < n_vocab) ? wi : 0; sel_logit[i] = wv; }
         }
+        __syncthreads();
     }
 
+    // ---- min_p: keep the descending prefix whose probability is at least `min_p` of the top token's.  The
+    // kept list is in selection order (descending), so the survivors are a PREFIX and the cut composes with
+    // top_p's below.  In logit space the threshold is `sel_logit[0] + logf(min_p)` - equivalent to
+    // `p >= min_p * p_max` without the overflow an exp of raw logits risks.  0 disables, and the head itself
+    // always survives (`expf(0) == 1 >= min_p` for min_p in 0..1), so the count never reaches zero.
+    int n_minp = k;
+    if (p.min_p > 0.0f) {
+        const float thresh = sel_logit[0] + logf(p.min_p);
+        for (int i = 0; i < k; ++i)
+            if (sel_logit[i] < thresh) { n_minp = i; break; }
+    }
 
-    if (p.greedy || p.temperature <= 0.0f) {
-        // DEAD CODE AS WRITTEN: the identical condition at the top of this function already returned, so this
-        // block cannot be reached.  It is kept and corrected rather than deleted because if that early return
-        // is ever narrowed, this is the path that would run - and it carried the same `* inv_t` bug (every
-        // logit 0.0f at temperature 0, so `best` would always be `keep_ids[0]`).  Temperature is monotonic and
-        // cannot reorder a maximum, so it must not appear here at all.
-        // argmax over the SURVIVORS, ties by smallest id (the same convention as the router and top_k)
-        int best = keep_ids[0];
-        float bv = keep_logit[0];
-        for (int i = 1; i < n_keep; ++i) {
-            const float v = keep_logit[i];
-            if (v > bv) { bv = v; best = keep_ids[i]; }
+    // ---- top_p over the survivors, in descending order (which the selection produced), then temperature and
+    // one Philox draw.  Every thread computes the same chain redundantly over `sel_*` - the arithmetic is the
+    // serial kernel's, instruction for instruction - so they agree on `pick` and thread 0 writes it.
+    int n_keep = n_minp;
+    float mx = sel_logit[0];
+    for (int i = 1; i < n_minp; ++i) mx = fmaxf(mx, sel_logit[i]);
+    if (p.top_p < 1.0f) {
+        double sum = 0.0;
+        for (int i = 0; i < n_minp; ++i) sum += exp((double) sel_logit[i] - (double) mx);
+        double cum = 0.0;
+        int cut = n_minp;
+        for (int i = 0; i < n_minp; ++i) {
+            cum += exp((double) sel_logit[i] - (double) mx) / sum;
+            if (cum >= (double) p.top_p) { cut = i + 1; break; }
         }
-        out[t] = best;
-        return;
+        if (cut < p.min_keep) cut = p.min_keep < n_minp ? p.min_keep : n_minp;
+        n_keep = cut;
     }
-
-    // ---- pick: temperature, then penalties, then softmax over the survivors, then one uniform draw.
-    // The ORDER is the chain's: the survivors were chosen on the RAW logits, temperature scales them, and the
-    // penalty is applied to the scaled value - which is why the `logit <= 0` branch inside the penalty sees a
-    // temperature-scaled logit and not the model's own.
-    for (int i = 0; i < n_keep; ++i) {
-        keep_logit[i] = apply_penalties(keep_logit[i] * inv_t,
-                                        hrow ? history_count(hrow, hlen, keep_ids[i]) : 0, p);
-    }
-    float mx = keep_logit[0];
-    for (int i = 1; i < n_keep; ++i) mx = fmaxf(mx, keep_logit[i]);
+    auto scaled = [&](int i) {
+        return apply_penalties(sel_logit[i] * inv_t, hit_count(sel_ids[i]), p);
+    };
+    float smx = scaled(0);
+    for (int i = 1; i < n_keep; ++i) smx = fmaxf(smx, scaled(i));
     double sum = 0.0;
-    for (int i = 0; i < n_keep; ++i) sum += exp((double) keep_logit[i] - (double) mx);
+    for (int i = 0; i < n_keep; ++i) sum += exp((double) scaled(i) - (double) smx);
     const float u = philox_uniform(p.seed, p.counter + (uint64_t) t);
     double cum = 0.0;
-    int pick = keep_ids[n_keep - 1];
+    int pick = sel_ids[n_keep - 1];
     for (int i = 0; i < n_keep; ++i) {
-        cum += exp((double) keep_logit[i] - (double) mx) / sum;
-        if ((double) u < cum) { pick = keep_ids[i]; break; }
+        cum += exp((double) scaled(i) - (double) smx) / sum;
+        if ((double) u < cum) { pick = sel_ids[i]; break; }
     }
-    out[t] = pick;
+    if (threadIdx.x == 0) out[t] = pick;
 }
 
 }  // namespace
@@ -298,16 +310,19 @@ void sample_tokens(const float* logits, int n_tokens, int n_vocab, const int* hi
                      p.penalty_last_n, (const void*) history, history_len);
         std::exit(1);
     }
-    const int threads = 64;
-    const unsigned grid = (unsigned) ((n_tokens + threads - 1) / threads);
+    const unsigned shmem = (history != nullptr && history_len > 0 && p.penalty_last_n > 0)
+                               ? (unsigned) ((n_vocab + 31) / 32) * sizeof(unsigned)   // the penalty bitmap
+                               : 0;
     if (p.greedy || p.temperature <= 0.0f) {
         // One block per token, 1,024 threads over the vocabulary.  See `sampler_greedy_kernel`.
         const int gthreads = 1024;
-        sampler_greedy_kernel<<<(unsigned) n_tokens, gthreads, 0, (cudaStream_t) stream>>>(
+        sampler_greedy_kernel<<<(unsigned) n_tokens, gthreads, shmem, (cudaStream_t) stream>>>(
             logits, n_vocab, history, history_len, p, p.penalty_last_n, p.penalty_last_n, out);
     } else {
-        sampler_kernel<<<grid, threads, 0, (cudaStream_t) stream>>>(logits, n_vocab, n_tokens, history,
-                                                                   history_len, p, out);
+        // The same block-per-token shape: the selection's k argmax rounds reduce inside the block.  See
+        // `sampler_kernel`'s header for what the old one-thread-per-token launch cost.
+        sampler_kernel<<<(unsigned) n_tokens, 1024, shmem, (cudaStream_t) stream>>>(
+            logits, n_vocab, n_tokens, history, history_len, p, out);
     }
     const cudaError_t e = cudaGetLastError();
     if (e != cudaSuccess) {

src/kernels/cuda/verify_kernels.cu

@@ -244,12 +244,14 @@ __global__ void ident_hits_kernel(const int32_t* __restrict__ ids, int n, int32_
     if (i == 0) *count = n;
 }
 
-__global__ void gather_rows_kernel(const uint4* __restrict__ src, long long row16, const int32_t* __restrict__ ids,
-                                   long long n, uint4* __restrict__ dst) {
-    const long long total = n * row16;
+// E = the widest element the row size divides into (16, 4 or 1 bytes): a Q6_K head row of 2560 values is 2100 bytes
+template<typename E>
+__global__ void gather_rows_kernel(const E* __restrict__ src, long long row_e, const int32_t* __restrict__ ids,
+                                   long long n, E* __restrict__ dst) {
+    const long long total = n * row_e;
     for (long long i = (long long) blockIdx.x * blockDim.x + threadIdx.x; i < total; i += (long long) gridDim.x * blockDim.x) {
-        const long long r = i / row16, o = i - r * row16;
-        dst[i] = src[(long long) ids[r] * row16 + o];
+        const long long r = i / row_e, o = i - r * row_e;
+        dst[i] = src[(long long) ids[r] * row_e + o];
     }
 }
 
@@ -334,8 +336,13 @@ void mtp_select(const float* R_src, int64_t R_stride, const int32_t* ids, const
 }
 
 void gather_rows(const uint8_t* src, int64_t row_bytes, const int32_t* ids, int64_t n, uint8_t* dst, void* stream) {
-    if (row_bytes % 16 != 0) { std::fprintf(stderr, "gather_rows: row size must be a multiple of 16\n"); std::exit(1); }
-    gather_rows_kernel<<<48 * 8, 256, 0, (cudaStream_t) stream>>>((const uint4*) src, row_bytes / 16, ids, n, (uint4*) dst);
+    cudaStream_t s = (cudaStream_t) stream;
+    if (row_bytes % 16 == 0)
+        gather_rows_kernel<<<48 * 8, 256, 0, s>>>((const uint4*) src, row_bytes / 16, ids, n, (uint4*) dst);
+    else if (row_bytes % 4 == 0)
+        gather_rows_kernel<<<48 * 8, 256, 0, s>>>((const uint32_t*) src, row_bytes / 4, ids, n, (uint32_t*) dst);
+    else
+        gather_rows_kernel<<<48 * 8, 256, 0, s>>>(src, row_bytes, ids, n, dst);
     check("gather_rows");
 }
 

src/kernels/kv_stream_parity.cpp

@@ -0,0 +1,227 @@
+// src/kernels/kv_stream_parity.cpp - KV streaming (kv_stream.hpp) against a fully resident pool (GPU, no model).
+//
+// Decodes a synthetic sequence twice: into a fully resident pool with the identity page table, and into a streamed
+// layer (host copy + a small VRAM slot pool that must evict constantly). Every few cells a batch of 1-8 queries
+// with selections shaped like the indexer's (whole 4-cell blocks, a partial threshold block, the tail block, half of
+// them recent) is attended in both; `kv_stream_resolve` runs before the streamed attention. Checks:
+//   1. the attention outputs are BITWISE equal (the streamed readers see exactly the resident values);
+//   2. the residency map is consistent after every call (slot_block and page_table invert each other);
+//   3. no call overflowed, and the hit/miss counters add up;
+//   4. a ring (the MTP drafter's layout) restored from the host copy reads the same values as the resident pool.
+// INT8 and FP16 pools both.
+#include "strata/kernels/kv_q8.hpp"
+#include "strata/kernels/kv_stream.hpp"
+#include "strata/kernels/qsa.hpp"
+#include "strata/kernels/qsa_decode_attn.hpp"
+
+#include <cuda_runtime.h>
+
+#include <algorithm>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <random>
+#include <set>
+#include <vector>
+
+namespace k = strata::kernels;
+
+namespace {
+int g_fail = 0;
+void ck(cudaError_t e, const char* w) {
+    if (e != cudaSuccess) { std::fprintf(stderr, "%s: %s\n", w, cudaGetErrorString(e)); std::exit(2); }
+}
+template <typename T> T* dalloc(size_t n) {
+    T* p = nullptr;
+    ck(cudaMalloc(&p, n * sizeof(T) + 64), "malloc");
+    ck(cudaMemset(p, 0, n * sizeof(T) + 64), "memset");
+    return p;
+}
+template <typename T> T* halloc(size_t n) {   // pinned, mapped; returns the device pointer
+    void* h = nullptr;
+    void* d = nullptr;
+    ck(cudaHostAlloc(&h, n * sizeof(T) + 64, cudaHostAllocMapped), "hostalloc");
+    std::memset(h, 0, n * sizeof(T) + 64);
+    ck(cudaHostGetDevicePointer(&d, h, 0), "devptr");
+    return (T*) d;
+}
+
+struct Pools {   // one K/V pool set of `pages` pages
+    k::KvHostPools p;   // reused as a plain pointer bundle
+    void alloc(int64_t pages, const k::QsaShapes& s, bool int8, bool host) {
+        const size_t rows = (size_t) pages * s.n_head_kv * s.page_size;
+        if (int8) {
+            p.k_q = host ? halloc<int8_t>(rows * s.head_dim) : dalloc<int8_t>(rows * s.head_dim);
+            p.v_q = host ? halloc<int8_t>(rows * s.head_dim) : dalloc<int8_t>(rows * s.head_dim);
+            p.k_scale = host ? halloc<uint16_t>(rows * 4) : dalloc<uint16_t>(rows * 4);
+            p.v_scale = host ? halloc<uint16_t>(rows * 4) : dalloc<uint16_t>(rows * 4);
+        } else {
+            p.k_pool = host ? halloc<uint16_t>(rows * s.head_dim) : dalloc<uint16_t>(rows * s.head_dim);
+            p.v_pool = host ? halloc<uint16_t>(rows * s.head_dim) : dalloc<uint16_t>(rows * s.head_dim);
+        }
+    }
+    k::QsaAttnPools attn(const int32_t* table) const {
+        k::QsaAttnPools a;
+        a.k_pool = p.k_pool; a.v_pool = p.v_pool; a.k_q = p.k_q; a.v_q = p.v_q; a.k_scale = p.k_scale;
+        a.v_scale = p.v_scale; a.page_table = table;
+        return a;
+    }
+};
+
+void append(const Pools& pl, const int32_t* table, const int32_t* step, const float* kc, const float* vc,
+            const k::QsaShapes& s, bool int8, const k::KvHostPools* host) {
+    if (int8) k::kv_append_q8_step(pl.p.k_q, pl.p.v_q, pl.p.k_scale, pl.p.v_scale, table, step, kc, vc, s, nullptr, host);
+    else k::kv_append_step(pl.p.k_pool, pl.p.v_pool, table, step, kc, vc, s, nullptr, host);
+}
+
+// A selection like qsa_block_topk's: ascending cells, whole blocks, the tail block's cells, and one partial block.
+std::vector<int32_t> selection(int64_t n_kv, int64_t width, std::mt19937& rng) {
+    std::vector<int32_t> ids;
+    if (n_kv <= width) {
+        for (int64_t c = 0; c < n_kv; ++c) ids.push_back((int32_t) c);
+        return ids;
+    }
+    const int64_t n_bid = n_kv / 4, tail = n_kv - n_bid * 4;
+    std::set<int64_t> blocks;
+    const int64_t want_full = (width - tail) / 4, part = (width - tail) % 4;
+    std::uniform_int_distribution<int64_t> any(0, n_bid - 1), recent(std::max<int64_t>(0, n_bid - 2048), n_bid - 1);
+    while ((int64_t) blocks.size() < want_full + (part ? 1 : 0)) blocks.insert(rng() % 2 ? recent(rng) : any(rng));
+    int64_t partial_block = part ? *std::next(blocks.begin(), (long) (rng() % blocks.size())) : -1;
+    for (int64_t b : blocks) {
+        const int64_t take = b == partial_block ? part : 4;
+        for (int64_t i = 0; i < take; ++i) ids.push_back((int32_t) (b * 4 + i));
+    }
+    for (int64_t c = n_bid * 4; c < n_kv; ++c) ids.push_back((int32_t) c);
+    std::sort(ids.begin(), ids.end());
+    return ids;
+}
+
+bool run(bool int8) {
+    k::QsaShapes s = k::qsa_real_shapes();
+    const int64_t N = 40000, n_blocks = (N + 3) / 4, n_slots = 8 * 516 + 700;   // must evict: slots < blocks
+    const int64_t cap = k::qsa_selection_width(k::kTopkMaxCells, s), NQ = 8;
+    const int64_t H = s.n_head_kv, D = s.head_dim, NH = s.n_head;
+    std::mt19937 rng(int8 ? 7 : 11);
+    std::normal_distribution<float> nd(0.f, 1.f);
+
+    Pools ref, slots, host;
+    ref.alloc(n_blocks, s, int8, false);
+    slots.alloc(n_slots, s, int8, false);
+    host.alloc(n_blocks, s, int8, true);
+    int32_t* ident = dalloc<int32_t>(n_blocks);
+    {
+        std::vector<int32_t> t(n_blocks);
+        for (int64_t i = 0; i < n_blocks; ++i) t[i] = (int32_t) i;
+        ck(cudaMemcpy(ident, t.data(), n_blocks * 4, cudaMemcpyHostToDevice), "ident");
+    }
+    k::KvStreamMap m;
+    m.page_table = dalloc<int32_t>(n_blocks);
+    m.slot_block = dalloc<int32_t>(n_slots); m.slot_stamp = dalloc<int32_t>(n_slots); m.slot_ref = dalloc<int32_t>(n_slots);
+    m.miss_block = dalloc<int32_t>(n_slots); m.miss_slot = dalloc<int32_t>(n_slots);
+    m.ctl = dalloc<int32_t>(k::kKvCtlInts);
+    m.n_blocks = n_blocks; m.n_slots = n_slots;
+    k::kv_stream_reset(m, nullptr);
+
+    float* kc = dalloc<float>(H * D);
+    float* vc = dalloc<float>(H * D);
+    int32_t* step = dalloc<int32_t>(k::kStepCount);
+    int32_t* steps = dalloc<int32_t>(NQ * k::kStepCount);
+    int32_t* ids = dalloc<int32_t>(NQ * cap);
+    float* q = dalloc<float>(NQ * NH * D);
+    const uint64_t scr = k::qsa_decode_attn_scratch_floats(cap, s);
+    float* scratch = dalloc<float>(NQ * scr);
+    float* out_ref = dalloc<float>(NQ * NH * D);
+    float* out_str = dalloc<float>(NQ * NH * D);
+    std::vector<float> hk(H * D), hv(H * D), hq(NQ * NH * D), a((size_t) NQ * NH * D), b2((size_t) NQ * NH * D);
+    std::vector<int32_t> pt(n_blocks), sb(n_slots);
+
+    int batches = 0, bad = 0;
+    for (int64_t pos = 0; pos < N; ++pos) {
+        for (auto& x : hk) x = nd(rng) * ((pos % 17 == 0) ? 30.f : 1.f);
+        for (auto& x : hv) x = nd(rng);
+        const int32_t st[4] = {(int32_t) pos, (int32_t) (pos + 1), (int32_t) ((pos + 1) / 4), 0};
+        ck(cudaMemcpy(kc, hk.data(), hk.size() * 4, cudaMemcpyHostToDevice), "k");
+        ck(cudaMemcpy(vc, hv.data(), hv.size() * 4, cudaMemcpyHostToDevice), "v");
+        ck(cudaMemcpy(step, st, sizeof(st), cudaMemcpyHostToDevice), "step");
+        append(ref, ident, step, kc, vc, s, int8, nullptr);
+        append(slots, m.page_table, step, kc, vc, s, int8, &host.p);
+        if (pos % 131 != 130 && pos != N - 1) continue;
+        // a batch: n_q queries at the last n_q positions (as a verify window), each with its own selection
+        const int n_q = 1 + (int) (rng() % NQ);
+        std::vector<int32_t> hids((size_t) NQ * cap, 0), hst(NQ * 4, 0);
+        for (int t = 0; t < n_q; ++t) {
+            const int64_t p = std::max<int64_t>(0, pos - (n_q - 1 - t)), n_kv = p + 1;
+            const int64_t width = std::min<int64_t>(n_kv, cap);
+            const std::vector<int32_t> sel = selection(n_kv, width, rng);
+            if ((int64_t) sel.size() != width) { std::fprintf(stderr, "selection size %zu != %lld\n", sel.size(), (long long) width); return false; }
+            std::copy(sel.begin(), sel.end(), hids.begin() + (size_t) t * cap);
+            hst[t * 4 + 0] = (int32_t) p; hst[t * 4 + 1] = (int32_t) n_kv; hst[t * 4 + 2] = (int32_t) (n_kv / 4);
+            hst[t * 4 + 3] = (int32_t) width;
+        }
+        for (auto& x : hq) x = nd(rng);
+        ck(cudaMemcpy(ids, hids.data(), hids.size() * 4, cudaMemcpyHostToDevice), "ids");
+        ck(cudaMemcpy(steps, hst.data(), hst.size() * 4, cudaMemcpyHostToDevice), "steps");
+        ck(cudaMemcpy(q, hq.data(), hq.size() * 4, cudaMemcpyHostToDevice), "q");
+        k::qsa_decode_attn_batch(q, ref.attn(ident), ids, steps, cap, s, scratch, out_ref, n_q, nullptr);
+        k::kv_stream_resolve(m, slots.attn(m.page_table), host.p, int8, ids, steps, n_q, cap, s, nullptr);
+        k::qsa_decode_attn_batch(q, slots.attn(m.page_table), ids, steps, cap, s, scratch, out_str, n_q, nullptr);
+        ck(cudaDeviceSynchronize(), "batch");
+        ck(cudaMemcpy(a.data(), out_ref, (size_t) n_q * NH * D * 4, cudaMemcpyDeviceToHost), "a");
+        ck(cudaMemcpy(b2.data(), out_str, (size_t) n_q * NH * D * 4, cudaMemcpyDeviceToHost), "b");
+        if (std::memcmp(a.data(), b2.data(), (size_t) n_q * NH * D * 4) != 0) {
+            if (bad++ < 5) std::fprintf(stderr, "  %s pos %lld n_q %d: streamed attention differs\n", int8 ? "int8" : "fp16", (long long) pos, n_q);
+        }
+        // the map inverts itself
+        ck(cudaMemcpy(pt.data(), m.page_table, n_blocks * 4, cudaMemcpyDeviceToHost), "pt");
+        ck(cudaMemcpy(sb.data(), m.slot_block, n_slots * 4, cudaMemcpyDeviceToHost), "sb");
+        int64_t resident = 0;
+        for (int64_t bk = 0; bk < n_blocks; ++bk) {
+            if (pt[bk] < -1 || pt[bk] >= n_slots || (pt[bk] >= 0 && sb[pt[bk]] != bk)) {
+                if (bad++ < 5) std::fprintf(stderr, "  map broken at block %lld (table %d)\n", (long long) bk, pt[bk]);
+                break;
+            }
+            resident += pt[bk] >= 0;
+        }
+        for (int64_t sl = 0; sl < n_slots; ++sl)
+            if (sb[sl] >= 0 && pt[sb[sl]] != sl) { if (bad++ < 5) std::fprintf(stderr, "  slot %lld not in the table\n", (long long) sl); break; }
+        ++batches;
+    }
+    const k::KvStreamCounters c = k::kv_stream_counters(m);
+    std::printf("  %s: %d batches, %llu block lookups, %llu misses (%.1f%% hit), overflow %d, %d failures\n",
+                int8 ? "int8" : "fp16", batches, (unsigned long long) c.lookups, (unsigned long long) c.misses,
+                c.lookups ? 100.0 * (double) (c.lookups - c.misses) / (double) c.lookups : 0.0, (int) c.overflow, bad);
+    if (c.overflow || c.calls != (uint64_t) batches || c.misses == 0 || c.misses >= c.lookups) ++bad;
+
+    // 4. a ring over the same host copy: blocks [b1 - R, b1) restored, read through `b % R`
+    {
+        const int64_t R = 1500, b1 = n_blocks, b0 = b1 - R;
+        Pools ring;
+        ring.alloc(R, s, int8, false);
+        int32_t* rt = dalloc<int32_t>(n_blocks);
+        k::kv_ring_table(rt, n_blocks, R, nullptr);
+        k::kv_ring_restore(ring.attn(rt), host.p, int8, b0, b1, R, s, nullptr);
+        const int64_t width = cap;
+        std::vector<int32_t> hids((size_t) cap), hst = {(int32_t) (N - 1), (int32_t) N, (int32_t) (N / 4), (int32_t) width};
+        for (int64_t i = 0; i < width; ++i) hids[i] = (int32_t) (N - width + i);   // the window's last cells
+        ck(cudaMemcpy(ids, hids.data(), hids.size() * 4, cudaMemcpyHostToDevice), "ids");
+        ck(cudaMemcpy(steps, hst.data(), 16, cudaMemcpyHostToDevice), "steps");
+        k::qsa_decode_attn_batch(q, ref.attn(ident), ids, steps, cap, s, scratch, out_ref, 1, nullptr);
+        k::qsa_decode_attn_batch(q, ring.attn(rt), ids, steps, cap, s, scratch, out_str, 1, nullptr);
+        ck(cudaDeviceSynchronize(), "ring");
+        ck(cudaMemcpy(a.data(), out_ref, (size_t) NH * D * 4, cudaMemcpyDeviceToHost), "a");
+        ck(cudaMemcpy(b2.data(), out_str, (size_t) NH * D * 4, cudaMemcpyDeviceToHost), "b");
+        const bool ok = std::memcmp(a.data(), b2.data(), (size_t) NH * D * 4) == 0;
+        std::printf("  %s ring restore: %s\n", int8 ? "int8" : "fp16", ok ? "identical" : "DIFFERS");
+        if (!ok) ++bad;
+    }
+    return bad == 0;
+}
+}  // namespace
+
+int main() {
+    std::printf("kv_stream_parity: streamed vs resident KV, bitwise\n");
+    const bool a = run(true), b = run(false);
+    if (!a || !b) ++g_fail;
+    std::printf(g_fail ? "FAIL\n" : "PASS\n");
+    return g_fail ? 1 : 0;
+}

src/kernels/sampler_parity.cpp

@@ -12,6 +12,7 @@
 
 #include <cuda_runtime.h>
 
+#include <algorithm>
 #include <cmath>
 #include <cstdio>
 #include <cstdlib>
@@ -126,6 +127,139 @@ int run(const char* name, const std::vector<float>& logits, int n_tokens, const
     return bad;
 }
 
+// The Philox draw, host side - a transcription of the kernel's `philox_uniform` so the SAMPLED pick (not
+// just the greedy argmax) can be pinned against a reference.  `__umulhi(a, b)` is the high half of a 32x32
+// multiply, spelled `(uint32_t)(((uint64_t) a * b) >> 32)` here.
+struct PhiloxRound {
+    uint32_t& c0; uint32_t& c1; uint32_t& c2; uint32_t& c3;
+    void step(uint32_t k0, uint32_t k1) const {
+        const uint32_t hi0 = (uint32_t) (((uint64_t) 0x9E3779B9u * c0) >> 32);
+        const uint32_t hi1 = (uint32_t) (((uint64_t) 0xBB67AE85u * c2) >> 32);
+        const uint32_t lo0 = 0x9E3779B9u * c0;
+        const uint32_t lo1 = 0xBB67AE85u * c2;
+        const uint32_t n0 = hi1 ^ c1 ^ k0;
+        const uint32_t n1 = lo1;
+        const uint32_t n2 = hi0 ^ c3 ^ k1;
+        const uint32_t n3 = lo0;
+        c0 = n0; c1 = n1; c2 = n2; c3 = n3;
+    }
+};
+
+float host_philox_uniform(uint64_t seed, uint64_t counter) {
+    uint32_t c0 = (uint32_t) counter, c1 = (uint32_t) (counter >> 32);
+    uint32_t c2 = (uint32_t) seed, c3 = (uint32_t) (seed >> 32);
+    PhiloxRound r{c0, c1, c2, c3};
+    for (int i = 0; i < 10; ++i) r.step((uint32_t) i, 0u);
+    return (float) (c0 >> 8) * (1.0f / 16777216.0f);
+}
+
+// The full SAMPLED chain, host side - the kernel's `sampler_kernel` in serial form: penalties on the raw
+// logits during the top_k selection (ties to the lowest index), the min_p prefix cut, top_p's cut in double,
+// then penalties AGAIN on the temperature-scaled survivors (docs/sampling.md's after-temperature pass), and
+// one Philox draw at (seed, counter + row).
+int sampled_reference(const std::vector<float>& l, const std::vector<int>& hist,
+                      const strata::kernels::SamplerParams& p, int row) {
+    auto penal = [&](float logit, int count) {
+        if (count <= 0) return logit;
+        if (logit <= 0.0f) logit *= p.penalty_repeat; else logit /= p.penalty_repeat;
+        logit -= (float) count * p.penalty_freq + (count > 0 ? 1.0f : 0.0f) * p.penalty_present;
+        return logit;
+    };
+    auto count = [&](int v) { int c = 0; for (int h : hist) if (h == v) ++c; return c; };
+    const int nv = (int) l.size();
+    const int KMAX = 64;
+    const int k = p.top_k > 0 ? (p.top_k < KMAX ? p.top_k : KMAX) : 0;
+    std::vector<int> sel_ids;
+    std::vector<float> sel_logit;
+    std::vector<char> taken((size_t) nv, 0);
+    for (int i = 0; i < k; ++i) {
+        int best = -1; float bv = 0;
+        for (int v = 0; v < nv; ++v) {
+            if (taken[(size_t) v]) continue;
+            const float s = penal(l[(size_t) v], count(v));
+            if (best < 0 || s > bv) { best = v; bv = s; }
+        }
+        taken[(size_t) best] = 1;
+        sel_ids.push_back(best); sel_logit.push_back(bv);
+    }
+    int n_minp = (int) sel_ids.size();
+    if (p.min_p > 0.0f) {
+        const float thresh = sel_logit[0] + std::log(p.min_p);
+        for (int i = 0; i < (int) sel_ids.size(); ++i)
+            if (sel_logit[(size_t) i] < thresh) { n_minp = i; break; }
+    }
+    int n_keep = n_minp;
+    if (p.top_p < 1.0f) {
+        double sum = 0.0;
+        for (int i = 0; i < n_minp; ++i) sum += std::exp((double) sel_logit[(size_t) i] - (double) sel_logit[0]);
+        double cum = 0.0;
+        int cut = n_minp;
+        for (int i = 0; i < n_minp; ++i) {
+            cum += std::exp((double) sel_logit[(size_t) i] - (double) sel_logit[0]) / sum;
+            if (cum >= (double) p.top_p) { cut = i + 1; break; }
+        }
+        if (cut < p.min_keep) cut = p.min_keep < n_minp ? p.min_keep : n_minp;
+        n_keep = cut;
+    }
+    const float inv_t = p.temperature > 0.0f ? 1.0f / p.temperature : 0.0f;
+    auto scaled = [&](int i) { return penal(sel_logit[(size_t) i] * inv_t, count(sel_ids[(size_t) i])); };
+    float smx = scaled(0);
+    for (int i = 1; i < n_keep; ++i) smx = std::fmax(smx, scaled(i));
+    double sum = 0.0;
+    for (int i = 0; i < n_keep; ++i) sum += std::exp((double) scaled(i) - (double) smx);
+    const float u = host_philox_uniform(p.seed, p.counter + (uint64_t) row);
+    double cum = 0.0;
+    int pick = sel_ids[(size_t) (n_keep - 1)];
+    for (int i = 0; i < n_keep; ++i) {
+        cum += std::exp((double) scaled(i) - (double) smx) / sum;
+        if ((double) u < cum) { pick = sel_ids[(size_t) i]; break; }
+    }
+    return pick;
+}
+
+// Survivors after the min_p + top_p cuts, in the sampled chain - the quantity an order or a threshold
+// actually changes, used to assert a fixture can SEE the feature before asserting the kernel matches.
+int sampled_cut(const std::vector<float>& l, const std::vector<int>& hist, const strata::kernels::SamplerParams& p) {
+    auto penal = [&](float logit, int count) {
+        if (count <= 0) return logit;
+        if (logit <= 0.0f) logit *= p.penalty_repeat; else logit /= p.penalty_repeat;
+        logit -= (float) count * p.penalty_freq + (count > 0 ? 1.0f : 0.0f) * p.penalty_present;
+        return logit;
+    };
+    auto count = [&](int v) { int c = 0; for (int h : hist) if (h == v) ++c; return c; };
+    const int nv = (int) l.size();
+    const int KMAX = 64;
+    const int k = p.top_k > 0 ? (p.top_k < KMAX ? p.top_k : KMAX) : 0;
+    std::vector<float> sel;
+    std::vector<char> taken((size_t) nv, 0);
+    for (int i = 0; i < k; ++i) {
+        int best = -1; float bv = 0;
+        for (int v = 0; v < nv; ++v) {
+            if (taken[(size_t) v]) continue;
+            const float s = penal(l[(size_t) v], count(v));
+            if (best < 0 || s > bv) { best = v; bv = s; }
+        }
+        taken[(size_t) best] = 1; sel.push_back(bv);
+    }
+    int n_minp = (int) sel.size();
+    if (p.min_p > 0.0f) {
+        const float thresh = sel[0] + std::log(p.min_p);
+        for (int i = 0; i < (int) sel.size(); ++i)
+            if (sel[(size_t) i] < thresh) { n_minp = i; break; }
+    }
+    if (p.top_p >= 1.0f) return n_minp;
+    double sum = 0.0;
+    for (int i = 0; i < n_minp; ++i) sum += std::exp((double) sel[(size_t) i] - (double) sel[0]);
+    double cum = 0.0;
+    int cut = n_minp;
+    for (int i = 0; i < n_minp; ++i) {
+        cum += std::exp((double) sel[(size_t) i] - (double) sel[0]) / sum;
+        if (cum >= (double) p.top_p) { cut = i + 1; break; }
+    }
+    if (cut < p.min_keep) cut = p.min_keep < n_minp ? p.min_keep : n_minp;
+    return cut;
+}
+
 }  // namespace
 
 int main(int argc, char** argv) {
@@ -326,6 +460,116 @@ int main(int argc, char** argv) {
         bad += run("T=1 greedy=true  is the argmax", l, NT3, p2, want);
     }
 
+    // ---- fixture 6: PENALTIES IN THE SAMPLED CHAIN.  Fixture 4 pins the greedy (argmax) path; the sampled
+    // chain applies penalties a SECOND time - after the temperature scaling, on the survivors - so it gets its
+    // own reference (the full chain with the host Philox) and its own observability check: with the penalties
+    // on, the history row's favourite must LOSE a pick it would win penalty-free.
+    {
+        const int NV2 = 8, NT2 = 2;
+        strata::kernels::SamplerParams p;
+        p.top_k = 5; p.top_p = 0.9f; p.temperature = 0.8f; p.seed = 9; p.counter = 0;
+        p.penalty_last_n = 4; p.penalty_repeat = 3.0f; p.penalty_freq = 0.2f; p.penalty_present = 0.6f;
+
+        std::vector<float> l((size_t) NV2 * NT2, -8.0f);
+        std::vector<int> hist((size_t) NT2 * 4, -1);
+        for (int t = 0; t < NT2; ++t) {
+            float* row = l.data() + (size_t) t * NV2;
+            row[0] = 9.0f; row[1] = 4.5f; row[2] = 4.4f; row[3] = 4.3f;   // token 0 leads clean (9 vs 4.5)
+            hist[(size_t) t * 4 + 0] = 0;                                  // and falls to 2.2/2.0 penalised
+            hist[(size_t) t * 4 + 1] = t == 1 ? 0 : -1;                    // (repeat 3, freq, presence)
+        }
+        std::vector<int> want((size_t) NT2), clean((size_t) NT2);
+        strata::kernels::SamplerParams clean_p = p;
+        clean_p.penalty_last_n = 0; clean_p.penalty_repeat = 1.0f;
+        clean_p.penalty_freq = 0.0f; clean_p.penalty_present = 0.0f;
+        for (int t = 0; t < NT2; ++t) {
+            const std::vector<float> row(l.begin() + (size_t) t * NV2, l.begin() + (size_t) (t + 1) * NV2);
+            const std::vector<int> h(hist.begin() + (size_t) t * 4, hist.begin() + (size_t) (t + 1) * 4);
+            want[(size_t) t] = sampled_reference(row, h, p, t);
+            clean[(size_t) t] = sampled_reference(row, h, clean_p, t);
+        }
+        const bool visible = want[0] != clean[0] || want[1] != clean[1];
+        std::printf("  %-34s %s (penalised picks %d/%d, clean %d/%d)\n",
+                    "sampled penalties are observable", visible ? "yes" : "*** NO ***", want[0], want[1],
+                    clean[0], clean[1]);
+        if (!visible) ++bad;
+        else bad += run("sampled chain: penalties + top_k/p", l, NT2, p, want, hist, 4);
+    }
+
+    // ---- fixture 7: MIN_P.  The cut is a PREFIX of the descending top_k list (logit >= max + log(min_p)),
+    // so the fixture asserts the survivor count moves with the threshold (the observability half) and that
+    // the kernel's pick equals the reference's through the full sampled chain (the correctness half).
+    {
+        const int NV3 = 8, NT3 = 2;
+        std::vector<float> l((size_t) NV3 * NT3, -8.0f);
+        for (int t = 0; t < NT3; ++t) {
+            float* row = l.data() + (size_t) t * NV3;
+            row[0] = 4.0f; row[1] = 3.5f; row[2] = 3.2f; row[3] = 3.1f;   // gaps keep the cut off the
+            row[4] = 2.0f;                                                // logf/rounding knife edge
+        }
+        strata::kernels::SamplerParams base;
+        base.top_k = 6; base.top_p = 1.0f; base.temperature = 0.9f; base.seed = 77;
+
+        int c0 = 0, c05 = 0, c09 = 0;
+        for (int t = 0; t < NT3; ++t) {
+            const std::vector<float> row(l.begin() + (size_t) t * NV3, l.begin() + (size_t) (t + 1) * NV3);
+            strata::kernels::SamplerParams q = base; q.min_p = 0.0f;
+            c0 += sampled_cut(row, {}, q);
+            q.min_p = 0.5f; c05 += sampled_cut(row, {}, q);
+            q.min_p = 0.9f; c09 += sampled_cut(row, {}, q);
+        }
+        const bool visible = c0 > c05 && c05 > c09 && c09 >= NT3;
+        std::printf("  %-34s %s (survivors: min_p 0 -> %d, 0.5 -> %d, 0.9 -> %d)\n",
+                    "min_p cut is observable", visible ? "yes" : "*** NO ***", c0, c05, c09);
+        if (!visible) ++bad;
+
+        for (float mp : {0.0f, 0.5f, 0.9f}) {
+            strata::kernels::SamplerParams q = base; q.min_p = mp;
+            std::vector<int> want((size_t) NT3);
+            for (int t = 0; t < NT3; ++t) {
+                const std::vector<float> row(l.begin() + (size_t) t * NV3, l.begin() + (size_t) (t + 1) * NV3);
+                want[(size_t) t] = sampled_reference(row, {}, q, t);
+            }
+            char name[64];
+            std::snprintf(name, sizeof name, "sampled chain: min_p=%.1f", (double) mp);
+            bad += run(name, l, NT3, q, want);
+        }
+    }
+
+    // ---- fixture 8: THE PENALTY WINDOW IS THE TAIL.  With an 8-entry history and penalty_last_n = 4, only
+    // the LAST four entries count: a token punished in the old half must come back to full strength, and one
+    // punished in the tail half stays down.  The reference counts the same tail; the observability check runs
+    // the reference once more WITHOUT the clamp (counting all 8) and requires the picks to differ.
+    {
+        const int NV4 = 8;
+        strata::kernels::SamplerParams p;
+        p.top_k = 0; p.top_p = 1.0f; p.temperature = 1.0f; p.greedy = true;
+        p.penalty_last_n = 4; p.penalty_repeat = 3.0f; p.penalty_freq = 0.3f; p.penalty_present = 0.5f;
+
+        std::vector<float> l((size_t) NV4, -8.0f);
+        l[0] = 6.0f; l[3] = 6.5f;                       // token 0 leads clean; token 3 is the tail offender
+        std::vector<int> hist = {0, 0, 0, 0, 3, 3, 3, 3};   // token 0 old (out), token 3 in the tail
+
+        auto pick_clamped = [&](bool clamp) {
+            int best = 0; float bv = 0; bool first = true;
+            for (int v = 0; v < NV4; ++v) {
+                int c = 0;
+                for (int i = 0; i < (clamp ? 4 : 8); ++i) if (hist[(size_t) (8 - (clamp ? 4 : 8) + i)] == v) ++c;
+                float logit = l[(size_t) v];
+                if (c > 0) { logit = logit <= 0.0f ? logit * p.penalty_repeat : logit / p.penalty_repeat;
+                             logit -= (float) c * p.penalty_freq + p.penalty_present; }
+                if (first || logit > bv) { bv = logit; best = v; first = false; }
+            }
+            return best;
+        };
+        const int want = pick_clamped(true), unclamped = pick_clamped(false);
+        const bool visible = want != unclamped;
+        std::printf("  %-34s %s (clamped pick %d, full-history pick %d)\n",
+                    "penalty window clamp is observable", visible ? "yes" : "*** NO ***", want, unclamped);
+        if (!visible) ++bad;
+        else bad += run("penalty window: tail only", {l.begin(), l.end()}, 1, p, {want}, hist, 8);
+    }
+
     // A continuous stream and individual decode calls consume the same draw counters.
     {
         constexpr int count = 32, vocab = 16;

src/prefill/kernels.cu

@@ -335,10 +335,12 @@ __global__ void gate_attn_kernel(const float* __restrict__ a, const float* __res
     o16[i] = hf(a[i] * (1.0f / (1.0f + expf(-qf[t * 24 * 512 + h * 512 + 256 + d]))));
 }
 
-// one block per (token, kv head, 64-value group); 64 threads
+// one block per (token, kv head, 64-value group); 64 threads. KV streaming: the pool page only if the block is
+// resident (table >= 0), and the host copy and the prompt path's staging pool (both identity layout) when given.
 __global__ void kv_append_kernel(const float* __restrict__ K, const float* __restrict__ V, int64_t pos0,
                                  const int32_t* __restrict__ table, int64_t page_size, uint16_t* k_pool,
-                                 uint16_t* v_pool, int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale) {
+                                 uint16_t* v_pool, int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale,
+                                 strata::kernels::KvHostPools host, strata::kernels::KvHostPools stage) {
     const int64_t t = blockIdx.x;
     const int kvh = blockIdx.y, g = blockIdx.z >> 1;
     const bool is_v = (blockIdx.z & 1) != 0;
@@ -347,8 +349,12 @@ __global__ void kv_append_kernel(const float* __restrict__ K, const float* __res
     const int64_t pos = pos0 + t;
     const int64_t page = table[pos / page_size];
     const int64_t row = (page * 2 + kvh) * page_size + pos % page_size;
+    const int64_t row_id = ((pos / page_size) * 2 + kvh) * page_size + pos % page_size;
     if (k_pool != nullptr) {
-        (is_v ? v_pool : k_pool)[row * 256 + d] = hf(x);
+        const uint16_t h = hf(x);
+        if (page >= 0) (is_v ? v_pool : k_pool)[row * 256 + d] = h;
+        if (host.k_pool != nullptr) (is_v ? host.v_pool : host.k_pool)[row_id * 256 + d] = h;
+        if (stage.k_pool != nullptr) (is_v ? stage.v_pool : stage.k_pool)[row_id * 256 + d] = h;
         return;
     }
     float a = fabsf(x);
@@ -361,13 +367,27 @@ __global__ void kv_append_kernel(const float* __restrict__ K, const float* __res
     const float sf = __half2float(__ushort_as_half(sb));
     int q = 0;
     if (sf > 0.0f) { q = __float2int_rn(x / sf); q = q < -127 ? -127 : (q > 127 ? 127 : q); }
-    (is_v ? v_q : k_q)[row * 256 + d] = (int8_t) q;
-    if (threadIdx.x == 0) (is_v ? v_scale : k_scale)[row * 4 + g] = sb;
+    if (page >= 0) {
+        (is_v ? v_q : k_q)[row * 256 + d] = (int8_t) q;
+        if (threadIdx.x == 0) (is_v ? v_scale : k_scale)[row * 4 + g] = sb;
+    }
+    if (host.k_q != nullptr) {
+        (is_v ? host.v_q : host.k_q)[row_id * 256 + d] = (int8_t) q;
+        if (threadIdx.x == 0) (is_v ? host.v_scale : host.k_scale)[row_id * 4 + g] = sb;
+    }
+    if (stage.k_q != nullptr) {
+        (is_v ? stage.v_q : stage.k_q)[row_id * 256 + d] = (int8_t) q;
+        if (threadIdx.x == 0) (is_v ? stage.v_scale : stage.k_scale)[row_id * 4 + g] = sb;
+    }
 }
 __global__ void to_f16_kernel(const float* __restrict__ x, uint16_t* __restrict__ y, int64_t n) {
     for (int64_t i = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; i < n; i += (int64_t) gridDim.x * blockDim.x)
         y[i] = hf(x[i]);
 }
+__global__ void round_f16_kernel(const float* __restrict__ x, float* __restrict__ y, int64_t n) {
+    for (int64_t i = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; i < n; i += (int64_t) gridDim.x * blockDim.x)
+        y[i] = __half2float(__float2half_rn(x[i]));
+}
 __global__ void to_bf16_kernel(const float* __restrict__ x, uint16_t* __restrict__ y, int64_t n) {
     for (int64_t i = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; i < n; i += (int64_t) gridDim.x * blockDim.x)
         y[i] = bf(x[i]);
@@ -377,10 +397,11 @@ __global__ void to_bf16_kernel(const float* __restrict__ x, uint16_t* __restrict
 
 void kv_append(const float* K, const float* V, int64_t T, int64_t pos0, const int32_t* page_table, int64_t page_size,
                uint16_t* k_pool, uint16_t* v_pool, int8_t* k_q, int8_t* v_q, uint16_t* k_scale, uint16_t* v_scale,
-               void* stream) {
+               void* stream, const strata::kernels::KvHostPools* host, const strata::kernels::KvHostPools* stage) {
     if (T <= 0) return;
-    kv_append_kernel<<<dim3((unsigned) T, 2, 8), 64, 0, (cudaStream_t) stream>>>(K, V, pos0, page_table, page_size, k_pool,
-                                                                                  v_pool, k_q, v_q, k_scale, v_scale);
+    kv_append_kernel<<<dim3((unsigned) T, 2, 8), 64, 0, (cudaStream_t) stream>>>(
+        K, V, pos0, page_table, page_size, k_pool, v_pool, k_q, v_q, k_scale, v_scale,
+        host ? *host : strata::kernels::KvHostPools{}, stage ? *stage : strata::kernels::KvHostPools{});
     check("kv_append");
 }
 void to_f16(const float* x, uint16_t* y, int64_t n, void* stream) {
@@ -388,6 +409,11 @@ void to_f16(const float* x, uint16_t* y, int64_t n, void* stream) {
     to_f16_kernel<<<(unsigned) ((n + 255) / 256 < 4096 ? (n + 255) / 256 : 4096), 256, 0, (cudaStream_t) stream>>>(x, y, n);
     check("to_f16");
 }
+void round_f16(const float* x, float* y, int64_t n, void* stream) {
+    if (n <= 0) return;
+    round_f16_kernel<<<(unsigned) ((n + 255) / 256 < 4096 ? (n + 255) / 256 : 4096), 256, 0, (cudaStream_t) stream>>>(x, y, n);
+    check("round_f16");
+}
 void to_bf16(const float* x, uint16_t* y, int64_t n, void* stream) {
     if (n <= 0) return;
     to_bf16_kernel<<<(unsigned) ((n + 255) / 256 < 4096 ? (n + 255) / 256 : 4096), 256, 0, (cudaStream_t) stream>>>(x, y, n);

src/prefill/prefill.cpp

@@ -9,6 +9,8 @@
 #include "strata/kernels/iq_kernels.hpp"
 #include "strata/kernels/cpu/expert_layout.hpp"
 #include "strata/kernels/qsa.hpp"
+#include "strata/kernels/kv_stream.hpp"
+#include "strata/core/layer.hpp"
 #include "strata/kernels/qsa_decode_attn.hpp"
 #include "strata/kernels/qsa_select.hpp"
 #include "strata/prefill/gemm.hpp"
@@ -111,8 +113,44 @@ struct Prefill::Impl {
     std::vector<uint32_t> ple_rows;
     float* ple_norm = nullptr;
     PrefillStats* stats = nullptr;
+    // KV streaming: one layer's whole K/V, staged from the host copy per layer and chunk (identity layout)
+    strata::kernels::KvHostPools stage;
+    int32_t* ident_table = nullptr;
 };
 
+namespace {
+// the staging pool of a streamed session: every page of one layer (same sequence in init and bytes_needed)
+// STRATA_KV_STAGE_OWN (A/B only): the staging pool gets its own allocation instead of borrowed expert slots, so a
+// streamed run lends the prompt path exactly the slots a resident one does (a lent expert runs on the CPU, which
+// rounds differently: without this an A/B compares two expert placements as well as two KV placements)
+bool stage_own() { static const bool v = std::getenv("STRATA_KV_STAGE_OWN") != nullptr; return v; }
+void take_stage(Alloc& o_borrowed, const core::SessionState& ss, const strata::kernels::QsaShapes& s,
+                strata::kernels::KvHostPools& st, bool& ok) {
+    const core::QsaState& q0 = ss.qsa_states[0];
+    if (q0.kv_mode != 1) return;
+    if (stage_own() && o_borrowed.count_only) return;
+    Alloc own;
+    own.owned = o_borrowed.owned;
+    Alloc& o = stage_own() ? own : o_borrowed;
+    const size_t rows = (size_t) q0.n_pages * s.n_head_kv * s.page_size;
+    if (q0.kv_int8) {
+        st.k_q = o.take<int8_t>(rows * s.head_dim, ok);
+        st.v_q = o.take<int8_t>(rows * s.head_dim, ok);
+        st.k_scale = o.take<uint16_t>(rows * (s.head_dim / 64), ok);
+        st.v_scale = o.take<uint16_t>(rows * (s.head_dim / 64), ok);
+    } else {
+        st.k_pool = o.take<uint16_t>(rows * s.head_dim, ok);
+        st.v_pool = o.take<uint16_t>(rows * s.head_dim, ok);
+    }
+}
+strata::kernels::QsaAttnPools pools_of(const strata::kernels::KvHostPools& h, const int32_t* table) {
+    strata::kernels::QsaAttnPools p;
+    p.k_pool = h.k_pool; p.v_pool = h.v_pool; p.k_q = h.k_q; p.v_q = h.v_q; p.k_scale = h.k_scale; p.v_scale = h.v_scale;
+    p.page_table = table;
+    return p;
+}
+}  // namespace
+
 Prefill::Prefill() : impl_(new Impl) {}
 Prefill::~Prefill() {
     if (!impl_) return;
@@ -189,6 +227,17 @@ bool Prefill::init(const core::WeightTable& wt, const core::ModelGeometry& g, co
     }
     m.ple_emb = o.take<float>(T * N, ok);
     m.ple_norm = o.take<float>((size_t) strata::kernels::NG_HC_DIM, ok);
+    take_stage(o, ss, s, m.stage, ok);
+    if (ok && m.stage.present()) {
+        const int64_t pages = ss.qsa_states[0].n_pages;
+        std::vector<int32_t> ident((size_t) pages);
+        for (int64_t i = 0; i < pages; ++i) ident[(size_t) i] = (int32_t) i;
+        if (cudaMalloc((void**) &m.ident_table, ident.size() * 4) != cudaSuccess ||
+            cudaMemcpy(m.ident_table, ident.data(), ident.size() * 4, cudaMemcpyHostToDevice) != cudaSuccess)
+            ok = false;
+        else
+            m.owned.push_back(m.ident_table);
+    }
     m.ids_host.resize(T * K); m.slot_host.resize(T * K); m.src_host.resize(T * K); m.cnt.resize(NE); m.off.resize(NE + 1);
     m.ple_emb_host.resize(T * N); m.ple_rows.resize(T * strata::kernels::PLE_N_HEADS);
     if (!ok) { err = "prefill: device buffers for a chunk of " + std::to_string(chunk) + " tokens do not fit"; return false; }
@@ -223,6 +272,8 @@ uint64_t Prefill::bytes_needed(const core::ModelGeometry& g, const core::Session
     for (int i = 0; i < STAGE; ++i) o.take<uint8_t>((size_t) MAXBLOB(), ok);
     f(T * N);
     f((size_t) strata::kernels::NG_HC_DIM);
+    strata::kernels::KvHostPools stage;
+    take_stage(o, ss, s, stage, ok);
     return o.used + (8u << 20);   // alignment slack
 }
 
@@ -369,8 +420,15 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                     if (!bf16_proj(m.gemm, wiq, m.mixed_bf, m.q_idx, T, v.name("indexer.q_proj.weight"), err)) return false;
                     rms_rows(m.Kc, (const float*) wkn->data, T * 2, 256, 256, EPS, m.cs);
                     rope(m.Kc, T, 2, 256, 512, p0, (float) strata::kernels::qsa_freq_base(), m.cs);
+                    // KV streaming: this layer's cells [0, p0) come in from the host copy to the staging pool, and the
+                    // chunk's cells go to the host copy, the staging pool, and the VRAM slots of resident blocks
+                    const bool staged = st.kv_mode == 1;
+                    if (staged)
+                        strata::kernels::kv_stage_from_host(pools_of(m.stage, m.ident_table), st.host, st.kv_int8,
+                                                            (p0 + s.page_size - 1) / s.page_size, s, m.cs);
                     kv_append(m.Kc, m.Vc, T, p0, 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.kv_int8 ? nullptr : st.v_pool, st.k_q, st.v_q, st.k_scale, st.v_scale, m.cs,
+                              &st.host, staged ? &m.stage : nullptr);
                     split_q(m.Qf, m.q, T, m.cs);
                     rms_rows(m.q, (const float*) wqn->data, T * 24, 256, 256, EPS, m.cs);
                     rope(m.q, T, 24, 256, 6144, p0, (float) strata::kernels::qsa_freq_base(), m.cs);
@@ -395,10 +453,73 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
                         strata::kernels::qsa_block_topk(m.sel_scores, steps0, nb, m.max_blocks, m.cap, s,
                                                         m.sel_ids + t0 * m.cap, m.cs);
                     }
-                    strata::kernels::QsaAttnPools pools;
-                    pools.page_table = st.page_table;
-                    if (st.kv_int8) { pools.k_q = st.k_q; pools.v_q = st.v_q; pools.k_scale = st.k_scale; pools.v_scale = st.v_scale; }
-                    else { pools.k_pool = st.k_pool; pools.v_pool = st.v_pool; }
+                    // STRATA_IDX_FP16_CHECK: would FP16 pooled indexer keys select the same cells? (the KV-streaming
+                    // design's last question). Every query is selected again from the pooled keys and `dead` rounded
+                    // to fp16 (exactly what an fp16 store reads back); the agreement with the fp32 selection is
+                    // printed cumulatively after each chunk's last QSA layer. Debug: syncs per layer.
+                    if (static const bool f16chk = std::getenv("STRATA_IDX_FP16_CHECK") != nullptr; f16chk) {
+                        static float *pooled16 = nullptr, *dead16 = nullptr;
+                        static int32_t* ids16 = nullptr;
+                        static double shared = 0, cells = 0;
+                        static long long queries = 0, same = 0, sel_queries = 0;
+                        const int64_t rows = st.idx_pooled_rows;
+                        if (pooled16 == nullptr &&
+                            (cudaMalloc((void**) &pooled16, (size_t) rows * s.idx_dim * 4) != cudaSuccess ||
+                             cudaMalloc((void**) &dead16, (size_t) s.idx_dim * 4) != cudaSuccess ||
+                             cudaMalloc((void**) &ids16, (size_t) (m.T * m.cap) * 4) != cudaSuccess)) {
+                            err = "STRATA_IDX_FP16_CHECK: no room for its buffers";
+                            return false;
+                        }
+                        round_f16(st.idx_pooled, pooled16, rows * s.idx_dim, m.cs);
+                        round_f16(st.idx_dead, dead16, s.idx_dim, m.cs);
+                        for (int64_t t0 = 0; t0 < T; t0 += m.sel_batch) {
+                            const int64_t nb = std::min(m.sel_batch, T - t0);
+                            const int32_t* steps0 = m.steps_dev + t0 * strata::kernels::kStepCount;
+                            strata::kernels::qsa_block_scores(pooled16, dead16, m.q_idx + t0 * 512, steps0, nb,
+                                                              m.max_blocks, s, m.sel_scores, m.cs);
+                            strata::kernels::qsa_block_topk(m.sel_scores, steps0, nb, m.max_blocks, m.cap, s,
+                                                            ids16 + t0 * m.cap, m.cs);
+                        }
+                        std::vector<int32_t> a((size_t) (T * m.cap)), b((size_t) (T * m.cap));
+                        cudaMemcpyAsync(a.data(), m.sel_ids, a.size() * 4, cudaMemcpyDeviceToHost, m.cs);
+                        cudaMemcpyAsync(b.data(), ids16, b.size() * 4, cudaMemcpyDeviceToHost, m.cs);
+                        cudaStreamSynchronize(m.cs);
+                        for (int64_t t = 0; t < T; ++t) {
+                            const int64_t w = m.steps_host[(size_t) (t * strata::kernels::kStepCount + strata::kernels::kStepWidth)];
+                            const int32_t *x = a.data() + t * m.cap, *y = b.data() + t * m.cap;
+                            int64_t i = 0, j = 0, c = 0;
+                            while (i < w && j < w) {
+                                if (x[i] == y[j]) { ++c; ++i; ++j; } else if (x[i] < y[j]) ++i; else ++j;
+                            }
+                            ++queries;
+                            same += c == w;
+                            if (p0 + t + 1 > m.cap) { ++sel_queries; shared += (double) c; cells += (double) w; }
+                        }
+                        if (qsa_index + 1 == g.n_qsa_layers())
+                            std::fprintf(stderr, "strata prefill: FP16 indexer keys: %lld of %lld selections identical; "
+                                                 "where the selection is sparse, %.4f%% of cells shared (%lld queries)\n",
+                                         same, queries, cells > 0 ? 100.0 * shared / cells : 100.0, sel_queries);
+                    }
+                    // STRATA_QSA_DUMP=<file>: append every QSA layer's selected cells for the prompt's last
+                    // STRATA_QSA_DUMP_LAST (4096) positions - records of int32 {qsa layer, pos0, T, cap} + T*cap cells,
+                    // for tools/qsa_locality.py (how local the sparse attention's reads are: the KV-streaming question)
+                    if (static const char* dump = std::getenv("STRATA_QSA_DUMP"); dump != nullptr) {
+                        static const long long last = std::getenv("STRATA_QSA_DUMP_LAST")
+                                                          ? std::atoll(std::getenv("STRATA_QSA_DUMP_LAST")) : 4096;
+                        if (p0 + T > pos0 + n - last) {
+                            std::vector<int32_t> h((size_t) (T * m.cap));
+                            cudaMemcpyAsync(h.data(), m.sel_ids, h.size() * 4, cudaMemcpyDeviceToHost, m.cs);
+                            cudaStreamSynchronize(m.cs);
+                            if (std::FILE* f = std::fopen(dump, "ab")) {
+                                const int32_t hdr[4] = {(int32_t) qsa_index, (int32_t) p0, (int32_t) T, (int32_t) m.cap};
+                                std::fwrite(hdr, 4, 4, f);
+                                std::fwrite(h.data(), 4, h.size(), f);
+                                std::fclose(f);
+                            }
+                        }
+                    }
+                    const strata::kernels::QsaAttnPools pools = staged ? pools_of(m.stage, m.ident_table)
+                                                                       : core::qsa_attn_pools(st);
                     for (int64_t t0 = 0; t0 < T; t0 += m.attn_batch) {
                         const int64_t nb = std::min(m.attn_batch, T - t0);
                         strata::kernels::qsa_decode_attn_batch(m.q + t0 * ZV, pools, m.sel_ids + t0 * m.cap,
@@ -533,6 +654,20 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
         return false;
     }
     stats_.ms_total += ms_since(t_start);
+    if (std::getenv("STRATA_STATE_HASH_GDN") != nullptr) {   // debug: the GDN states as the prompt path leaves them
+        cudaStreamSynchronize(m.cs);
+        std::vector<uint8_t> b((size_t) gdn_floats * 4);
+        std::string line;
+        char h[8];
+        for (int64_t i = 0; i < g.n_gdn_layers(); ++i) {
+            cudaMemcpy(b.data(), ss.gdn_state + (size_t) i * gdn_floats, b.size(), cudaMemcpyDeviceToHost);
+            uint64_t x = 1469598103934665603ull;
+            for (uint8_t c : b) x = (x ^ c) * 1099511628211ull;
+            std::snprintf(h, sizeof(h), "%04llx ", (unsigned long long) (x & 0xffff));
+            line += h;
+        }
+        std::fprintf(stderr, "strata prefill: GDN_HASH %s\n", line.c_str());
+    }
     return true;
 }
 

src/program/generate.cpp

@@ -41,6 +41,8 @@
 #include "strata/prefill/prefill.hpp"
 #include "strata/core/native_dense.hpp"
 #include "strata/program/logits_selection.hpp"
+#include "strata/spec/draft_policy.hpp"
+#include "strata/spec/suffix_drafter.hpp"
 
 #include <cuda_runtime.h>
 
@@ -116,6 +118,7 @@ struct Options {
     double ple_delay_us = 0;           ///< fault injection: every row read completes no earlier than this
     bool ple_sync_submit = false;      ///< A/B arm: submit reads on the token thread, no I/O worker
     std::string kv = "fp16";           ///< plan v0.3 P7: KV storage, fp16 (default) or int8 (half the VRAM)
+    int64_t kv_resident = 0;           ///< KV streaming: resident cells per QSA layer (0: all in VRAM)
     std::string dump_residual;
     /// The head input, `bb.mixed`.  It exists so the head can be SPLIT: steps 1-4 (the per-stream norm, the two
     /// bf16 projections and the stream mean) recompute cheaply in Python, and only the 794 MB GEMV does not.
@@ -235,6 +238,26 @@ struct Options {
     /// The vision path: keep a per-cell (t, h, w) rotary position table so --serve can take GENI requests.
     bool vision = false;
     int adapt_swaps = 96;
+    /// --serve: how many conversation checkpoints to keep between requests (0 = every request reads its whole
+    /// prompt again, the v0.1.2 behaviour).  One is the GDN recurrence of the 36 layers, the QSA indexer tails and
+    /// the PLE history (~118 MB of host RAM); the KV cache itself is positional and stays where it is.
+    int prompt_cache = 6;
+    /// --serve: also keep a checkpoint every N freshly read prompt tokens (0 = only at the last turn boundary)
+    int64_t prompt_cache_every = 16384;
+    /// --serve: the token that opens a chat turn (<|im_start|>).  The last one in a prompt is where the chat's
+    /// history ends and the new assistant turn begins, which is the checkpoint the next request can reuse.
+    int64_t turn_token = 248045;
+    /// --serve: a text part of the prompt of at most N tokens (a chat message, the assistant header, a short tool
+    /// result) goes through the verify windows, S tokens at a time, instead of the batched prompt path (0 = always
+    /// the batched path)
+    int64_t short_read = 64;
+    /// The suffix drafter (prompt lookup): when the text being written repeats an earlier stretch of the context (code
+    /// edits, quoted input, tool-call JSON) by at least this many tokens, the window may be filled with what followed
+    /// it there instead of the MTP's drafts, where the MTP's own first guess agrees and the draft policy expects it to
+    /// pay (strata/spec/draft_policy.hpp).  On by default; 0 = MTP only.
+    int suffix_draft = 3;
+    /// The MTP's own window cap (0 = --spec): with --spec 6 --mtp-max-t 4 the long windows come from suffix matches.
+    int mtp_max_t = 0;
 };
 
 void usage() {
@@ -254,6 +277,9 @@ void usage() {
                  "  --ple-sync-submit    A/B arm: submit table reads on the token thread (default: an I/O thread)\n"
                  "  --kv fp16|int8       KV storage (plan v0.3 P7): int8 codes + fp16 scale per 64 values, half the\n"
                  "                       VRAM; default fp16 until gate G-C accepts int8\n"
+                 "  --kv-resident N      KV streaming: keep N cells of each QSA layer in VRAM (min 20480) and the\n"
+                 "                       whole K/V in pinned RAM; the freed VRAM goes to expert slots. 0 (default):\n"
+                 "                       all of it in VRAM. A context of N cells or fewer is not streamed\n"
                  "  --stream-token       enqueue token work on the session stream (experimental)\n"
                  "  --check-logits       copy and check all logits in the stream-token path\n"
                  "  --gr-fp32-activations  experimental CUDA-oracle GR activation precision\n"
@@ -292,6 +318,15 @@ void usage() {
                  "  --shared-late        A/B: shared expert after the CPU pool (default: overlapped with it)\n"
                  "  --keep-canonical     A/B: also load canonical copies of natively served tensors (more VRAM)\n"
                  "  --vision             --serve takes images too (GENI requests; embeddings from strata-vision)\n"
+                 "  --prompt-cache N     --serve: keep N conversation checkpoints between requests (default 6, ~118 MB\n"
+                 "                       of RAM each; 0 = read every prompt from the start)\n"
+                 "  --prompt-cache-every N  --serve: also checkpoint every N fresh prompt tokens (default 16384, 0 = off)\n"
+                 "  --turn-token ID      --serve: the token that opens a chat turn (default 248045, <|im_start|>)\n"
+                 "  --short-read N       --serve: read at most N fresh text tokens through the decode windows instead\n"
+                 "                       of the batched prompt path (default 64, 0 = off)\n"
+                 "  --suffix-draft N     prompt lookup: draft from an earlier repeat of the last N+ tokens of context\n"
+                 "                       when it pays (default 3; 0 = MTP only)\n"
+                 "  --mtp-max-t M        cap the MTP's windows at M tokens (0 = --spec; longer ones come from suffixes)\n"
                  "  --no-token-graph     A/B: two graphs per layer (the host launches each) instead of one per token\n"
                  "  --no-fused-gr        A/B: the six-kernel hyper-connection read and a separate write (native)\n"
                  "  --prefill CHUNK      batched prompt processing in chunks of CHUNK tokens (needs --native)\n"
@@ -442,6 +477,84 @@ int argmax(const std::vector<float>& v) {
     return best;
 }
 
+// ---- --serve's conversation cache.  A chat or an agent sends the whole conversation again with every request, and
+// reading it again is what made a long session wait minutes for every turn.  What a sequence leaves behind splits in
+// two, and only one half needs copying:
+//   * POSITIONAL state - the KV cache of the 12 QSA layers and their pooled indexer keys, the draft layer's KV.  A
+//     cell is written once for its position and read only by later positions (the block scores take `dead` for the
+//     block being filled, never its pooled row), so rewinding to a position just means writing from there again.
+//   * RUNNING state - the 36 GDN recurrences and conv histories, each QSA layer's indexer tail (the unfinished
+//     block's raw keys) and the PLE's normalized history.  Each describes "everything so far" and cannot be
+//     rewound, so a checkpoint is a copy of exactly these: ~118 MB, the same set the verifier snapshots to roll
+//     back rejected drafts.
+// A checkpoint is only valid while the positional cells below it still hold ITS tokens, so the serve loop keeps
+// just the checkpoints that are prefixes of the tokens the session holds now.
+struct ImgKey {
+    int64_t start = 0;      ///< the image's first <|image_pad|> position
+    uint64_t hash = 0;      ///< its embeddings and grid: the pad tokens alone are the same for every picture
+    bool operator==(const ImgKey& o) const { return start == o.start && hash == o.hash; }
+};
+
+struct ConvCheckpoint {
+    std::vector<int32_t> ids;     ///< the tokens this state has consumed
+    std::vector<ImgKey> imgs;     ///< the images among them
+    std::vector<uint8_t> gdn, ple, tails;
+};
+
+uint64_t fnv1a(const void* data, size_t n, uint64_t h = 1469598103934665603ull) {
+    const uint8_t* p = (const uint8_t*) data;
+    for (size_t i = 0; i < n; ++i) { h ^= p[i]; h *= 1099511628211ull; }
+    return h;
+}
+
+struct ConvStateSizes {
+    size_t gdn = 0, ple = 0, tail = 0;
+};
+
+ConvStateSizes conv_state_sizes(const strata::core::ModelGeometry& g) {
+    ConvStateSizes z;
+    z.gdn = (size_t) g.n_gdn_layers() *
+            ((size_t) g.ssm_state_size * (size_t) g.ssm_v_heads * (size_t) g.ssm_state_size +
+             (size_t) g.ssm_conv_channels * (size_t) (g.ssm_d_conv - 1)) * sizeof(float);
+    z.ple = (size_t) strata::kernels::NG_HIST * (size_t) strata::kernels::NG_HC_DIM * sizeof(float);
+    z.tail = (size_t) (strata::kernels::qsa_real_shapes().idx_block - 1) * (size_t) g.idx_key_dim * sizeof(float);
+    return z;
+}
+
+/// Copies the running state out.  The caller has synchronized the device.
+bool checkpoint_save(ConvCheckpoint& c, const strata::core::SessionState& ss, const strata::core::ModelGeometry& g) {
+    const ConvStateSizes z = conv_state_sizes(g);
+    c.gdn.resize(z.gdn);
+    c.ple.resize(ss.ple_hist != nullptr ? z.ple : 0);
+    c.tails.resize(z.tail * (size_t) g.n_qsa_layers());
+    if (cudaMemcpy(c.gdn.data(), ss.gdn_state, z.gdn, cudaMemcpyDeviceToHost) != cudaSuccess) return false;
+    if (!c.ple.empty() && cudaMemcpy(c.ple.data(), ss.ple_hist, z.ple, cudaMemcpyDeviceToHost) != cudaSuccess)
+        return false;
+    for (int64_t i = 0; i < g.n_qsa_layers(); ++i)
+        if (cudaMemcpy(c.tails.data() + (size_t) i * z.tail, ss.qsa_states[i].idx_tail, z.tail, cudaMemcpyDeviceToHost) !=
+            cudaSuccess)
+            return false;
+    return true;
+}
+
+/// Puts a checkpoint's running state back; the positional cells below it are the caller's to guarantee.
+bool checkpoint_restore(const ConvCheckpoint& c, strata::core::SessionState& ss, const strata::core::ModelGeometry& g) {
+    const ConvStateSizes z = conv_state_sizes(g);
+    if (c.gdn.size() != z.gdn || c.tails.size() != z.tail * (size_t) g.n_qsa_layers()) return false;
+    if (cudaMemcpy(ss.gdn_state, c.gdn.data(), z.gdn, cudaMemcpyHostToDevice) != cudaSuccess) return false;
+    if (!c.ple.empty() && cudaMemcpy(ss.ple_hist, c.ple.data(), z.ple, cudaMemcpyHostToDevice) != cudaSuccess)
+        return false;
+    for (int64_t i = 0; i < g.n_qsa_layers(); ++i)
+        if (cudaMemcpy(ss.qsa_states[i].idx_tail, c.tails.data() + (size_t) i * z.tail, z.tail, cudaMemcpyHostToDevice) !=
+            cudaSuccess)
+            return false;
+    // the PLE's token window is the last two tokens, OLDEST FIRST, -1 where there is none (as session_zero leaves it)
+    const size_t L = c.ids.size();
+    ss.ple_prev[0] = L >= 2 ? c.ids[L - 2] : -1;
+    ss.ple_prev[1] = L >= 1 ? c.ids[L - 1] : -1;
+    return cudaDeviceSynchronize() == cudaSuccess;
+}
+
 }  // namespace
 
 int main(int argc, char** argv) {
@@ -506,6 +619,7 @@ int main(int argc, char** argv) {
         else if (a == "--ple-delay-us") o.ple_delay_us = std::atof(next("--ple-delay-us"));
         else if (a == "--ple-sync-submit") o.ple_sync_submit = true;
         else if (a == "--kv") o.kv = next("--kv");
+        else if (a == "--kv-resident") o.kv_resident = std::atoll(next("--kv-resident"));
         else if (a == "--stream-token") o.stream_token = true;
         else if (a == "--check-logits") o.check_logits = true;
         else if (a == "--gr-fp32-activations") o.gr_fp32_activations = true;
@@ -557,6 +671,12 @@ int main(int argc, char** argv) {
         else if (a == "--pcie-mode") o.pcie_mode = next("--pcie-mode");
         else if (a == "--serve") o.serve = true;
         else if (a == "--vision") o.vision = true;
+        else if (a == "--prompt-cache") o.prompt_cache = std::max(0, std::atoi(next("--prompt-cache")));
+        else if (a == "--prompt-cache-every") o.prompt_cache_every = std::max(0LL, std::atoll(next("--prompt-cache-every")));
+        else if (a == "--turn-token") o.turn_token = std::atoll(next("--turn-token"));
+        else if (a == "--short-read") o.short_read = std::max(0LL, std::atoll(next("--short-read")));
+        else if (a == "--suffix-draft") o.suffix_draft = std::max(0, std::atoi(next("--suffix-draft")));
+        else if (a == "--mtp-max-t") o.mtp_max_t = std::max(0, std::atoi(next("--mtp-max-t")));
         else if (a == "--no-spec-split") o.spec_split = false;
         else if (a == "--eos-ids") {
             std::string e;
@@ -610,6 +730,18 @@ int main(int argc, char** argv) {
         return 2;
     }
     strata::core::qsa_set_kv_int8(o.kv == "int8");
+    if (o.kv_resident < 0) {
+        std::fprintf(stderr, "strata generate: --kv-resident must be >= 0\n");
+        return 2;
+    }
+    strata::core::qsa_set_kv_resident(o.kv_resident);
+    // Prompt lookup (the suffix drafter, on by default): the MTP keeps its --spec windows and a lookup window may be
+    // up to 2 tokens longer; the draft policy (strata/spec/draft_policy.hpp) takes one only where it pays. Code
+    // edits +6-11%, ordinary text unchanged (bench/results/2026-09-27-spec). --suffix-draft 0 turns it off.
+    if (o.suffix_draft > 0 && o.spec >= 2 && o.mtp_max_t == 0) {
+        o.mtp_max_t = o.spec;
+        o.spec = std::min(o.spec + 2, 8);   // kVerifyMaxT
+    }
     strata::core::layer_set_shared_early(!o.shared_late);
     if (!o.native_preset.empty()) {
         if (o.no_ple || o.ple_gguf.empty()) {
@@ -840,6 +972,10 @@ int main(int argc, char** argv) {
         std::fprintf(stderr, "strata generate: session_init failed\n");
         return 1;
     }
+    if (g.n_qsa_layers() > 0 && ss.qsa_states[0].kv_mode == 1)
+        std::fprintf(stderr, "strata generate: KV streaming: %lld of %lld cells per QSA layer in VRAM, the K/V in "
+                             "%.2f GiB of pinned RAM\n", (long long) (ss.qsa_states[0].n_slots * 4),
+                     (long long) o.max_context, (double) strata::core::qsa_kv_host_bytes() / 1073741824.0);
 
     // ---- **THE HALF-LEVEL DUMP HAS TO BE ARMED BEFORE `session_capture`, AND THE LADDER MUST NOT BE.**  The
     // half copies are issued from inside `block_layer_pre`/`block_layer_post`, so they are only ever enqueued
@@ -1812,11 +1948,15 @@ int main(int argc, char** argv) {
     //
     // and each generated token is written to stdout as `T <id>` as soon as its verify window is done, followed by
     //
-    //     DONE <generated> <prompt_tokens> <prompt_ms> <decode_ms> <stop|length>
+    //     DONE <generated> <prompt_tokens> <prompt_ms> <decode_ms> <stop|length|cancel> <drafts accepted>
+    //          <drafts offered> <prompt tokens reused>
     //
-    // (`ERR <message>` instead when a request cannot run; `QUIT` ends the process).  Every request starts from an
-    // empty sequence (`session_zero`): the prompt goes through the batched prompt path and its last token through
-    // the first verify window - the path all three model files share.  Decoding is greedy.
+    // Before that, `RESUME <n>` (n prompt tokens are not read again), `PP <position> <prompt_tokens> <ms> <tok/s>`
+    // after every prompt chunk, and `REUSED <n>` once the prompt is read.  (`ERR <message>` instead when a request
+    // cannot run; `STOP` ends the running request at its next step; `QUIT` ends the process.)  A request continues
+    // from the live session or the longest conversation checkpoint its prompt starts with (see ConvCheckpoint),
+    // otherwise from an empty sequence (`session_zero`); the rest of the prompt goes through the batched prompt path
+    // and its last token through the first verify window - the path all three model files share.  Decoding is greedy.
     if (o.serve) {
         if (o.spec < 2 || o.mtp.empty() || o.prefill_chunk <= 0 || thits.d_res == nullptr || host_res.empty()) {
             std::fprintf(stderr, "strata serve: needs --spec T, --mtp DIR, --prefill CHUNK, --expert-profile P and "
@@ -1861,6 +2001,16 @@ int main(int argc, char** argv) {
             return 1;
         }
         mem_mark("the head and the prompt path");
+        // the penalty-history buffer: ONE row, the request's last `penalty_last_n` tokens, -1 padded in
+        // front.  Allocated once at the cap; a request without penalties gets a null row and takes the
+        // byte-for-byte neutral path (no upload, no buffer handed to the sampler).
+        constexpr int kPenaltyWindowCap = 4096;
+        int32_t* d_hist = nullptr;
+        std::vector<int32_t> hist_stage((size_t) kPenaltyWindowCap, -1);
+        if (cudaMalloc(&d_hist, (size_t) kPenaltyWindowCap * sizeof(int32_t)) != cudaSuccess) {
+            std::fprintf(stderr, "strata serve: the penalty-history allocation failed\n");
+            return 1;
+        }
         strata::core::Verifier ver;
         strata::core::VerifyHits vh;
         vh.d_res = thits.d_res;
@@ -1876,10 +2026,47 @@ int main(int argc, char** argv) {
         ver.set_pcie_mode(o.pcie_mode == "dma" ? 0 : o.pcie_mode == "direct" ? 1 : o.pcie_mode == "kernel" ? 2
                           : native_pack ? 0 : 2);   // auto: DMA for the native packs, the copy kernel for Q2_0
         std::vector<int64_t> cur;
+        // ---- the conversation cache (see ConvCheckpoint).  `live` is what the session holds right now: the tokens
+        // it has consumed, so a request that starts with exactly them continues without any copy.  `checks` are the
+        // saved points; every one of them is a prefix of `live` (the loop drops the rest).
+        std::vector<int32_t> live;
+        std::vector<ImgKey> live_imgs, req_imgs;
+        bool live_ok = false;
+        std::vector<ConvCheckpoint> checks;
+        int64_t pp_total = 0, pp_from = 0, pp_next_check = 0;
+        Clock::time_point pp_t0 = Clock::now();
+        auto imgs_below = [&](const std::vector<ImgKey>& all, int64_t L) {
+            std::vector<ImgKey> v;
+            for (const ImgKey& k : all) if (k.start < L) v.push_back(k);
+            return v;
+        };
+        // a checkpoint of the state after `cur[0, L)`; false only when the copy itself failed
+        auto checkpoint_at = [&](int64_t L) -> bool {
+            if (o.prompt_cache <= 0 || L < 1) return true;
+            for (const ConvCheckpoint& c : checks) if ((int64_t) c.ids.size() == L) return true;
+            ConvCheckpoint c;
+            c.ids.assign(cur.begin(), cur.begin() + L);
+            c.imgs = imgs_below(req_imgs, L);
+            if (cudaDeviceSynchronize() != cudaSuccess || !checkpoint_save(c, ss, g)) return false;
+            checks.push_back(std::move(c));
+            while ((int) checks.size() > o.prompt_cache) checks.erase(checks.begin());   // the oldest goes first
+            return true;
+        };
         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)];
-            return mtp.prefill(R_rows, nxt.data(), T, p0, e);
+            if (!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();
+            std::printf("PP %lld %lld %.0f %.1f\n", (long long) done, (long long) pp_total, ms,
+                        ms > 0.0 ? 1000.0 * (double) (done - pp_from) / ms : 0.0);
+            std::fflush(stdout);
+            if (o.prompt_cache_every > 0 && done >= pp_next_check) {
+                if (!checkpoint_at(done)) { e = "saving a conversation checkpoint failed"; return false; }
+                pp_next_check = done + o.prompt_cache_every;
+            }
+            return true;
         };
         drive.d.plan = ver.plan_sink();
         drive.d.pcie_num = std::max(0, std::min(256, (int) (o.pcie_frac * 256.0 + 0.5)));
@@ -1994,6 +2181,10 @@ int main(int argc, char** argv) {
         std::string line;
         int64_t rounds = 0;
         const int S = o.spec;
+        const int S_mtp = o.mtp_max_t > 0 ? std::min(o.mtp_max_t, S) : S;   // the MTP's windows; suffixes go up to S
+        if (S_mtp < S) mtp.set_max_drafts(S_mtp - 1);
+        strata::spec::SuffixDrafter sfx(std::max(1, o.suffix_draft), 64, (size_t) o.max_context + 4096);
+        strata::spec::DraftPolicy policy(S);   // MTP or lookup window, learned over the whole process
         // The vision path (--vision): GENI <max_new> <embeddings file> <id,id,...> carries images.  The file is one
         // or more strata-vision records (int32 'SVE1', n, nx, ny, n_embd, then n x n_embd floats) in prompt order;
         // each image's rows go to its run of <|image_pad|> tokens, whose M-RoPE positions are mtmd's: t = p,
@@ -2012,6 +2203,37 @@ int main(int argc, char** argv) {
             }
             char* endp = nullptr;
             const long long max_new = std::strtoll(line.c_str() + (geni ? 5 : 4), &endp, 10);
+            // optional sampling keys between max_new and the ids: temperature=F, top_p=F, top_k=N, min_p=F,
+            // penalty_last_n=N, penalty_repeat=F, penalty_freq=F, penalty_present=F, seed=N (text requests
+            // only).  Absent keys keep today's behavior: greedy, no penalties.
+            float req_temperature = 0.0f, req_top_p = 1.0f;
+            int req_top_k = 20;   // the sampler's own default; the sampled path REQUIRES top_k in 1..64
+            unsigned long long req_seed = 0;
+            float req_min_p = 0.0f, req_penalty_repeat = 1.0f, req_penalty_freq = 0.0f, req_penalty_present = 0.0f;
+            int req_penalty_last_n = 0;
+            if (!geni && endp != nullptr) {
+                for (;;) {
+                    while (*endp == ' ') ++endp;
+                    const char* start = endp;
+                    while (*endp != '\0' && *endp != ' ') ++endp;
+                    if (endp == start) break;
+                    const std::string tok(start, (size_t) (endp - start));
+                    const size_t eq = tok.find('=');
+                    if (eq == std::string::npos) { endp = const_cast<char*>(start); break; }
+                    const std::string key = tok.substr(0, eq);
+                    const float fv = std::strtof(tok.c_str() + eq + 1, nullptr);
+                    if (key == "temperature") req_temperature = fv;
+                    else if (key == "top_p") req_top_p = fv;
+                    else if (key == "top_k") req_top_k = std::atoi(tok.c_str() + eq + 1);
+                    else if (key == "min_p") req_min_p = fv;
+                    else if (key == "penalty_last_n") req_penalty_last_n = std::atoi(tok.c_str() + eq + 1);
+                    else if (key == "penalty_repeat") req_penalty_repeat = fv;
+                    else if (key == "penalty_freq") req_penalty_freq = fv;
+                    else if (key == "penalty_present") req_penalty_present = fv;
+                    else if (key == "seed") req_seed = std::strtoull(tok.c_str() + eq + 1, nullptr, 10);
+                    // unknown keys are skipped: the ids start at the first token without '='
+                }
+            }
             std::string emb_path;
             if (geni && endp != nullptr) {
                 while (*endp == ' ') ++endp;
@@ -2025,6 +2247,7 @@ int main(int argc, char** argv) {
                 continue;
             }
             const int64_t n = (int64_t) ids.size();
+            req_imgs.clear();
             if (geni && !o.vision) { std::printf("ERR this engine was started without --vision\n"); continue; }
             if (geni || !mrope_identity) {
                 // positions for every cell this request can reach; the identity again for a text request
@@ -2066,6 +2289,12 @@ int main(int argc, char** argv) {
                         if (ids[(size_t) i] != kImagePad) { put(i, p, p, p); ++p; ++i; continue; }
                         if (k >= imgs.size()) { ve = "the prompt has more images than the embeddings file"; break; }
                         const Img& im = imgs[k++];
+                        {   // what the conversation cache compares: a picture is its grid and its embeddings
+                            const int64_t grid[3] = {im.n, im.nx, im.ny};
+                            uint64_t h = fnv1a(grid, sizeof grid);
+                            h = fnv1a(img_rows.data() + im.off, (size_t) im.n * (size_t) g.n_embd * sizeof(float), h);
+                            req_imgs.push_back({i, h});
+                        }
                         for (int64_t j = 0; j < im.n && ve.empty(); ++j)
                             if (i + j >= n || ids[(size_t) (i + j)] != kImagePad)
                                 ve = "image " + std::to_string(k) + " has " + std::to_string((long long) im.n) +
@@ -2110,67 +2339,264 @@ int main(int argc, char** argv) {
             if (bad) { std::printf("ERR a token id is outside the vocabulary\n"); continue; }
             cur = ids;
             const Clock::time_point r0 = Clock::now();
-            strata::core::session_zero(ss, g, nullptr, main_cs);
-            cudaStreamSynchronize(main_stream);
+            // ---- where this request starts reading: the live session, or a checkpoint, whose tokens AND pictures are
+            // exactly the start of this prompt - at most n - 1 of them, the last token is always the first window
+            auto starts_with = [&](const std::vector<int32_t>& pre, const std::vector<ImgKey>& pre_imgs) -> bool {
+                const int64_t L = (int64_t) pre.size();
+                if (L < 1 || L > n - 1) return false;
+                for (int64_t i = 0; i < L; ++i)
+                    if ((int32_t) ids[(size_t) i] != pre[(size_t) i]) return false;
+                return imgs_below(req_imgs, L) == pre_imgs;
+            };
+            int64_t resume = 0;
+            bool from_live = false;
+            if (o.prompt_cache > 0) {
+                if (live_ok && starts_with(live, live_imgs)) { resume = (int64_t) live.size(); from_live = true; }
+                for (const ConvCheckpoint& c : checks)
+                    if ((int64_t) c.ids.size() > resume && starts_with(c.ids, c.imgs)) {
+                        resume = (int64_t) c.ids.size();
+                        from_live = false;
+                    }
+            }
+            // this request rewrites every cell from `resume` on, so a checkpoint past it (or not on this prompt's
+            // path) no longer has its cells; the ones kept are prefixes of both the old tokens and the new
+            checks.erase(std::remove_if(checks.begin(), checks.end(), [&](const ConvCheckpoint& c) {
+                             return (int64_t) c.ids.size() > resume || !starts_with(c.ids, c.imgs);
+                         }), checks.end());
+            live_ok = false;   // until this request has finished, the session is in between
+            int64_t reread_to = -1;   // STRATA_CKPT_REREAD only: read [0, reread_to) again instead of restoring
+            if (resume == 0) {
+                strata::core::session_zero(ss, g, nullptr, main_cs);
+                cudaStreamSynchronize(main_stream);
+                checks.clear();
+            } else if (!from_live) {
+                const ConvCheckpoint* c = nullptr;
+                for (const ConvCheckpoint& k : checks) if ((int64_t) k.ids.size() == resume) c = &k;
+                static const bool reread = std::getenv("STRATA_CKPT_REREAD") != nullptr;
+                if (reread && c != nullptr) {
+                    // THE CHECK OF THE CHECKPOINT: instead of restoring it, read its tokens again from position 0 in
+                    // one run (below, with the prompt path's slots lent like any read) - the same chunks the request
+                    // that saved it read them in, when that request started at 0.  With the VRAM expert set fixed
+                    // (--adapt-swaps 0) the answer must match the restored one token for token; anything the
+                    // checkpoint missed shows up as a difference.
+                    strata::core::session_zero(ss, g, nullptr, main_cs);
+                    cudaStreamSynchronize(main_stream);
+                    reread_to = resume;
+                    std::fprintf(stderr, "strata serve: STRATA_CKPT_REREAD: reading %lld tokens again instead of "
+                                         "restoring\n", (long long) resume);
+                } else if (c == nullptr || !checkpoint_restore(*c, ss, g)) {
+                    std::printf("ERR restoring a conversation checkpoint failed\n");
+                    return 1;
+                }
+            }
+            // KV streaming: the drafter's ring may hold cells past `resume` from a longer turn; the main layers'
+            // host copies and slots are always current (every writer writes both), so they need nothing
+            if (resume > 0 && reread_to <= 0) mtp.kv_restore(resume);
             tr("request", n, geni ? 1 : 0);
             mtp.set_prompt_len(n);
+            const int64_t read_from = reread_to > 0 ? 0 : resume;
+            pp_total = n;
+            pp_from = read_from;
+            pp_t0 = r0;
+            pp_next_check = reread_to > 0 ? INT64_MAX : resume + o.prompt_cache_every;
+            std::printf("RESUME %lld\n", (long long) resume);   // before reading: this many prompt tokens are reused
+            std::fflush(stdout);
+            // A SHORT PART OF THE PROMPT - the new message of a chat that continues from a checkpoint, the assistant
+            // header - goes through the verify windows, S tokens at a time, as decode reads them.  The batched path
+            // costs ~300 ms per run however few tokens it has (it streams every expert the chunk routes to that is
+            // not in VRAM over PCIe), and it borrows slots it must refill after (~180 ms); a window costs ~16 ms a
+            // token, with the misses on the CPU.  Each part below is decided on its own, so a long first message is
+            // read batched and its header still goes through the windows.  Picture rows need the batched path.
+            // STRATA_CKPT_REREAD compares a restored checkpoint with a batched re-read, so it keeps every read batched.
+            static const bool no_short = std::getenv("STRATA_CKPT_REREAD") != nullptr;
+            auto windows_ok = [&](int64_t a, int64_t b) -> bool {
+                if (no_short || b - a > o.short_read) return false;
+                if (sp.embd_rows != nullptr)
+                    for (int64_t i = a; i < b; ++i)
+                        if (sp.embd_rows[i] != nullptr) return false;
+                return true;
+            };
+            // 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 {
+                std::vector<int32_t> win((size_t) S), outw((size_t) S), nxt((size_t) S);
+                for (int64_t q = a; q < b;) {
+                    if (stop_req.load()) { e = "cancelled"; return false; }
+                    const int T = (int) std::min<int64_t>(S, b - q);
+                    for (int t = 0; t < T; ++t) {
+                        win[(size_t) t] = (int32_t) cur[(size_t) (q + t)];
+                        nxt[(size_t) t] = (int32_t) cur[(size_t) (q + t + 1)];
+                    }
+                    drive.d.layers = 0;
+                    drive.d.experts = 0;
+                    drive.d.failed = false;
+                    if (!ver.run(T, win.data(), q, &drive_pool_multi, &drive, outw.data(), e) || drive.d.failed) {
+                        if (drive.d.failed && drive.d.fail) e = drive.d.fail;
+                        return false;
+                    }
+                    if (!ver.commit(T, e) || !mtp.prefill(ver.final_R_all(), nxt.data(), T, q, e)) return false;
+                    q += T;
+                }
+                const double ms = std::chrono::duration<double, std::milli>(Clock::now() - pp_t0).count();
+                std::printf("PP %lld %lld %.0f %.1f\n", (long long) b, (long long) pp_total, ms,
+                            ms > 0.0 ? 1000.0 * (double) (b - pp_from) / ms : 0.0);
+                std::fflush(stdout);
+                return true;
+            };
+            // the batched path's slots are lent just before its first run and given back (refilled) before a window
+            // reads - so the windows always see the whole expert cache - or once the prompt is read
             std::vector<std::pair<int32_t, int32_t>> lent_now;
-            apply_pending(true);
-            if (lend_first >= 0) {
+            auto lend = [&] {
+                if (lend_first < 0 || !lent_now.empty()) return;
                 for (size_t i = 0; i < host_res.size(); ++i)
                     if (host_res[i] >= lend_first) {
                         lent_now.emplace_back((int32_t) i, host_res[i]);
                         host_res[i] = strata::core::kNotResident;
                     }
                 cudaMemcpy(d_res, host_res.data(), host_res.size() * sizeof(int32_t), cudaMemcpyHostToDevice);
-            }
+            };
+            auto refill = [&](std::string& e) -> bool {
+                if (lent_now.empty()) return true;
+                tr("refill start", (long long) lent_now.size());
+                for (const auto& [i, slot] : lent_now) {
+                    const uint8_t* b = srcp->blob(i / g.n_expert, i % g.n_expert);
+                    if (b == nullptr || !xcache.fill_slot_blocking(slot, b, e,
+                            (int64_t) strata::kernels::cpu::expert_layout().blob_bytes(i / g.n_expert)))
+                        return false;
+                    host_res[(size_t) i] = slot;
+                }
+                cudaMemcpy(d_res, host_res.data(), host_res.size() * sizeof(int32_t), cudaMemcpyHostToDevice);
+                lent_now.clear();
+                return true;
+            };
+            apply_pending(true);
+            // per-request sampling for the verify window's head (greedy when temperature is absent)
+            strata::kernels::SamplerParams req_sp;
+            req_sp.greedy = req_temperature <= 0.0f;
+            req_sp.temperature = req_temperature;
+            req_sp.top_p = req_top_p;
+            req_sp.top_k = req_top_k;
+            req_sp.seed = req_seed ? req_seed
+                                   : (unsigned long long) std::chrono::steady_clock::now().time_since_epoch().count();
+            req_sp.min_p = std::clamp(req_min_p, 0.0f, 1.0f);
+            req_sp.penalty_last_n = std::max(req_penalty_last_n, 0);
+            req_sp.penalty_repeat = req_penalty_repeat;
+            req_sp.penalty_freq = req_penalty_freq;
+            req_sp.penalty_present = req_penalty_present;
+            req_sp.counter = 0;
+            ver.set_sampling(req_sp);
+            const int hist_n = std::min(req_sp.penalty_last_n, kPenaltyWindowCap);
+            ver.set_history(hist_n > 0 ? d_hist : nullptr, hist_n);
             bool cancelled = false;
             tr("prompt start", n - 1);
-            if (n > 1 && !sp.run(ids.data(), n - 1, 0, err)) {
-                if (!stop_req.load()) {
-                    std::fprintf(stderr, "strata serve: %s\n", err.c_str());
-                    std::printf("ERR %s\n", err.c_str());
-                    return 1;
-                }
-                cancelled = true;   // stopped while reading the prompt: refill the lent slots below, then DONE cancel
-            }
-            for (const auto& [i, slot] : lent_now) {
-                const uint8_t* b = srcp->blob(i / g.n_expert, i % g.n_expert);
-                if (b == nullptr || !xcache.fill_slot_blocking(slot, b, err,
-                        (int64_t) strata::kernels::cpu::expert_layout().blob_bytes(i / g.n_expert))) {
+            // The prompt is read in two parts when it has a turn boundary past `resume`: up to the last <|im_start|>
+            // (the conversation so far), a checkpoint there, then the new turn's header.  The next request of the same
+            // chat renders the same history - but not always the same header or the thinking of this reply - so that
+            // checkpoint is the one it reuses.
+            int64_t turn_at = -1;
+            if (o.prompt_cache > 0 && o.turn_token >= 0)
+                for (int64_t i = n - 1; i > resume; --i)
+                    if (ids[(size_t) i] == o.turn_token) { turn_at = i; break; }
+            int64_t at = read_from;
+            for (const int64_t to : {reread_to, turn_at, n - 1}) {
+                if (to <= at) continue;
+                const bool win = windows_ok(at, to);
+                if (win && !refill(err)) {
                     std::printf("ERR refilling a lent slot failed: %s\n", err.c_str());
                     return 1;
                 }
-                host_res[(size_t) i] = slot;
+                if (!win) lend();
+                const auto tsp = Clock::now();
+                const bool sp_ok = win ? read_windows(at, to, err) : sp.run(ids.data() + at, to - at, at, err);
+                if (trace) {
+                    std::fprintf(stderr, "strata trace: read %lld tokens (%s) in %.1f ms\n", (long long) (to - at),
+                                 win ? "windows" : "batched",
+                                 std::chrono::duration<double, std::milli>(Clock::now() - tsp).count());
+                    std::fflush(stderr);
+                }
+                if (!sp_ok) {
+                    if (!stop_req.load()) {
+                        std::fprintf(stderr, "strata serve: %s\n", err.c_str());
+                        std::printf("ERR %s\n", err.c_str());
+                        return 1;
+                    }
+                    cancelled = true;   // stopped while reading the prompt: refill the lent slots below, then DONE cancel
+                    break;
+                }
+                at = to;
+                if (to == turn_at && !checkpoint_at(to)) {
+                    std::printf("ERR saving a conversation checkpoint failed\n");
+                    return 1;
+                }
+            }
+            if (!refill(err)) {
+                std::printf("ERR refilling a lent slot failed: %s\n", err.c_str());
+                return 1;
             }
-            if (!lent_now.empty())
-                cudaMemcpy(d_res, host_res.data(), host_res.size() * sizeof(int32_t), cudaMemcpyHostToDevice);
             tr("prompt done (slots refilled)");
             const double prompt_ms = std::chrono::duration<double, std::milli>(Clock::now() - r0).count();
+            std::printf("REUSED %lld\n", (long long) resume);   // the prompt is read; the first window comes next
+            std::fflush(stdout);
             // the verify windows: the first holds the last prompt token alone
             int64_t p = n - 1;
             int32_t x = (int32_t) ids[(size_t) (n - 1)];
             std::vector<int32_t> drafts((size_t) S, 0), window((size_t) S), outv((size_t) S);
             std::vector<float> dprob((size_t) S, 0.0f);
+            std::vector<int32_t> sbuf((size_t) S, 0);
+            if (o.suffix_draft > 0) {
+                sfx.reset();
+                for (int64_t t : ids) sfx.append((int32_t) t);
+            }
             bool first_window = true;
-            int64_t produced_n = 0;
+            int64_t produced_n = 0, sfx_windows = 0, sfx_drafts = 0, sfx_ok = 0;
+            int64_t draft_offered = 0, draft_accepted = 0;
+            // what the session holds once this request is done: the prompt read so far, then every committed token
+            std::vector<int32_t> consumed;
+            consumed.reserve((size_t) (n + max_new + S));
+            for (int64_t i = 0; i < n - 1; ++i) consumed.push_back((int32_t) ids[(size_t) i]);
             const char* finish = "length";
             const Clock::time_point d0 = Clock::now();
             if (cancelled) finish = "cancel";
             while (!cancelled && produced_n < max_new) {
-                int T = S;
+                int T = S_mtp;
                 if (o.spec_min_p > 0.0) {
                     T = 1;
-                    while (T < S && dprob[(size_t) T - 1] >= (float) o.spec_min_p) ++T;
+                    while (T < S_mtp && dprob[(size_t) T - 1] >= (float) o.spec_min_p) ++T;
                 }
                 if (first_window) T = 1;
+                // a repeat of earlier context (prompt lookup) where the MTP's own first guess agrees: the policy takes it
+                // when its expected tokens per ms, from the measured acceptance and window costs, beat the MTP window's
+                bool from_sfx = false;
+                int sfx_match = 0;
+                if (o.suffix_draft > 0 && !first_window) {
+                    const int k = sfx.propose(S - 1, sbuf.data());
+                    sfx_match = sfx.last_match();
+                    if (k > 0 && sbuf[0] == drafts[0]) {
+                        const strata::spec::DraftPolicy::Pick pk = policy.choose(T, k, sfx_match);
+                        if (pk.lookup) { T = pk.t; from_sfx = true; }
+                    }
+                }
+                const bool timed_round = !first_window;
+                const Clock::time_point round0 = Clock::now();
                 if (p + T > o.max_context) break;
                 window[0] = x;
-                for (int i = 1; i < T; ++i) window[(size_t) i] = drafts[(size_t) i - 1];
+                for (int i = 1; i < T; ++i) window[(size_t) i] = from_sfx ? sbuf[(size_t) i - 1] : drafts[(size_t) i - 1];
                 drive.d.layers = 0;
                 drive.d.experts = 0;
                 drive.d.failed = false;
                 apply_pending(false);
+                if (hist_n > 0) {
+                    // the tail the penalties count over: the tokens the state has consumed plus the fed-back
+                    // head `x` (it joins `consumed` only after this window commits).  Most recent LAST,
+                    // -1 in the unused front slots - the sampler's row layout.
+                    const int64_t avail = (int64_t) consumed.size() + 1;
+                    const int take = (int) std::min<int64_t>(hist_n, avail);
+                    std::fill(hist_stage.begin(), hist_stage.begin() + hist_n, -1);
+                    for (int j = 0; j < take - 1; ++j)
+                        hist_stage[(size_t) (hist_n - take + j)] =
+                            consumed[(size_t) ((int64_t) consumed.size() - (take - 1) + j)];
+                    hist_stage[(size_t) (hist_n - 1)] = (int32_t) x;
+                    cudaMemcpy(d_hist, hist_stage.data(), (size_t) hist_n * sizeof(int32_t), cudaMemcpyHostToDevice);
+                }
                 tr("window", p, T);
                 if (!ver.run(T, window.data(), p, &drive_pool_multi, &drive, outv.data(), err) || drive.d.failed) {
                     std::printf("ERR %s\n", drive.d.failed && drive.d.fail ? drive.d.fail : err.c_str());
@@ -2178,6 +2604,7 @@ int main(int argc, char** argv) {
                 }
                 int a = 0;
                 while (a < T - 1 && window[(size_t) a + 1] == outv[(size_t) a]) ++a;
+                if (from_sfx) { ++sfx_windows; sfx_drafts += T - 1; sfx_ok += a; }
                 std::thread adapt_thr;   // the adaptive tier beside the commit and the draft (as in generate)
                 bool adapt_ok = true;
                 if (!drive.d.usage.empty() && ((rounds + 1) % o.adapt_every) == 0)
@@ -2187,11 +2614,16 @@ int main(int argc, char** argv) {
                     std::printf("ERR %s\n", err.c_str());
                     return 1;
                 }
+                // the window's first a + 1 tokens are in the session now (the last output is not: it is next x)
+                for (int i = 0; i <= a; ++i) consumed.push_back(window[(size_t) i]);
+                draft_offered += T - 1;
+                draft_accepted += a;
                 first_window = false;
                 bool eos = false;
                 for (int i = 0; i <= a && produced_n < max_new && !eos; ++i) {
                     std::printf("T %d\n", (int) outv[(size_t) i]);
                     ++produced_n;
+                    if (o.suffix_draft > 0) sfx.append(outv[(size_t) i]);
                     eos = std::find(o.eos_ids.begin(), o.eos_ids.end(), (int64_t) outv[(size_t) i]) != o.eos_ids.end();
                 }
                 std::fflush(stdout);
@@ -2207,18 +2639,122 @@ int main(int argc, char** argv) {
                     std::printf("ERR %s\n", err.c_str());
                     return 1;
                 }
+                if (timed_round && !eos)
+                    policy.observe(from_sfx, T, a, sfx_match,
+                                   std::chrono::duration<double, std::milli>(Clock::now() - round0).count());
                 if (eos) { finish = "stop"; break; }
                 if (stop_req.load()) { finish = "cancel"; break; }
                 x = outv[(size_t) a];
                 p += a + 1;
             }
             const double decode_ms = std::chrono::duration<double, std::milli>(Clock::now() - d0).count();
-            std::printf("DONE %lld %lld %.1f %.1f %s\n", (long long) produced_n, (long long) n, prompt_ms, decode_ms,
-                        finish);
+            if (!cancelled) {
+                // a prompt stopped halfway leaves the session somewhere between two chunks: nothing to continue from
+                // (the checkpoints taken while reading it are still good)
+                live.swap(consumed);
+                live_imgs = imgs_below(req_imgs, (int64_t) live.size());
+                live_ok = o.prompt_cache > 0;
+            }
+            static const bool state_hash = std::getenv("STRATA_STATE_HASH") != nullptr;
+            if (state_hash && live_ok) {
+                // DEBUG: a fingerprint of every part of the session over the positions it holds ([0, L)), and
+                // separately of what lies past them in the last KV page (stale cells, fine unless something reads them)
+                cudaDeviceSynchronize();
+                const int64_t L = (int64_t) live.size();
+                const strata::kernels::QsaShapes qs = [&] {
+                    strata::kernels::QsaShapes s = strata::kernels::qsa_real_shapes();
+                    s.n_head_kv = g.n_head_kv; s.head_dim = g.head_dim; s.idx_dim = g.idx_key_dim;
+                    return s;
+                }();
+                auto hash_dev = [&](const void* p, size_t bytes, uint64_t h) {
+                    std::vector<uint8_t> b(bytes);
+                    if (bytes) cudaMemcpy(b.data(), p, bytes, cudaMemcpyDefault);   // VRAM or a streamed host copy
+                    return fnv1a(b.data(), b.size(), h);
+                };
+                // the cells [c0, c1) of one int8 K or V pool ([page][kv_head][page_size][head_dim]), bytes per value `w`
+                auto hash_cells = [&](const void* pool, int64_t per_cell, int64_t c0, int64_t c1, uint64_t h) {
+                    const int64_t ps = qs.page_size;
+                    for (int64_t pg = c0 / ps; pg * ps < c1; ++pg)
+                        for (int64_t hd = 0; hd < qs.n_head_kv; ++hd) {
+                            const int64_t a = std::max(c0, pg * ps) - pg * ps, e = std::min(c1, (pg + 1) * ps) - pg * ps;
+                            const size_t off = (size_t) (((pg * qs.n_head_kv + hd) * ps + a) * per_cell);
+                            h = hash_dev((const uint8_t*) pool + off, (size_t) ((e - a) * per_cell), h);
+                        }
+                    return h;
+                };
+                const ConvStateSizes z = conv_state_sizes(g);
+                uint64_t h_gdn = hash_dev(ss.gdn_state, z.gdn, 1469598103934665603ull);
+                if (std::getenv("STRATA_STATE_HASH_GDN") != nullptr) {   // per GDN layer: which one differs first
+                    const size_t per = z.gdn / (size_t) g.n_gdn_layers();
+                    std::string s;
+                    char b[8];
+                    for (int64_t i = 0; i < g.n_gdn_layers(); ++i) {
+                        std::snprintf(b, sizeof(b), "%04llx ", (unsigned long long) (hash_dev((const uint8_t*) ss.gdn_state + i * per, per, 1469598103934665603ull) & 0xffff));
+                        s += b;
+                    }
+                    std::fprintf(stderr, "strata serve: STATE_HASH_GDN %s\n", s.c_str());
+                }
+                uint64_t h_ple = hash_dev(ss.ple_hist, z.ple, 1469598103934665603ull);
+                uint64_t h_tail = 1469598103934665603ull, h_pool = h_tail, h_kv = h_tail, h_stale = h_tail;
+                const int64_t kvb = qs.head_dim, scb = (qs.head_dim / 64) * 2;
+                const int64_t end_cell = std::min<int64_t>(((L + qs.page_size - 1) / qs.page_size) * qs.page_size,
+                                                           ss.qsa_states[0].max_cells);
+                for (int64_t i = 0; i < g.n_qsa_layers(); ++i) {
+                    const strata::core::QsaState& st = ss.qsa_states[i];
+                    h_tail = hash_dev(st.idx_tail, z.tail, h_tail);
+                    h_pool = hash_dev(st.idx_pooled, (size_t) (L / qs.idx_block) * qs.idx_dim * 4, h_pool);
+                    // KV streaming: the host copy is the identity layout and holds every cell
+                    const bool hs = st.kv_mode != 0;
+                    for (const auto& [pool, w] : {std::pair<const void*, int64_t>{hs ? st.host.k_q : st.k_q, kvb},
+                                                  {hs ? st.host.v_q : st.v_q, kvb},
+                                                  {hs ? st.host.k_scale : st.k_scale, scb},
+                                                  {hs ? st.host.v_scale : st.v_scale, scb}}) {
+                        h_kv = hash_cells(pool, w, 0, L, h_kv);
+                        h_stale = hash_cells(pool, w, L, end_cell, h_stale);
+                    }
+                }
+                const strata::core::QsaState& ms = mtp.kv_state();
+                uint64_t h_mtp = 1469598103934665603ull;
+                const int64_t mL = std::min<int64_t>(L, ms.max_cells);
+                const bool mh = ms.kv_mode != 0;
+                for (const auto& [pool, w] : {std::pair<const void*, int64_t>{mh ? ms.host.k_q : ms.k_q, kvb},
+                                              {mh ? ms.host.v_q : ms.v_q, kvb},
+                                              {mh ? ms.host.k_scale : ms.k_scale, scb},
+                                              {mh ? ms.host.v_scale : ms.v_scale, scb}})
+                    if (pool != nullptr) h_mtp = hash_cells(pool, w, 0, mL, h_mtp);
+                std::fprintf(stderr, "strata serve: STATE_HASH L=%lld gdn=%016llx ple=%016llx tail=%016llx pooled=%016llx "
+                                     "kv=%016llx mtp=%016llx stale=%016llx ple_prev=%d,%d\n", (long long) L,
+                             (unsigned long long) h_gdn, (unsigned long long) h_ple, (unsigned long long) h_tail,
+                             (unsigned long long) h_pool, (unsigned long long) h_kv, (unsigned long long) h_mtp,
+                             (unsigned long long) h_stale, ss.ple_prev[0], ss.ple_prev[1]);
+            }
+            // DONE <generated> <prompt> <prompt ms> <decode ms> <finish> <drafts accepted> <drafts offered> <reused>
+            std::printf("DONE %lld %lld %.1f %.1f %s %lld %lld %lld\n", (long long) produced_n, (long long) n, prompt_ms,
+                        decode_ms, finish, (long long) draft_accepted, (long long) draft_offered, (long long) resume);
             std::fflush(stdout);
-            std::fprintf(stderr, "strata serve: %lld prompt tokens in %.0f ms (%.1f tok/s), %lld generated in %.0f ms "
-                                 "(%.1f tok/s)\n", (long long) n, prompt_ms, prompt_ms > 0 ? 1000.0 * n / prompt_ms : 0.0,
-                         (long long) produced_n, decode_ms, decode_ms > 0 ? 1000.0 * produced_n / decode_ms : 0.0);
+            const int64_t fresh = n - resume;
+            std::fprintf(stderr, "strata serve: prompt %lld tokens = %lld reused + %lld read in %.0f ms (%.1f tok/s), "
+                                 "%lld generated in %.0f ms (%.1f tok/s), drafts accepted %lld of %lld, %zu checkpoints%s\n",
+                         (long long) n, (long long) resume, (long long) fresh, prompt_ms,
+                         prompt_ms > 0 ? 1000.0 * fresh / prompt_ms : 0.0, (long long) produced_n, decode_ms,
+                         decode_ms > 0 ? 1000.0 * produced_n / decode_ms : 0.0, (long long) draft_accepted,
+                         (long long) draft_offered, checks.size(), cancelled ? " (cancelled)" : "");
+            if (g.n_qsa_layers() > 0 && ss.qsa_states[0].kv_mode == 1) {
+                // KV streaming, cumulative over the process: blocks the selections named vs blocks read from RAM
+                uint64_t miss = 0, look = 0;
+                bool over = false;
+                for (int64_t i = 0; i < g.n_qsa_layers(); ++i) {
+                    const strata::kernels::KvStreamCounters c = strata::kernels::kv_stream_counters(ss.qsa_states[i].map);
+                    miss += c.misses; look += c.lookups; over = over || c.overflow;
+                }
+                std::fprintf(stderr, "strata serve: KV streaming: %.2f%% of %llu block reads hit VRAM, %.1f MiB read "
+                                     "from RAM%s\n", look ? 100.0 * (double) (look - miss) / (double) look : 100.0,
+                             (unsigned long long) look, (double) miss * 4224.0 / 1048576.0,
+                             over ? " - OVERFLOW (too few resident cells)" : "");
+            }
+            if (sfx_windows > 0)
+                std::fprintf(stderr, "strata serve: suffix drafts: %lld windows, %lld of %lld drafts accepted\n",
+                             (long long) sfx_windows, (long long) sfx_ok, (long long) sfx_drafts);
         }
         return 0;
     }
@@ -2663,16 +3199,37 @@ int main(int argc, char** argv) {
         std::vector<int32_t> window((size_t) o.spec), outv((size_t) o.spec);
         std::vector<int64_t> accepted_hist((size_t) o.spec, 0);
         int64_t rounds = 0, drafts_total = 0, drafts_ok = 0, corrupt_counter = 0;
+        const int S_mtp = o.mtp_max_t > 0 ? std::min(o.mtp_max_t, o.spec) : o.spec;
+        if (use_mtp && S_mtp < o.spec) mtp.set_max_drafts(S_mtp - 1);
+        strata::spec::SuffixDrafter sfx(std::max(1, o.suffix_draft), 64, (size_t) o.max_context + 4096);
+        strata::spec::DraftPolicy policy(o.spec);   // MTP or lookup window (see draft_policy.hpp)
+        std::vector<int32_t> sbuf((size_t) o.spec, 0);
+        int64_t sfx_windows = 0, sfx_drafts = 0, sfx_ok = 0;
+        if (o.suffix_draft > 0) {
+            for (int64_t t : o.tokens) sfx.append((int32_t) t);
+            for (int64_t t : produced) sfx.append((int32_t) t);
+        }
         const double pool_ms0 = drive.cpu_ms;
         const int64_t misses0 = drive.d.multi_misses, entries0 = drive.d.multi_entries;
         while ((int64_t) produced.size() < o.max_new) {
             const Clock::time_point t0 = Clock::now();
-            int T = o.spec;
+            int T = S_mtp;
             if (use_mtp && o.spec_min_p > 0.0) {
                 T = 1;
-                while (T < o.spec && dprob[(size_t) T - 1] >= (float) o.spec_min_p) ++T;
+                while (T < S_mtp && dprob[(size_t) T - 1] >= (float) o.spec_min_p) ++T;
             }
             if (first_window) T = 1;
+            bool from_sfx = false;
+            int sfx_match = 0;
+            if (o.suffix_draft > 0 && !first_window) {
+                const int k = sfx.propose(o.spec - 1, sbuf.data());
+                sfx_match = sfx.last_match();
+                if (k > 0 && (!use_mtp || sbuf[0] == drafts[0])) {
+                    const strata::spec::DraftPolicy::Pick pk = policy.choose(T, k, sfx_match);
+                    if (pk.lookup) { T = pk.t; from_sfx = true; }
+                }
+            }
+            const bool timed_round = !first_window;
             ++window_hist[(size_t) T];
             if (p + T > o.max_context) {
                 std::fprintf(stderr, "strata generate: ran out of context at position %lld\n", (long long) p);
@@ -2681,7 +3238,8 @@ int main(int argc, char** argv) {
             window[0] = x;
             for (int i = 1; i < T; ++i) {
                 const size_t at = produced.size() - 1 + (size_t) i;
-                int32_t d = use_mtp ? drafts[(size_t) i - 1] : at < oracle.size() ? (int32_t) oracle[at] : 0;
+                int32_t d = from_sfx ? sbuf[(size_t) i - 1] : use_mtp ? drafts[(size_t) i - 1]
+                                                    : at < oracle.size() ? (int32_t) oracle[at] : 0;
                 if (o.spec_corrupt > 0 && (++corrupt_counter % o.spec_corrupt) == 0) d = (d + 1) % (int32_t) n_vocab;
                 window[(size_t) i] = d;
             }
@@ -2720,9 +3278,11 @@ int main(int argc, char** argv) {
             drafts_total += T - 1;
             drafts_ok += a;
             ++accepted_hist[(size_t) a];
+            if (from_sfx) { ++sfx_windows; sfx_drafts += T - 1; sfx_ok += a; }
             bool eos = false;
             for (int i = 0; i <= a && (int64_t) produced.size() < o.max_new && !eos; ++i) {
                 produced.push_back(outv[(size_t) i]);
+                if (o.suffix_draft > 0) sfx.append(outv[(size_t) i]);
                 eos = o.stop_eos && std::find(o.eos_ids.begin(), o.eos_ids.end(), (int64_t) outv[(size_t) i]) != o.eos_ids.end();
             }
             if (eos) {
@@ -2740,7 +3300,9 @@ int main(int argc, char** argv) {
             }
             x = outv[(size_t) a];
             p += a + 1;
-            total_ms += std::chrono::duration<double, std::milli>(Clock::now() - t0).count();
+            const double round_ms = std::chrono::duration<double, std::milli>(Clock::now() - t0).count();
+            total_ms += round_ms;
+            if (timed_round) policy.observe(from_sfx, T, a, sfx_match, round_ms);
             if (rounds % 64 == 0)
                 std::fprintf(stderr, "strata generate: position %lld, %lld tokens, %lld rounds\n", (long long) p,
                              (long long) produced.size(), (long long) rounds);
@@ -2754,6 +3316,9 @@ int main(int argc, char** argv) {
             for (size_t i = 1; i < window_hist.size(); ++i) std::printf(" T%zu:%lld", i, (long long) window_hist[i]);
             std::printf("  (min draft probability %.2f)\n", o.spec_min_p);
         }
+        if (o.suffix_draft > 0)
+            std::printf("%-24s %lld windows, drafts accepted %lld of %lld\n", "suffix drafts", (long long) sfx_windows,
+                        (long long) sfx_ok, (long long) sfx_drafts);
         std::printf("%-24s", "accepted per round");
         for (size_t i = 0; i < accepted_hist.size(); ++i) std::printf(" %zu:%lld", i, (long long) accepted_hist[i]);
         std::printf("\n");

src/spec/draft_policy.cpp

@@ -0,0 +1,101 @@
+// src/spec/draft_policy.cpp - see include/strata/spec/draft_policy.hpp.
+#include "strata/spec/draft_policy.hpp"
+
+#include <algorithm>
+
+namespace strata::spec {
+namespace {
+
+// The shape of a round's cost by window size, relative to one token, used only for sizes not measured yet (the
+// measured round times replace it). Between the RTX 5070's measured curves: +10 ms per token with every missed expert
+// on the CPU (bench/results/2026-09-27-spec/window-cost), flatter with the default CPU/DMA split.
+constexpr double kShape[DraftPolicy::kMaxT + 1] = {0.0, 1.0, 1.35, 1.7, 2.05, 2.45, 2.85, 3.25, 3.6};
+constexpr double kCostAlpha = 0.1;    // EMA weight of a new round time
+constexpr double kTokAlpha = 0.05;    // EMA weight of a new MTP window outcome
+constexpr double kDecay = 0.97;       // lookup counts: older windows fade
+// Before a bucket has data: the longer the match, the likelier its continuation (llama.cpp's lookup decoding gates
+// on the same thing); worth 4 observations, so a few real windows override it
+constexpr double kPriorQ[DraftPolicy::kBuckets] = {0.75, 0.88, 0.93, 0.96};
+constexpr double kPriorN = 4.0;
+constexpr int kProbes = 3;       // a lookup window size is tried this often before its guessed cost can veto it
+
+}  // namespace
+
+DraftPolicy::DraftPolicy(int max_t, double margin)
+    : max_t_(std::clamp(max_t, 1, kMaxT)), margin_(margin) {}
+
+int DraftPolicy::bucket(int match) {
+    return match < 6 ? 0 : match < 12 ? 1 : match < 24 ? 2 : 3;
+}
+
+double DraftPolicy::lookup_rate(int match) const {
+    const int b = bucket(match);
+    return (ok_[b] + kPriorN * kPriorQ[b]) / (ok_[b] + bad_[b] + kPriorN);
+}
+
+double DraftPolicy::cost_ms(int t) const {
+    t = std::clamp(t, 1, kMaxT);
+    if (cost_n_[t] > 0) return cost_[t];
+    // scale from the measured sizes, weighting each by how often it was seen
+    double num = 0.0, den = 0.0;
+    for (int u = 1; u <= kMaxT; ++u)
+        if (cost_n_[u] > 0) {
+            const double w = std::min(cost_n_[u], 20.0);
+            num += w * cost_[u] * kShape[t] / kShape[u];
+            den += w;
+        }
+    return den > 0 ? num / den : kShape[t];
+}
+
+double DraftPolicy::mtp_tokens(int t) const {
+    if (mtp_n_[t] > 0) return mtp_tok_[t];
+    return 1.0 + 0.7 * (t - 1);       // before any MTP window of this size: a typical acceptance
+}
+
+DraftPolicy::Pick DraftPolicy::choose(int t_mtp, int lookup_k, int match) const {
+    Pick p;
+    p.t = std::clamp(t_mtp, 1, max_t_);
+    if (lookup_k <= 0) return p;
+    const double base = mtp_tokens(p.t) / cost_ms(p.t);
+    const double q = lookup_rate(match);
+    double e = 1.0, qi = 1.0, best = 0.0;
+    int best_t = 0;
+    for (int k = 1; k <= std::min(lookup_k, max_t_ - 1); ++k) {
+        qi *= q;
+        e += qi;
+        const double r = e / cost_ms(k + 1);
+        if (r > best) { best = r; best_t = k + 1; }
+    }
+    if (best_t > 0 && best > base * (1.0 + margin_)) {
+        p.lookup = true;
+        p.t = best_t;
+        return p;
+    }
+    // a guessed cost can keep the policy from ever measuring a size: the first few times a confident lookup would
+    // need a size not measured yet, it is tried (verification keeps the output; only the one round's speed is at stake)
+    const int t_full = std::min(lookup_k, max_t_ - 1) + 1;
+    if (t_full > p.t && cost_n_[t_full] < kProbes && q >= 0.85) {
+        p.lookup = true;
+        p.t = t_full;
+    }
+    return p;
+}
+
+void DraftPolicy::observe(bool lookup, int t, int accepted, int match, double round_ms) {
+    t = std::clamp(t, 1, kMaxT);
+    if (round_ms > 0) {
+        cost_[t] = cost_n_[t] > 0 ? (1.0 - kCostAlpha) * cost_[t] + kCostAlpha * round_ms : round_ms;
+        cost_n_[t] += 1.0;
+    }
+    if (lookup) {
+        const int b = bucket(match);
+        ok_[b] = kDecay * ok_[b] + accepted;
+        bad_[b] = kDecay * bad_[b] + (accepted < t - 1 ? 1.0 : 0.0);
+    } else {
+        const double got = accepted + 1.0;
+        mtp_tok_[t] = mtp_n_[t] > 0 ? (1.0 - kTokAlpha) * mtp_tok_[t] + kTokAlpha * got : got;
+        mtp_n_[t] += 1.0;
+    }
+}
+
+}  // namespace strata::spec

src/spec/draft_policy_test.cpp

@@ -0,0 +1,63 @@
+// src/spec/draft_policy_test.cpp - DraftPolicy: when does a lookup window beat the MTP's?
+//
+// Simulated rounds with the costs measured on the RTX 5070 (window-cost: ~10 ms more per token) check that
+//   1. with no lookup proposal the MTP window is kept;
+//   2. lookup drafts that are mostly rejected stop being taken (their bucket's rate falls);
+//   3. lookup drafts that are always accepted are taken, and the window grows with them;
+//   4. match-length buckets learn separately (short matches failing does not stop long ones);
+//   5. the policy never proposes a window beyond its cap.
+#include "strata/spec/draft_policy.hpp"
+
+#include <cstdio>
+
+using strata::spec::DraftPolicy;
+
+namespace {
+int g_fail = 0;
+void check(bool ok, const char* what) {
+    std::printf("  %-66s %s\n", what, ok ? "ok" : "FAIL");
+    if (!ok) ++g_fail;
+}
+double cost(int t) { return 19.0 + 10.5 * (t - 1); }   // ms per round, measured shape
+}  // namespace
+
+int main() {
+    std::printf("draft_policy_test\n");
+    {
+        DraftPolicy p(6);
+        for (int i = 0; i < 50; ++i) p.observe(false, 4, 2, 0, cost(4));   // MTP windows of 4: 3 tokens each
+        const DraftPolicy::Pick k = p.choose(4, 0, 0);
+        check(!k.lookup && k.t == 4, "no proposal: the MTP window");
+    }
+    {
+        DraftPolicy p(6);
+        for (int i = 0; i < 50; ++i) p.observe(false, 4, 2, 0, cost(4));
+        for (int t = 2; t <= 6; ++t) p.observe(false, t, 0, 0, cost(t));
+        for (int i = 0; i < 40; ++i) p.observe(true, 6, 0, 4, cost(6));      // short matches, all rejected
+        check(p.lookup_rate(4) < 0.15, "rejected short-match drafts: their rate falls below 0.15");
+        check(!p.choose(4, 5, 4).lookup, "rejected short-match drafts: no longer taken");
+        for (int i = 0; i < 40; ++i) p.observe(true, 6, 5, 30, cost(6));     // long matches, all accepted
+        check(p.lookup_rate(30) > 0.9, "accepted long-match drafts: their rate rises above 0.9");
+        const DraftPolicy::Pick k = p.choose(4, 5, 30);
+        check(k.lookup && k.t == 6, "accepted long matches: the full lookup window is taken");
+        check(!p.choose(4, 5, 4).lookup, "buckets are separate: short matches still not taken");
+        check(p.choose(4, 20, 30).t <= 6, "never beyond the window cap");
+    }
+    {
+        DraftPolicy p(8);
+        for (int i = 0; i < 50; ++i) p.observe(false, 3, 2, 0, cost(3));      // a very good MTP: 3 of 3 tokens
+        for (int t = 2; t <= 8; ++t) p.observe(false, t, t - 1, 0, cost(t));
+        for (int i = 0; i < 40; ++i) p.observe(true, 4, 2, 8, cost(4));       // lookup at q ~ 0.67
+        check(!p.choose(3, 7, 8).lookup, "a mediocre lookup does not replace a strong MTP window");
+    }
+    {
+        DraftPolicy p(6);
+        for (int i = 0; i < 50; ++i) p.observe(false, 4, 3, 0, cost(4));      // a near-perfect MTP, only size 4 seen
+        const DraftPolicy::Pick k = p.choose(4, 5, 40);
+        check(k.lookup && k.t == 6, "an unmeasured size is probed for a confident lookup");
+        for (int i = 0; i < 3; ++i) p.observe(true, 6, 5, 40, 3.0 * cost(6));   // it turns out very expensive
+        check(!p.choose(4, 5, 40).lookup, "after the probes, the measured cost decides");
+    }
+    std::printf(g_fail ? "FAIL\n" : "PASS\n");
+    return g_fail ? 1 : 0;
+}

tools/_paths.py

@@ -14,7 +14,7 @@ ENGINE = Path(__file__).resolve().parents[1]
 
 def gguf_py() -> str:
     cands = [os.environ.get("STRATA_GGUF_PY"), ENGINE / "third_party" / "llama.cpp" / "gguf-py",
-             ENGINE.parents[1] / ".ref" / "llama.cpp" / "gguf-py"]
+             ENGINE.parents[min(1, len(ENGINE.parents) - 1)] / ".ref" / "llama.cpp" / "gguf-py"]   # the dev tree; a shallow install (G:\Strata) has no grandparent
     for c in cands:
         if c and (Path(c) / "gguf").is_dir():
             return str(c)

tools/mtp_rt.py

@@ -19,7 +19,6 @@ from pathlib import Path
 
 import numpy as np
 
-STRATA = Path(__file__).resolve().parents[3]
 sys.path.insert(0, str(Path(__file__).resolve().parent))
 from _paths import add_gguf_py  # noqa: E402
 add_gguf_py()

tools/pack_layer.py

@@ -34,7 +34,7 @@ sys.path.insert(0, str(pathlib.Path(__file__).resolve().parent))
 from gguf_reader import GGUFFile  # noqa: E402
 
 REPO = pathlib.Path(__file__).resolve().parents[1]
-STRATA = REPO.parents[1]          # the development layout's default shards (setup.py always passes --gguf)
+STRATA = REPO.parents[min(1, len(REPO.parents) - 1)]   # the dev layout's default shards (setup.py passes --gguf)
 SHARD1 = STRATA / "Q2_0" / "Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf"
 
 H, FF, NE = 2560, 640, 512

tools/strata_pack.py

@@ -41,7 +41,8 @@ from canonical_xcheck import (MAPPINGS, KV_IQ4NL, open_shard,             # noqa
 
 FORMAT_VERSION = 1
 ALIGN = 64                       # every plane starts on a 64-byte boundary (P1.S7)
-STRATA = pathlib.Path(__file__).resolve().parents[3]   # the development layout; build takes the shard from --gguf
+_HERE = pathlib.Path(__file__).resolve()
+STRATA = _HERE.parents[min(3, len(_HERE.parents) - 1)]   # the development layout; build takes the shard from --gguf
 SHARD2 = STRATA / "Q2_0" / "Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf"
 
 EXPERT_RE = ("ffn_gate_exps.weight", "ffn_up_exps.weight", "ffn_down_exps.weight")