Merge branch 'f-139b-par' into f-139-rel

# Conflicts:
#	src/core/verify.cpp
#	src/program/generate.cpp

README.md

@@ -136,7 +136,8 @@ Sizes, downloads and what fits where: [docs/MODELS.md](docs/MODELS.md). You can
 - **Pictures:** say yes to "Images?" in setup, then click **Picture** in the chat, or attach them in your app
   (AMD cards: on Linux through the processor, not on Windows yet).
 - **From your phone or another PC:** `START-HERE.bat --setup --host 0.0.0.0 --api-key <secret>` - always with a key.
-- **Good to know:** it answers one request at a time. The first message of a chat is read in full (about 1 minute
+- **Good to know:** it answers one request at a time (several at once is opt-in: `"parallel": 2`,
+  [BATCHING.md](docs/BATCHING.md)). The first message of a chat is read in full (about 1 minute
   per 30,000 tokens); follow-ups start in seconds.
 
 More: [where your chats are stored](docs/INSTALL.md#where-things-are-stored), [the API](docs/DETAILS.md#using-it).

docs/AI_SETUP.md

@@ -174,7 +174,8 @@ line when it is ready; the engine log is `strata-<model>.log` in the Strata fold
 - **Claude Code:** `ANTHROPIC_BASE_URL=http://127.0.0.1:8080`, `ANTHROPIC_MODEL` set to a Claude model name it knows
   (Strata ignores the name), and any `ANTHROPIC_AUTH_TOKEN` (or the configured key).
 - **Thinking level:** `"reasoning_effort": "none" | "low" | "medium" | "high"` (default high).
-- Strata answers one request at a time. API details: [DETAILS.md](DETAILS.md#using-it).
+- Strata answers one request at a time; `"parallel": N` in the model's config (or setup `--parallel N`) decodes up to
+  N together ([BATCHING.md](BATCHING.md)). API details: [DETAILS.md](DETAILS.md#using-it).
 
 ## 9. When something fails
 

docs/BATCHING.md

@@ -0,0 +1,215 @@
+# Several requests at once (batch slots)
+
+By default Strata serves **one request at a time**: the others wait in the server's queue. With `"parallel": N`
+(the engine's `--batch N`, also spelled `--slots N`) the engine keeps up to N conversations open and decodes them
+**together**: every verify window then carries one token of each conversation, so the dense weights, the shared
+expert, the head and every routed expert two conversations share are read once per window for all of them.
+Combined with a layer split across several GPUs and `--batch-groups`, the cards also work on different
+conversations at the same time instead of waiting for each other.
+
+It is opt-in and changes nothing when the options are absent (#465; the engine part is PR #559).
+
+## Turning it on
+
+One GPU: add `"parallel": 2` to the model's config (`strata-<model>.json`) and restart, or run setup with
+`--parallel 2`. Setup recommends it only where it does not cost speed (below); any number you ask for is kept as
+asked, with a note when it is more than setup would recommend.
+
+```
+"parallel": 2
+```
+
+With a layer split, the engine options go into the config's `args`:
+
+```
+"args": [ ..., "--batch", "8", "--batch-groups", "4", "--trim-stage-weights" ],
+"layer_split": "12,24,36"
+```
+
+| Option | What it does |
+| --- | --- |
+| `"parallel": N` / `--batch N` / `--slots N` (2..8) | up to N conversations decoded together; more requests wait for a free slot. Each slot gets its own state (a session carved like the stage's own: GDN recurrence, QSA K/V and indexer, PLE history) on every GPU of the split. |
+| `--batch-groups G` | with a layer split: the N slots in G groups that flow through the GPUs as a pipeline (GPU k runs one group while GPU k+1 runs another). G must divide N. 1 = all slots in one window, GPU after GPU. |
+| `--trim-stage-weights` | with an **explicit** `--layer-split` (e.g. `12,24,36`, not `auto`): every GPU loads only the dense weights of its own layers instead of the whole model's (the same as `STRATA_STAGE_TRIM=1`, PR #639). The VRAM this frees goes to the expert cache. Useful without `--batch` too. |
+
+The engine never refuses a count it cannot run: it says so in its log and runs what it can - at most 8 slots (a
+window holds 8 rows), as many as fit in VRAM, or none (one request at a time) when not two fit. The server reads
+the count the engine reports (`INFO batch_slots=N`), and `GET /v1/status` says it (`concurrency.serving`).
+
+### What a slot costs, and what setup recommends
+
+Every slot's session takes VRAM that the expert cache would otherwise hold: 0.56 GiB at a 32K context with 8-bit
+KV, more with a longer context unless the KV cache streams (`--kv-resident`: then only the attention's 32K window
+stays in VRAM, and each slot's whole KV cache takes pinned RAM - 1.6 GB at 128K). On a card whose experts mostly run
+on the CPU, a batch also reads about as many distinct experts as the requests one by one (different conversations
+route to different experts), so the gain is in **latency** (nobody waits for a whole answer), and a request alone
+runs slower (the smaller expert cache): 11-24 % on a 12 GB card, see the measurements below.
+
+So setup recommends `"parallel"` **only where the experts mostly fit in VRAM**: the expert cache (each card's VRAM
+less ~5 GB, every card of a layer split counted) must still hold at least half of the model's experts beside the
+slots, and the slots may take at most a fifth of it, up to 4 slots. With Q2_0 at 32K that is 3 slots on a 24 GB
+card, 4 from 32 GB or on a split such as 2 x 16 GB; IQ3_S needs 32 GB or a split. Everywhere else (any 12 or 16 GB
+card alone) it stays at one at a time and setup says: "parallel N reduces waiting for several users but costs
+about 10-25% speed per request on this card". `--parallel N` is honoured as asked either way.
+
+## How the server uses the slots
+
+- **One request alone** runs on the usual solo path (verify windows with MTP drafts): the fastest single stream.
+- **When a second request arrives**, the first is stopped (`STOP`) and continues in a batch slot with its prompt
+  plus what it generated so far - the engine's prompt cache holds exactly that, so nothing is read again - and the
+  new request is admitted next to it. A request in a slot decodes **without MTP drafts** (one token per window).
+- **A request left alone in a slot** (the others finished, nobody waits) goes back to the solo path: the slot is
+  stopped, the engine copies its sessions back and decodes with MTP drafts again (at most twice per request; with
+  `--prompt-cache 0` it stays in the slot; `STRATA_PARALLEL_SOLO=0` turns it off). The draft layer's own K/V was
+  built for another conversation then, but measured it accepted as many drafts (140 of 172) as a draft layer that
+  read the conversation (140 of 173).
+- **More requests than slots** wait for a free one (`/metrics` -> `live.slots` shows each slot: idle, reading or
+  decoding, its tokens and tok/s; `live.running` the requests in flight).
+- **Each admission** reads the request's prompt through the usual prompt path (prompt cache and conversation
+  checkpoints included) and produces its first token there; the state is then copied into the slot. Admissions
+  are taken one at a time, and **the slots decode between the prompt's chunks**: after each chunk (`--prefill`,
+  2048-8192 tokens) they decode for half as long as the chunk took (`STRATA_BATCH_DECODE_SHARE`, default 0.5), so
+  a long prompt slows the others down instead of stopping them. The chunks are the ones one uninterrupted read
+  takes, so the prompt's arithmetic is unchanged.
+- **A long prompt gives way to a short one** (#656's cooperative preemption): when a request with a prompt under
+  half as long is waiting, the server sends `BYIELD`; at its next chunk boundary the long read stops, the part read
+  so far is copied into a slot, the short request is admitted, and the long one then goes on from its slot with the
+  same chunks (at most twice per request).
+- **Each slot is a conversation cache.** A finished slot keeps what it holds (the prompt, the answer, and the
+  checkpoint at the prompt's last turn boundary); the next turn of that conversation goes to that slot and the
+  engine copies its state back (50-60 ms for a short conversation) instead of reading the history again - also for
+  a client that drops the reply's thinking from the history (the checkpoint matches up to the new turn). A new
+  conversation takes an empty slot, else the one used longest ago.
+- Slots are assigned so that consecutive requests land in different pipeline groups (`--batch-groups`).
+- A client that disconnects stops its slot (`BSTOP`); the others go on.
+
+## Exactness
+
+A batch row's arithmetic is the single-token window's, so with greedy decoding **every conversation of a batch
+produces exactly the tokens it produces alone** - verified token by token for 8 concurrent conversations of 150
+tokens, with and without the pipeline (`tools/batch_test.py`), and on one RTX 5070 for 4 conversations, for a long
+prompt read while two others decode, for a prompt that gave way and went on, for a next turn continued from its
+slot (from all it holds, and from its turn checkpoint without the reply's thinking), and for a request stopped in
+its slot and continued on the solo path (`tools/batch_interleave_test.py`). These
+settings make the comparison exact:
+
+- `STRATA_IQ_MT_MIN=1` (the multi-token CPU expert kernels for every group, as for the solo path's own
+  exactness tests: by default an expert's rows round differently alone than in a group, so the output depends on
+  how many rows of a window share an expert - which differs between a batch and a request alone),
+- `--pcie-frac 0`: the PCIe share of the missed experts is chosen per window from the window's misses, so the
+  same expert can run on the GPU in one window and on the CPU in another, which rounds differently, and
+- `--adapt-every 1000000` (the VRAM tier fixed), and `--no-prefill-borrow` while a prompt is read beside decoding
+  slots: their windows then see the expert cache without the slots the prompt borrowed (those experts run on the
+  CPU).
+
+With the default settings the outputs stay coherent but drift apart after some tokens, as two solo runs whose
+windows differ can. Measured on the RTX 5070 (Q2_0, 4 conversations of 200 tokens, `tools/batch_test.py` without
+the settings above): one equal to its solo run, the others apart from token 0, 99 and 174 - the first token already
+differed for one, because the adaptive VRAM tier had moved experts between the solo runs and the batch; all four
+texts read as well as their solo runs.
+
+Sampled requests (temperature, top_p, top_k, min_p, seed) are drawn row by row with the solo window's
+counter-based draw (Philox(seed, position)).
+
+## Limits (for now)
+
+- Batch windows carry no MTP drafts: a conversation in a slot decodes one token per window (the solo path keeps
+  its drafts, which is why a request alone is not put in a slot, and goes back to it when left alone).
+- Repetition / frequency / presence penalties are not applied in batch windows.
+- A prompt shorter than one chunk is read in one piece (the slots wait for it); a read gives way only at a chunk
+  boundary, and not for pictures.
+- Admissions are one at a time: two new long prompts are read one after the other.
+- `--batch-groups` needs every stage on its own GPU; a pipelined slot is not kept as a conversation cache.
+- The slot sessions take VRAM (above) and, with KV streaming, pinned RAM.
+
+## Measured
+
+One RTX 5070 (12 GB), Ryzen 5 7600, 64 GB DDR5, Q2_0, 32K context, through the HTTP server: C different requests
+sent at once (an 800-word essay each, 256 tokens per answer, greedy, thinking off), median of 3 rounds:
+
+| Concurrent | Setting | Total tok/s | vs one at a time | Per request tok/s | First token: median / last of the round |
+| ---: | --- | ---: | ---: | ---: | ---: |
+| 1 | one at a time (default; what setup recommends on this card) | 74.3 | - | 83.6 | 0.4 s / 0.4 s |
+| 1 | `"parallel": 2` | 67.3 | -9 % | 74.8 | 0.4 s / 0.4 s |
+| 1 | `"parallel": 4` | 57.9 | -22 % | 63.8 | 0.4 s / 0.4 s |
+| 2 | one at a time | 71.6 | - | 77.3 | 2.1 s / 4.1 s |
+| 2 | `"parallel": 2` | 61.0 | -15 % | 32.6 | 0.5 s / 0.9 s |
+| 2 | `"parallel": 4` | 54.6 | -24 % | 28.7 | 0.6 s / 0.7 s |
+| 4 | one at a time | 70.7 | - | 79.6 | 6.0 s / 11.2 s |
+| 4 | `"parallel": 2` | 61.2 | -13 % | 32.2 | 4.5 s / 9.3 s |
+| 4 | `"parallel": 4` | 63.1 | -11 % | 16.9 | 1.0 s / 1.8 s |
+
+On this card the slots buy **waiting time, not speed**: the fourth of four requests starts after 1.8 s instead of
+11.2 s, but together they decode 11-24 % slower than one after the other, and a request alone loses 11 % (2 slots)
+or 24 % (4 slots), because the slots' sessions (0.56 GiB each) come out of the expert cache and most experts run
+on the CPU: a batch window over 4 conversations reads 24 CPU experts per layer against ~8 for one, so it costs about
+what the 4 tokens cost one after the other (`strata batch:` in the engine log: 54 ms per 4-row window, ~20 ms per
+1-row window). This is why setup leaves a 12 GB card at one at a time. Cards that hold most experts in VRAM, and a
+layer split, are where the slots also add speed (below).
+
+A 4-GPU layer split (4 x 16 GB, PCIe Gen3), IQ3_S, `--batch 8 --batch-groups 4 --trim-stage-weights`, through
+the HTTP server, 400 tokens per answer, temperature 0.7 (PR #559):
+
+| Concurrent requests | Per request | Total |
+| ---: | ---: | ---: |
+| 1 | 123 tok/s (solo path) | 120 tok/s |
+| 2 | 57 tok/s | 113 tok/s |
+| 4 | 51 tok/s | 205 tok/s |
+| 8 | 45 tok/s | 360 tok/s |
+
+With the patches below on engine 0.1.38 and parking on, through the service: 8 requests at temperature 0 -> 369
+tok/s, at 0.7 -> 358 tok/s.
+
+`--trim-stage-weights` alone raised the share of experts held in VRAM on that machine from 76-85 % to 84-100 %
+per card.
+
+## Together with conversation parking
+
+`--conversation-cache-mib N --conversation-cache-slots K` (DETAILS.md) works with the layer split too: a request
+whose conversation was parked is restored on every stage before its admission, so an agent and its sub-agents, or
+several chats that alternate, come back without reading their history again. Measured on the same 4-GPU split,
+two long conversations alternating through the HTTP server: the first turns took 3.9 s and 5.7 s to the first token
+(their prompts read), the follow-ups 0.53 s and 0.46 s.
+
+## Testing
+
+Four scripts drive a built engine or a running server; each exits non-zero on a failure. `serve/test_parallel.py`
+tests the server's side with a scripted engine (no GPU).
+
+| Script | What it checks |
+| --- | --- |
+| `tools/batch_test.py` | the same prompts alone (`GEN`) and together in the batch slots (`BGEN`): every slot's greedy tokens equal its solo tokens; prints the aggregate rate. `--batch-groups` in `--extra` tests the pipeline, `--keys "temperature=0.7"` the sampled rows. |
+| `tools/batch_interleave_test.py` | a long prompt read while two slots decode, a prompt that gives way (`BYIELD`) and goes on, and a next turn continued from its slot: each equal to its solo tokens. |
+| `tools/parking_test.py` | a follow-up to a conversation decodes the same tokens whether its state stayed live or came back from the parking cache (with a layer split: every stage's image). |
+| `tools/early_close_test.py` | a client that stops reading a streamed answer early (alone, and with a second request running) does not leave its tokens to the next request (server). |
+
+For exact comparisons pass `--pcie-frac 0 --adapt-every 1000000` (and the scripts set `STRATA_IQ_MT_MIN=1`):
+
+```
+python3 tools/batch_test.py --exe engine/strata --config strata-<model>.json --batch 8 --n 8 \
+    --extra "--layer-split 12,24,36 --trim-stage-weights --batch-groups 4 --pcie-frac 0 --adapt-every 1000000"
+python3 tools/batch_interleave_test.py --exe engine/strata --config strata-<model>.json \
+    --extra "--pcie-frac 0 --adapt-every 1000000 --no-prefill-borrow"
+python3 tools/parking_test.py --exe engine/strata --config strata-<model>.json \
+    --extra "--layer-split 12,24,36 --conversation-cache-mib 8192 --conversation-cache-slots 4 --pcie-frac 0"
+STRATA_KEY=<key> python3 tools/early_close_test.py http://127.0.0.1:8080
+```
+
+## Engine protocol (`--serve`)
+
+On top of `GEN` / `GENI`:
+
+| Line | Direction | Meaning |
+| --- | --- | --- |
+| `BGEN <slot> <max_new> [keys] <ids>` | in | read the prompt (as `GEN 1`), then continue in `<slot>` |
+| `BGENI <slot> <max_new> [keys] <file> <ids>` | in | the same with images |
+| `BADM <slot> <1/0>` | out | after the admission's `DONE`: 1 = it continues in the slot, 0 = it ended |
+| `BT <slot> <id>` | out | a token of that slot |
+| `BDONE <slot> <generated> <stop/length/cancel> <ms>` | out | the slot is free again (it keeps its conversation) |
+| `BSTOP <slot>` | in | end that slot at its next window |
+| `BYIELD <slot>` | in | the prompt being read gives way at its next chunk boundary; its part read waits in `<slot>` (the admission's own, or a free slot for a solo request) |
+| `YIELDED <slot> <tokens>` | out | before the `DONE cancel` of a read that gave way: the request is sent again later and goes on from there |
+| `INFO ... batch_slots=N` | out | the slots the engine runs (only with `--batch`) |
+
+`tools/batch_test.py` drives the engine directly: the same prompts alone, then together, compared token by token,
+and the aggregate rate.

docs/DETAILS.md

@@ -642,9 +642,11 @@ in a bounded 8 GiB host-RAM cache. This preserves controller/worker histories wh
 their requests alternate; it does not execute requests concurrently. No client session
 ID is required: only exact token/image prefixes with matching steering mode are reused.
 The default budget is 0 (disabled); `--prompt-cache 0` also disables parking.
-The initial shared-core integration supports a single session GPU: combining
-enabled parking with `--layer-split` is rejected before model loading. Ordinary
-upstream layer-split checkpoints remain available with parking disabled. FP16,
+With `--layer-split`, a parked conversation holds one image per stage (each later stage's
+session and its share of the conversation checkpoints), every image is validated before
+anything is overwritten, and each is restored on its own GPU. Verified on a 4-GPU split:
+a 9,276-token conversation restored in 51 ms, and its follow-up decodes exactly the tokens
+it decodes when the conversation never left (`tools/parking_test.py`). FP16,
 INT8, Q4_0 and identity-layout K8V4 snapshots are supported; the K8V4 draft ring
 remains INT8, as in upstream. Windows/HIP and multi-GPU runtime coverage must be
 reported separately from Linux/CUDA evidence.
@@ -681,8 +683,9 @@ sizes and K/V bytes reused during capture. `STRATA_SNAPSHOT_FULL_CAPTURE=1` disa
 retention for diagnostic comparisons. Snapshots are not
 persisted across restarts.
 
-**Current limits (v1):** one request at a time, and one conversation cached at a time (switching between two chats
-re-reads the other one unless the opt-in cache above is enabled); images only when set up with them (below); no video. **Temperature / top_p / top_k / min_p /
+**Current limits (v1):** one request at a time unless `"parallel": N` is set (opt-in batch slots, up to N requests
+decoded together: [BATCHING.md](BATCHING.md)), and one conversation cached at a time (switching between two chats
+re-reads the other one unless the opt-in cache above is enabled, or each conversation keeps its own batch slot); 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

docs/MULTI_GPU.md

@@ -167,3 +167,8 @@ The Coder on an RTX 5080 + RTX 3090 (Ryzen 9 9950X3D), 32K context; details in
 - Leave out a much slower card when two already hold the model. An RTX 2080 Ti as a third card made the 5080 +
   3090 pair slower (68 / 90 tok/s decode): every extra card costs its own round per window.
 - More cards pay off when the model's routed experts do not fit the faster ones.
+
+## Several conversations at once
+
+With a layer split, `--batch N --batch-groups G --trim-stage-weights` decodes several conversations together and
+pipelines them through the cards: see [BATCHING.md](BATCHING.md).

include/strata/core/conversation_cache.hpp

@@ -53,13 +53,72 @@ struct ConversationKvReuse {
     std::vector<ConversationKv> kv;
     // Original image extent for validation, and the earliest subsequent rewrite.
     int64_t captured_tokens = 0, unchanged_tokens = 0;
+    // With a layer split: the later stages' own retained K/V, one per stage, same extents (empty: none)
+    std::vector<ConversationKvReuse> stages;
     size_t bytes() const {
-        size_t n = kv.capacity() * sizeof(ConversationKv);
+        size_t n = kv.capacity() * sizeof(ConversationKv) + stages.capacity() * sizeof(ConversationKvReuse);
         for (const auto& layer : kv) n += layer.bytes();
+        for (const auto& s : stages) n += s.bytes();
         return n;
     }
 };
 
+/// Layer-split parking: each running checkpoint holds the first stage's state and, in `stage_parts`, one part per
+/// later stage; a parked conversation keeps them in one image per stage.  `split` MOVES the checkpoints apart (no
+/// running state is copied): the complete ones (a part for each of `stages` stages) into `stage0` and `parts[k]`,
+/// every part carrying the checkpoint's ids / images / LRU stamp, the others aside; `merge` puts them back in their
+/// original order.  merge(split(x)) == x.
+struct ConversationCheckpointSplit {
+    std::vector<ConversationCheckpoint> stage0;
+    std::vector<std::vector<ConversationCheckpoint>> parts;   ///< [stage][checkpoint]
+    std::vector<ConversationCheckpoint> rest;                 ///< not split (no complete set of stage parts)
+    std::vector<bool> complete;                               ///< per original position: in stage0/parts or in rest
+};
+
+inline ConversationCheckpointSplit conversation_checkpoints_split(std::vector<ConversationCheckpoint>&& checks,
+                                                                  size_t stages) {
+    ConversationCheckpointSplit out;
+    out.parts.resize(stages);
+    for (auto& c : checks) {
+        const bool whole = c.stage_parts.size() == stages;
+        out.complete.push_back(whole);
+        if (!whole) { out.rest.push_back(std::move(c)); continue; }
+        for (size_t k = 0; k < stages; ++k) {
+            ConversationCheckpoint part = std::move(c.stage_parts[k]);
+            part.ids = c.ids; part.imgs = c.imgs; part.used = c.used; part.stage_parts.clear();
+            out.parts[k].push_back(std::move(part));
+        }
+        c.stage_parts.clear();
+        out.stage0.push_back(std::move(c));
+    }
+    checks.clear();
+    return out;
+}
+
+/// false (and `out` empty) when the parts do not line up: a stage with another number of checkpoints, or a part
+/// whose ids differ from its checkpoint's.
+inline bool conversation_checkpoints_merge(ConversationCheckpointSplit&& split, std::vector<ConversationCheckpoint>& out) {
+    out.clear();
+    for (const auto& p : split.parts)
+        if (p.size() != split.stage0.size()) return false;
+    size_t whole = 0, other = 0;
+    for (bool c : split.complete) (c ? whole : other) += 1;
+    if (!split.complete.empty() && (whole != split.stage0.size() || other != split.rest.size())) return false;
+    if (split.complete.empty()) split.complete.assign(split.stage0.size(), true);   // a restored image: all whole
+    for (size_t j = 0; j < split.stage0.size(); ++j)
+        for (const auto& p : split.parts)
+            if (p[j].ids != split.stage0[j].ids) return false;
+    size_t w = 0, r = 0;
+    for (bool c : split.complete) {
+        if (!c) { out.push_back(std::move(split.rest[r++])); continue; }
+        ConversationCheckpoint cp = std::move(split.stage0[w]);
+        for (auto& p : split.parts) cp.stage_parts.push_back(std::move(p[w]));
+        ++w;
+        out.push_back(std::move(cp));
+    }
+    return true;
+}
+
 struct SavedConversation {
     // Runtime compatibility only; NOT a model/weights identity or disk schema.
     std::array<int64_t, 18> geometry{};
@@ -69,10 +128,13 @@ struct SavedConversation {
     std::vector<ConversationCheckpoint> checkpoints;
     std::vector<ConversationKv> kv; // main layers followed by the draft layer
     bool cvec = true;
+    // with a layer split, the later stages' own images, one per stage, in stage order
+    std::vector<SavedConversation> stage_images;
 
     size_t bytes() const {
         size_t n = live.bytes() + checkpoints.capacity() * sizeof(ConversationCheckpoint) +
                    kv.capacity() * sizeof(ConversationKv);
+        for (const auto& s : stage_images) n += s.bytes();
         for (const auto& c : checkpoints) n += c.bytes();
         for (const auto& k : kv) n += k.bytes();
         return n;
@@ -110,13 +172,17 @@ public:
 
     // Retain only the restored K/V buffers, not duplicate running checkpoints.
     // This optimization never evicts a parked conversation to make itself fit.
-    void retain(std::vector<ConversationKv>&& kv, int64_t tokens) {
+    // `stage_kv`: with a layer split, the later stages' restored K/V (one per stage), retained with the first's.
+    void retain(std::vector<ConversationKv>&& kv, int64_t tokens,
+                std::vector<std::vector<ConversationKv>>&& stage_kv = {}) {
         reuse_ = {};
-        ConversationKvReuse candidate{std::move(kv), tokens, tokens};
+        ConversationKvReuse candidate{std::move(kv), tokens, tokens, {}};
+        for (auto& k : stage_kv) candidate.stages.push_back(ConversationKvReuse{std::move(k), tokens, tokens, {}});
         if (enabled() && candidate.bytes() <= budget_ - bytes_) reuse_ = std::move(candidate);
     }
     void limit_reuse(int64_t first_dirty) {
         reuse_.unchanged_tokens = std::min(reuse_.unchanged_tokens, first_dirty);
+        for (auto& s : reuse_.stages) s.unchanged_tokens = std::min(s.unchanged_tokens, first_dirty);
         if (reuse_.unchanged_tokens <= 0) reuse_ = {};
     }
     ConversationKvReuse take_reuse() { return std::exchange(reuse_, {}); }

include/strata/core/conversation_snapshot.hpp

@@ -76,4 +76,21 @@ ConversationRestore conversation_snapshot_restore(const SavedConversation& image
                                                    const ModelGeometry& g, const QsaState& draft,
                                                    std::string& error);
 
+// The same with `draft == nullptr`: an image WITHOUT the draft layer's K/V (kv holds the session's own QSA layers
+// only).  A layer split's later stages park this way; the draft ring is saved once, with the first stage's image.
+// An image is validated and restored with the same kind of call it was saved with (a K/V layer count mismatch is
+// rejected as invalid).
+bool conversation_snapshot_bytes(const ConversationView& view, const SessionState& session, const ModelGeometry& g,
+                                 const QsaState* draft, size_t& bytes, std::string& error);
+bool conversation_snapshot_capture_bytes(const ConversationKvReuse& reuse, const ConversationView& view,
+                                         const SessionState& session, const ModelGeometry& g, const QsaState* draft,
+                                         size_t& bytes, std::string& error);
+bool conversation_snapshot_save(SavedConversation& image, const ConversationView& view, const SessionState& session,
+                                const ModelGeometry& g, const QsaState* draft, std::string& error,
+                                ConversationKvReuse reuse = {}, size_t* reused_bytes = nullptr);
+bool conversation_snapshot_validate(const SavedConversation& image, const SessionState& session,
+                                    const ModelGeometry& g, const QsaState* draft, std::string& error);
+ConversationRestore conversation_snapshot_restore(const SavedConversation& image, SessionState& session,
+                                                   const ModelGeometry& g, const QsaState* draft, std::string& error);
+
 } // namespace strata::core

include/strata/core/native_dense.hpp

@@ -19,8 +19,10 @@ public:
     ~NativeDense();
     NativeDense(const NativeDense&) = delete;
     NativeDense& operator=(const NativeDense&) = delete;
+    /// With layer_hi >= 0, only the `blk.<l>.` matrices with layer_lo <= l < layer_hi are uploaded (a layer
+    /// split's stage holds its own layers' projections, not the whole model's); the others keep data == nullptr.
     bool load(const std::vector<std::string>& shards, WeightTable& table, std::string& err,
-              bool include_ple_key = false);
+              bool include_ple_key = false, int64_t layer_lo = 0, int64_t layer_hi = -1);
     /// Plan v0.3 P1: the canonical tensor names `load` would serve natively from these shards (eligible name,
     /// supported type, 2-D), read from the GGUF headers only - so the canonical arena can skip them.
     static bool served_names(const std::vector<std::string>& shards, bool include_ple_key,

include/strata/core/verify.hpp

@@ -31,6 +31,8 @@
 #include <cuda_runtime.h>
 
 #include <atomic>
+#include <chrono>
+#include <map>
 #include <cstdint>
 #include <string>
 #include <vector>
@@ -125,6 +127,50 @@ public:
 
     /// 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);
+
+    // ================================ SEVERAL SEQUENCES IN ONE WINDOW ================================
+    //
+    // A batch window holds S INDEPENDENT sequences, one token each: row s is slot s, at slot s's own position,
+    // reading and writing slot s's own state (GDN recurrence and conv history, QSA K/V and indexer, PLE history),
+    // which lives in `slots[s]` - a session carved like this verifier's own (same layer range, same max_cells).
+    // Everything that is per row already (hyper-connections, dense projections, router, shared expert, the
+    // routed experts on the GPU and the CPU, the head) runs once over the S rows, so the weights are read once
+    // per window for all the sequences.  Row s's arithmetic is the single-token window's, so a slot's greedy
+    // tokens are its solo greedy tokens (modulo the multi-token CPU kernel choice: STRATA_IQ_MT_MIN=1).
+    //
+    // Greedy only, no drafts (MTP) in a batch window.  `init_slots` once after `init` (S <= max_t); a layer
+    // split's stages each get their own sessions, and run_slots/commit_slots continue into the next stage.
+    bool init_slots(const std::vector<SessionState*>& slots, std::string& err);
+    int n_slots() const { return (int) slots_.size(); }
+    /// One batch window over slots [0, S): tokens[s] at positions pos[s]; out[s] = the greedy pick after it.
+    bool run_slots(int S, const int32_t* tokens, const int64_t* pos, PoolMultiFn pool, void* user, int32_t* out,
+                   std::string& err);
+    /// The same over the S slots `rows` (row t is slot rows[t], any distinct slots in any order): the slots not
+    /// listed are not touched, so an idle slot keeps its state (a finished conversation it may continue later).
+    bool run_slot_rows(const int* rows, int S, const int32_t* tokens, const int64_t* pos, PoolMultiFn pool, void* user,
+                       int32_t* out, std::string& err);
+    /// Keep every row of the last batch window: each slot's state advances by its one token.
+    bool commit_slots(std::string& err);
+
+    // ---- The stages of a layer split as a PIPELINE.  A batch window over the slot GROUP
+    // [base, base + S) is launched on ONE stage with its commit right behind it on the stage's stream (a batch window
+    // keeps every row, so the commit needs no host decision), and the host serves the rings of every stage that has
+    // a window in flight from one thread (batch_poll does not block).  Stage k can then run group g while stage k+1
+    // runs group g-1.  Rows of group `base` use hand-off rows [base, base + S), so groups never share a hand-off row.
+    bool batch_launch(int base, int S, const int32_t* tokens, const int64_t* pos, std::string& err);
+    /// 1 = this stage's window and commit are done (the last stage's picks are in batch_out), 0 = still running,
+    /// -1 = an error (err).  Serves every layer that has rung so far.
+    int batch_poll(PoolMultiFn pool, void* user, std::string& err);
+    bool batch_busy() const { return b_running_; }
+    /// A slot's sampling (temperature / top_p / top_k / min_p / seed; penalties are not applied in batch windows):
+    /// its row is drawn again on the last stage with Philox(seed, position), as a solo window draws it.  Greedy by
+    /// default.  Set on the first stage, it reaches the last.
+    void set_slot_sampling(int slot, const strata::kernels::SamplerParams& sp) {
+        if (slot >= 0 && slot < (int) slot_sp_.size()) slot_sp_[(size_t) slot] = sp;
+        if (next_) next_->set_slot_sampling(slot, sp);
+    }
+    const int32_t* batch_out() const { return b_out_; }
+    bool last_stage() const { return g_ != nullptr && le_ == g_->n_layers; }
     /// commit() returns without waiting for its graph (a single-GPU session sets it): the next window follows it on
     /// the same stream and the drafter reads nothing it writes, so it overlaps the draft. Whoever reads or writes
     /// the session from another stream or the host afterwards (a new request, a checkpoint, a snapshot, the prompt
@@ -169,6 +215,35 @@ public:
 private:
     RemoteExpertOpt* remote_opt_ = nullptr;
     bool capture(int T, std::string& err);
+    // batch windows (see init_slots)
+    std::vector<SessionState*> slots_;
+    bool batch_rec_ = false;               ///< record_window is capturing a batch window
+    int row_base_ = 0;                     ///< ... its hand-off rows start here (a pipeline group's own rows)
+    int brow_[8] = {};                     ///< ... and row t is slot brow_[t]
+    bool last_batch_ = false;              ///< the last run was a batch window (set_plan_slot: one group)
+    std::map<uint64_t, cudaGraphExec_t> exec_bm_, commit_bm_;   ///< key: batch_key(rows, S, hand-off base)
+    int last_rows_[8] = {};                ///< the slots of the last batch window's rows
+    static uint64_t batch_key(const int* rows, int S, int hbase) {
+        uint64_t k = (uint64_t) hbase << 40 | (uint64_t) S << 32;
+        for (int t = 0; t < S; ++t) k |= (uint64_t) (rows[t] & 15) << (4 * t);
+        return k;
+    }
+    // batch_launch / batch_poll
+    bool b_running_ = false;
+    int64_t b_k_ = 0, b_steps_ = 0;
+    std::chrono::steady_clock::time_point b_last_;
+    int32_t b_out_[8] = {};
+    std::vector<strata::kernels::SamplerParams> slot_sp_;
+    bool sample_rows(int S, std::string& err);   ///< the sampled slots' rows of the last batch window
+    int32_t* h_commitb_ = nullptr; int32_t* m_commitb_ = nullptr;   // per slot [1, 0, pos, -1 ..], stride 2 + max_t
+    int32_t* commitb_ = nullptr;
+    float* tail_snap_b_ = nullptr;         ///< per (slot, QSA layer) indexer tail snapshot
+    void* arena_b_ = nullptr;
+    int64_t last_pos_b_[8] = {};
+    bool capture_batch(const int* rows, int S, int hbase, std::string& err);
+    void collect_profile();   ///< STRATA_VERIFY_PROFILE: add the last window's stamps to prof_sum_
+    bool capture_commit_batch(const int* rows, int S, int hbase, std::string& err);
+    bool stage_batch(const int* rows, int S, int hbase, const int32_t* tokens, const int64_t* pos, std::string& err);
     strata::kernels::SamplerParams sampling_ = [] {
         strata::kernels::SamplerParams s;
         s.greedy = true;

serve/server.py

@@ -21,6 +21,7 @@ talks to it over stdin/stdout; `MockEngine` is a scripted stand-in that makes ev
 from __future__ import annotations
 
 import argparse
+import contextlib
 import collections
 import base64
 import hashlib
@@ -376,6 +377,17 @@ def conversation_cache_view(info: dict, hist: list, totals: dict, parked: dict)
             "last_prompt": last.get("prompt_tokens") if last else None}
 
 
+_BTRACE = bool(os.environ.get("STRATA_BATCH_TRACE"))
+
+
+def btrace(*a):
+    if _BTRACE:
+        print("[batch-trace]", threading.get_ident() % 10000, *a, file=sys.stderr, flush=True)
+
+
+EOS_IDS = {248044, 248046}   # <|endoftext|>, <|im_end|>: the engine's default --eos-ids
+
+
 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.
@@ -386,6 +398,25 @@ class StrataEngine:
     """
     silence_s = ENGINE_SILENCE_S         # #481: main() sets the config's engine_silence_s (survives restart())
     silent_note = None                   # #481: why the server ended a silent engine (death_note says it)
+    batch = 0                            # --batch: the engine's batch slots (0: one request at a time)
+
+    # `last`: the figures of the engine's last DONE line.  With batch slots several requests run at once, each on its
+    # own thread, so each one reads its own request's figures (Service.run records them); one at a time, the one dict.
+    @property
+    def last(self):
+        if self.batch:
+            tl = self.__dict__.setdefault("_tl", threading.local())
+            if not hasattr(tl, "last"):
+                tl.last = {}
+            return tl.last
+        return self.__dict__.get("_last", {})
+
+    @last.setter
+    def last(self, value):
+        self.__dict__["_last"] = value
+        if self.batch:
+            tl = self.__dict__.setdefault("_tl", threading.local())
+            tl.last = value
 
     def __init__(self, exe: str, args: list[str], cwd: str | None = None, log: str | None = None,
                  env: dict | None = None, lazy: bool = False):
@@ -454,16 +485,51 @@ class StrataEngine:
         self.ended = False                              # READY: alive from here (restart() set it True, #344)
         # the engine's stdout on a thread, so a request can wait with a timeout (heartbeats, cancel checks)
         self.lines: queue.Queue = queue.Queue()
+        # --batch: concurrent requests in the engine's batch slots (see generate_batched).  The engine says how many
+        # it runs (INFO batch_slots=N: it may have fewer than asked, or none, when they do not fit)
+        asked = next((int(args[args.index(k) + 1]) for k in ("--batch", "--slots") if k in args), 0)
+        self.batch = int(self.info.get("batch_slots") or 0)
+        if asked and self.batch != asked:
+            print(f"[strata] parallel requests: {asked} asked, the engine runs {self.batch or 'one at a time'} "
+                  "(its log says why)", flush=True)
+        groups = int(args[args.index("--batch-groups") + 1]) if "--batch-groups" in args else 1
+        groups = groups if self.batch and groups > 0 and self.batch % groups == 0 else 1
+        gs = self.batch // groups if self.batch else 0
+        # slots in the order that spreads requests over the pipeline's groups first: 0, gs, 2gs, .., 1, gs+1, ..
+        self.slot_order = [g * gs + t for t in range(gs) for g in range(groups)]
+        self.slot_q = [queue.Queue() for _ in range(self.batch)]
+        self.slot_busy = [False] * self.batch
+        # what each slot's sessions hold (prompt + every token a window fed), so a conversation's next turn goes to
+        # the slot that has its start (the engine checks it again); when the slot was last used
+        self.slot_held: list[list[int]] = [[] for _ in range(self.batch)]
+        self.slot_used = [0.0] * self.batch
+        self.slot_live: list[dict | None] = [None] * self.batch   # /metrics: the request in each slot
+        self.slot_cv = threading.Condition()
+        self.waiting = 0                                # requests waiting for the control lines (ctl)
+        self.wait_lens: list[list[int]] = []            # ... their prompt lengths (a long read gives way to short ones)
+        self.ctl_epoch = 0                              # how often the control lines were taken
+        self._yielded = None                            # (slot, tokens read): the last request on them gave way
+        self.ctl = threading.Lock()                     # one admission or solo request on the control lines at a time
+        self.wlock = threading.Lock()                   # stdin writes from several request threads
         self.pump = threading.Thread(target=self._pump, daemon=True)
         self.pump.start()
 
     def _pump(self):
         proc, lines = self.proc, self.lines             # this process's: a restart replaces both (#344)
+        slot_q = self.slot_q
         for line in proc.stdout:
+            if line.startswith(("BT ", "BDONE ")) and slot_q:   # --batch: a batch slot's own lines
+                try:
+                    slot_q[int(line.split()[1])].put(line)
+                    continue
+                except (IndexError, ValueError):
+                    pass
             lines.put(line)
         if self.proc is proc:                           # a killed engine's pump must not mark its successor dead
             self.ended = True                           # its output closed: it is gone, even before the OS says so
         lines.put(None)
+        for q in slot_q:
+            q.put(None)
 
     def death_note(self) -> str:
         """Why the engine most likely ended, from the end of its log: its own watchdog (issue #29), else RAM."""
@@ -639,10 +705,414 @@ class StrataEngine:
         on = sampling.get("experimental_speed_projection")
         return f" cvec={int(on)}" if isinstance(on, bool) else ""
 
+    def _send(self, text: str):
+        btrace("send>", text[:60])
+        try:
+            with self.wlock:
+                self.proc.stdin.write(text + "\n")
+                self.proc.stdin.flush()
+        except OSError:                                  # the pipe is gone: the engine died (not the client)
+            raise EngineDied(f"the engine stopped unexpectedly (exit code {self.exit_code()})") from None
+
+    def _control(self, cancel, on_token, stop_when=None):
+        """Reads the control lines of the request on them (GEN / BGEN), yielding None heartbeats.  Calls on_token(id)
+        for each `T`; returns ("done", None) at DONE, or ("badm", continues) at BADM (after DONE).  `stop_when()`
+        true sends STOP once (the request is then read to its DONE)."""
+        stopped = False
+        while True:
+            try:
+                line = self.lines.get(timeout=10.0)
+            except queue.Empty:
+                if cancel.is_set() and not stopped:
+                    self._send("STOP")
+                    stopped = True
+                yield None
+                continue
+            if line is None:
+                raise EngineDied(f"the engine stopped unexpectedly (exit code {self.exit_code()})")
+            if not line.startswith(("PP ", "INFO")):
+                btrace("ctl<", self._ctl_mode, line.strip()[:60])
+            if line.startswith("T "):
+                on_token(int(line[2:]))
+                if not stopped and (cancel.is_set() or (stop_when is not None and stop_when())):
+                    self._send("STOP")
+                    stopped = True
+                yield False                              # a token is pending (not a heartbeat)
+            elif line.startswith("PP "):
+                f = line.split()
+                if len(f) >= 3 and f[1].isdigit() and f[2].isdigit():
+                    self.progress = (int(f[1]), int(f[2]))
+                    self.prefill_tok_s_mean = float(f[4]) if len(f) >= 5 else None
+                yield None
+            elif line.startswith("DONE"):
+                self._parse_done(line)
+                self._last_done = line
+            elif line.startswith("BADM "):
+                f = line.split()
+                self._ctl_result = ("badm", len(f) >= 3 and f[2] == "1")
+                return
+            elif line.startswith("YIELDED "):            # the read gave way (BYIELD): <slot> <tokens read>
+                f = line.split()
+                if len(f) >= 3 and f[1].lstrip("-").isdigit() and f[2].isdigit():
+                    self._yielded = (int(f[1]), int(f[2]))
+            elif line.startswith("ERR"):
+                raise ValueError(line[4:].strip())
+            if line.startswith("DONE") and self._ctl_mode == "solo":
+                self._ctl_result = ("done", None)
+                return
+
+    def _drain_control(self, until: str, timeout: float = 300.0):
+        """After a consumer left early: read the control lines up to the next `until` line (DONE or BADM) so the next
+        request does not read this one's leftovers.  Returns that line (None: the engine ended or never answered)."""
+        end = time.monotonic() + timeout
+        while time.monotonic() < end:
+            try:
+                line = self.lines.get(timeout=max(0.1, end - time.monotonic()))
+            except queue.Empty:
+                break
+            if line is None:
+                return None
+            if line.startswith("DONE"):
+                self._parse_done(line)
+            if line.startswith(until) or line.startswith("ERR"):
+                return line
+        return None
+
+    def _release_slot_when_done(self, slot: int, stream: list[int] | None = None):
+        """A slot whose consumer left (a stop token, a stop string, a disconnect): BSTOP it and free it once the engine
+        says BDONE (in the background).  `stream`: the prompt and every token of it so far - with the tokens still to
+        come before BDONE, all but the last are what the slot holds then (the next turn of its conversation)."""
+        try:
+            self._send(f"BSTOP {slot}")
+        except EngineDied:
+            pass
+        def wait():
+            end = time.monotonic() + 600.0
+            tail = list(stream or [])
+            while time.monotonic() < end:
+                try:
+                    line = self.slot_q[slot].get(timeout=5.0)
+                except queue.Empty:
+                    continue
+                if line is None:
+                    tail = []
+                    break
+                if line.startswith("BT "):
+                    try:
+                        tail.append(int(line.split()[2]))
+                    except (IndexError, ValueError):
+                        tail = []
+                if line.startswith("BDONE "):
+                    break
+            self.slot_held[slot] = tail[:-1] if stream and tail else []
+            with self.slot_cv:
+                self.slot_busy[slot] = False
+                self.slot_cv.notify_all()
+        threading.Thread(target=wait, daemon=True).start()
+
+    YIELDS_MAX = 2   # #656: how often one request's prompt read gives way to a shorter waiting one
+
+    def _take_control(self, cancel, plen: int, after_epoch: int | None = None):
+        """Waits for the control lines (one prompt read at a time), yielding None heartbeats; False when cancelled.
+        `after_epoch`: a request whose read gave way lets the requests waiting then go first."""
+        entry = [plen]
+        with self.slot_cv:
+            self.waiting += 1
+            self.wait_lens.append(entry)
+        beat = time.monotonic()
+        try:
+            while True:
+                with self.slot_cv:
+                    turn = after_epoch is None or self.ctl_epoch > after_epoch or self.waiting <= 1
+                if turn and self.ctl.acquire(timeout=0.5):
+                    break
+                if not turn:
+                    with self.slot_cv:
+                        self.slot_cv.wait(timeout=0.5)
+                if cancel.is_set():
+                    return False
+                if time.monotonic() - beat >= 10.0:     # a heartbeat every 10 s, as the engine's own wait
+                    beat = time.monotonic()
+                    yield None
+        finally:
+            with self.slot_cv:
+                self.waiting -= 1
+                self.wait_lens.remove(entry)
+        with self.slot_cv:
+            self.ctl_epoch += 1
+            self.slot_cv.notify_all()
+        return True
+
+    SOLO_AGAIN_MAX = 2      # how often a request left alone in a slot goes back to the solo path
+    SOLO_AGAIN_MIN_LEFT = 32  # ... only with at least this many tokens still allowed (max_tokens)
+
+    def _may_go_solo(self, left: int, times: int, embeddings) -> bool:
+        """A request decoding in a slot that is alone now (no other slot busy, nobody waiting) goes back to the solo
+        path with its MTP drafts: the engine continues it from the slot's sessions (its slot cache; INFO
+        slot_cache=1).  STRATA_PARALLEL_SOLO=0 keeps it in the slot."""
+        if (embeddings or times >= self.SOLO_AGAIN_MAX or left < self.SOLO_AGAIN_MIN_LEFT or
+                not (self.info or {}).get("slot_cache") or os.environ.get("STRATA_PARALLEL_SOLO") == "0"):
+            return False
+        with self.slot_cv:
+            return sum(1 for b in self.slot_busy if b) == 1 and self.waiting == 0
+
+    def _shorter_waiting(self, plen: int) -> bool:
+        """#656: someone waits for the control lines with a prompt under half this one's: worth giving way to."""
+        with self.slot_cv:
+            return any(e[0] * 2 <= plen for e in self.wait_lens)
+
+    def generate_batched(self, ids, max_new, sampling, cancel, embeddings=None):
+        """--batch.  Alone (no slot busy, nobody waiting): the solo path (GEN, with drafts: the fastest stream)
+        - and when another request arrives meanwhile, this one is STOPped and continues in a batch slot (BGEN with
+        its prompt + what it generated: the engine reuses that prefix).  Otherwise: BGEN into a free slot, then the
+        slot's own BT lines until BDONE.  Several requests run at once; the control lines (prompt reading, admission)
+        are taken one request at a time - and a long prompt read gives way at a chunk boundary to a waiting request
+        with a much shorter prompt (#656: `BYIELD <slot>`; the part read waits in a slot and the read goes on after).
+        A consumer that stops early leaves the engine in step: the solo request is STOPped and read to its DONE, an
+        admission to its BADM, a slot is BSTOPped and freed at its BDONE."""
+        self.progress = None
+        keys = self.sampling_keys(sampling or {})
+        out: list[int] = []
+        pending: list[int] = []
+        prompt, left = list(ids), int(max_new)
+        ok = yield from self._take_control(cancel, len(prompt))
+        if not ok:
+            return
+        holding = True
+        btrace("ctl acquired")
+        slot, gen0, reserved = None, None, None
+        phase = "none"            # solo -> (admit -> slot) ; "done" once the engine has finished with this request
+        stop_sent = False
+        yields, solo_again = 0, 0
+        try:
+            while True:   # a request in a slot that is left alone goes back to the solo path
+                if not holding:
+                    ok = yield from self._take_control(cancel, len(prompt))
+                    if not ok:
+                        return
+                    holding = True
+                with self.slot_cv:
+                    alone = not any(self.slot_busy) and self.waiting == 0
+                if alone and left > 1:
+                    head = f"GENI {left}{keys} {embeddings}" if embeddings else f"GEN {left}{keys}"
+                    self._send(f"{head} {','.join(str(int(t)) for t in prompt)}")
+                    phase = "solo"
+                    self._ctl_mode, self._ctl_result, self._yielded = "solo", None, None
+                    def others():
+                        with self.slot_cv:
+                            return self.waiting > 0
+                    for x in self._control(cancel, pending.append, stop_when=others):
+                        while pending:
+                            t = pending.pop(0)
+                            out.append(t)
+                            yield t
+                        if x is None:                       # a heartbeat (False: a token, flushed above)
+                            if (reserved is None and not out and not embeddings and yields < self.YIELDS_MAX and
+                                    self._shorter_waiting(len(prompt))):
+                                with self.slot_cv:          # the slot the part read will wait in
+                                    reserved = self.pick_slot(prompt)
+                                    if reserved is not None:
+                                        self.slot_busy[reserved] = True
+                                if reserved is not None:
+                                    self._send(f"BYIELD {reserved}")
+                            yield None
+                    phase = "none"                          # its DONE is read
+                    while pending:
+                        t = pending.pop(0)
+                        out.append(t)
+                        yield t
+                    if self._yielded is not None and reserved is not None and self._yielded[0] == reserved:
+                        slot, reserved = reserved, None     # gave way: the read goes on in that slot (below)
+                    elif reserved is not None:              # it did not give way: the slot is free again
+                        with self.slot_cv:
+                            self.slot_busy[reserved] = False
+                            self.slot_cv.notify_all()
+                        reserved = None
+                    if slot is None:
+                        finish = (self.last or {}).get("finish") if isinstance(self.last, dict) else None
+                        left = int(max_new) - len(out)
+                        if cancel.is_set() or left <= 0 or finish in ("stop", "length") or (out and out[-1] in EOS_IDS):
+                            return
+                        prompt = list(ids) + out            # promoted: it continues in a batch slot from here
+                while True:
+                    if slot is None:
+                        # a free slot (they free themselves at BDONE, which needs no control lines): the one that holds
+                        # the start of this prompt (its conversation's last turn), else the one used longest ago
+                        with self.slot_cv:
+                            while True:
+                                slot = self.pick_slot(prompt)
+                                if slot is not None:
+                                    self.slot_busy[slot] = True
+                                    break
+                                self.slot_cv.wait(timeout=10.0)
+                                if cancel.is_set():
+                                    return
+                    if self._yielded is not None:           # it gave way: the others waiting then go first
+                        self.slot_held[slot] = list(prompt[:self._yielded[1]])
+                        self._yielded = None
+                        yields += 1
+                        self.ctl.release()
+                        holding = False
+                        with self.slot_cv:
+                            epoch = self.ctl_epoch
+                        ok = yield from self._take_control(cancel, len(prompt), after_epoch=epoch)
+                        if not ok:
+                            return
+                        holding = True
+                    while not self.slot_q[slot].empty():
+                        self.slot_q[slot].get_nowait()
+                    head = f"BGENI {slot} {left}{keys} {embeddings}" if embeddings else f"BGEN {slot} {left}{keys}"
+                    live = {"slot": slot, "state": "reading", "prompt_tokens": len(prompt), "generated": 0,
+                            "started": time.time(), "first_token": None}
+                    self.slot_live[slot] = live
+                    self._send(f"{head} {','.join(str(int(t)) for t in prompt)}")
+                    self.slot_held[slot] = []               # the admission overwrites what the slot held
+                    phase = "admit"
+                    self._ctl_mode, self._ctl_result, self._yielded = "batch", None, None
+                    asked = False
+                    for x in self._control(cancel, pending.append):
+                        while pending:
+                            t = pending.pop(0)
+                            out.append(t)
+                            yield t
+                        if x is None:
+                            if (not asked and not embeddings and yields < self.YIELDS_MAX and
+                                    self._shorter_waiting(len(prompt))):
+                                self._send(f"BYIELD {slot}")
+                                asked = True
+                            yield None
+                    cont = bool(self._ctl_result and self._ctl_result[1])
+                    phase = "slot" if cont else "none"
+                    while pending:
+                        t = pending.pop(0)
+                        out.append(t)
+                        yield t
+                    if not cont and self._yielded is not None and self._yielded[0] == slot and not cancel.is_set():
+                        continue                            # gave way: again once the shorter request is in
+                    break
+                self.ctl.release()
+                holding = False
+                if not cont:
+                    return
+                live.update(state="decoding", first_token=time.time(), generated=len(out))
+                gen0 = len(out) - 1                         # the admission's own token: the slot feeds it first
+                going_solo = False
+                while True:
+                    try:
+                        line = self.slot_q[slot].get(timeout=10.0)
+                    except queue.Empty:
+                        if cancel.is_set() and not stop_sent:
+                            self._send(f"BSTOP {slot}")
+                            stop_sent = True
+                        yield None
+                        continue
+                    if line is None:
+                        phase = "none"
+                        raise EngineDied(f"the engine stopped unexpectedly (exit code {self.exit_code()})")
+                    if line.startswith("BT "):
+                        t = int(line.split()[2])
+                        out.append(t)
+                        live["generated"] = len(out)
+                        if cancel.is_set():
+                            if not stop_sent:
+                                self._send(f"BSTOP {slot}")
+                                stop_sent = True
+                            continue
+                        yield t
+                        if not going_solo and not stop_sent and self._may_go_solo(int(max_new) - len(out), solo_again,
+                                                                                  embeddings):
+                            self._send(f"BSTOP {slot}")         # alone now: on with the drafts (below)
+                            going_solo = True
+                    elif line.startswith("BDONE "):
+                        phase = "none"
+                        f = line.split()
+                        if len(f) >= 5 and isinstance(self.last, dict):
+                            self.last = {**self.last, "finish": f[3], "decode_ms": float(f[4]),
+                                         "generated": int(f[2]) if f[2].isdigit() else self.last.get("generated")}
+                        # what the slot's sessions hold now: the prompt and every token fed (all but the last one)
+                        self.slot_held[slot] = list(prompt) + out[gen0:-1] if len(out) > gen0 else []
+                        if (going_solo and f[3:4] == ["cancel"] and not cancel.is_set() and len(out) < int(max_new)
+                                and not (out and out[-1] in EOS_IDS)):
+                            # the solo path continues it: the engine copies the slot's sessions back (all but the
+                            # last token are in them) and decodes with MTP drafts again
+                            self.slot_live[slot] = None
+                            self.slot_used[slot] = time.time()
+                            with self.slot_cv:
+                                self.slot_busy[slot] = False
+                                self.slot_cv.notify_all()
+                            slot, gen0 = None, None
+                            prompt, left = list(ids) + out, int(max_new) - len(out)
+                            solo_again += 1
+                            btrace("back to the solo path")
+                            break
+                        return
+        finally:
+            # a consumer that left early (or an error): keep the engine and this server in step
+            btrace("finally phase", phase, "slot", slot, "holding", holding)
+            try:
+                if phase == "solo":
+                    self._send("STOP")
+                    self._drain_control("DONE")
+                elif phase == "admit":
+                    line = self._drain_control("BADM")
+                    if line and line.startswith("BADM ") and line.split()[2:3] == ["1"]:
+                        phase = "slot"
+            except EngineDied:
+                pass
+            if holding:
+                self.ctl.release()
+            if reserved is not None:
+                with self.slot_cv:
+                    self.slot_busy[reserved] = False
+                    self.slot_cv.notify_all()
+            if slot is not None:
+                self.slot_live[slot] = None
+                self.slot_used[slot] = time.time()
+                if phase == "slot":
+                    self.slot_held[slot] = []
+                    stream = list(prompt) + out[gen0:] if gen0 is not None and len(out) > gen0 else None
+                    self._release_slot_when_done(slot, stream)    # freed at its BDONE
+                else:
+                    with self.slot_cv:
+                        self.slot_busy[slot] = False
+                        self.slot_cv.notify_all()
+
+    def pick_slot(self, prompt: list[int]) -> int | None:
+        """A free slot for `prompt` (the caller holds slot_cv): the one whose held tokens are the longest start of the
+        prompt (the engine then reads only the rest), else an empty one, else the one used longest ago - so the
+        conversations other slots hold stay for their next turns.  None: every slot is busy."""
+        free = [b for b in self.slot_order if not self.slot_busy[b]]
+        if not free:
+            return None
+        def held_prefix(b):
+            h = self.slot_held[b]
+            return len(h) if h and len(h) < len(prompt) and prompt[:len(h)] == h else 0
+        best = max(free, key=held_prefix)
+        if held_prefix(best) > 0:
+            return best
+        return min(free, key=lambda b: (bool(self.slot_held[b]), self.slot_used[b]))
+
+    def slots_view(self) -> list[dict]:
+        """/metrics: each batch slot - idle (with the tokens it holds for a next turn), reading or decoding."""
+        now, view = time.time(), []
+        for b in range(self.batch):
+            r = self.slot_live[b]
+            if r is None:
+                view.append({"slot": b, "state": "idle", "held_tokens": len(self.slot_held[b])})
+                continue
+            ft = r.get("first_token")
+            view.append({"slot": b, "state": r["state"], "prompt_tokens": r["prompt_tokens"],
+                         "generated": r["generated"], "elapsed_s": round(now - r["started"], 1),
+                         "tok_s": round(r["generated"] / max(1e-6, now - ft), 1) if ft else None})
+        return view
+
     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."""
+        if getattr(self, "batch", 0):
+            yield from self.generate_batched(ids, max_new, sampling, cancel, embeddings)
+            return
         self.progress = None
         self.prefill_tok_s_mean = None
         # an image request takes the same sampling keys as text (#75: it used to decode greedily whatever was asked)
@@ -1005,9 +1475,34 @@ def engine_args(cfg: dict) -> list[str]:
         seg = cfg.get("vram_segment_mib")
         if isinstance(seg, int) and not isinstance(seg, bool) and seg > 0 and "--vram-segment-mib" not in args:
             args += ["--vram-segment-mib", str(seg)]
+    args += parallel_args(cfg, args)
     return learned_profile_args(cfg, args)
 
 
+PARALLEL_MAX = 8      # the engine's batch window holds at most 8 rows (kVerifyMaxT)
+
+
+def parallel_args(cfg: dict, args: list[str]) -> list[str]:
+    """#465 (opt-in): "parallel": N in the config - up to N requests decode together in the engine's batch slots
+    (--batch N); more wait their turn.  1 or absent: one request at a time, as before.  A value the engine cannot use
+    is said and passed on: the engine warns and adjusts (recommend, never force)."""
+    if "--batch" in args or "--slots" in args:
+        return []
+    n = cfg.get("parallel")
+    if n is None or n is False:
+        return []
+    if isinstance(n, bool) or not isinstance(n, int):
+        print(f'[strata] "parallel" must be a whole number of requests (2..{PARALLEL_MAX}), not {n!r}: ignored',
+              flush=True)
+        return []
+    if n <= 1:
+        return []
+    if n > PARALLEL_MAX:
+        print(f'[strata] "parallel": {n} - the engine runs at most {PARALLEL_MAX} requests together; it will use '
+              f"{PARALLEL_MAX}", flush=True)
+    return ["--batch", str(n)]
+
+
 def profile_shape(path: str) -> tuple[int, int] | None:
     """An expert profile's (layers, experts per layer), from its header (tools/make_profile.py's format), or None
     when the file is missing, is not one or is shorter than the pairs its header promises."""
@@ -1183,6 +1678,9 @@ class Service:
         self.host_names: set[str] = set(LOOPBACK_NAMES)
         self.status = {"busy": False, "queued": 0}      # GET /status: what the model is doing right now
         self.rate = collections.deque(maxlen=32)        # (time, generated) samples for the live tok/s window
+        # "parallel" (the engine's batch slots): every running request's own status and rate window; self.status
+        # then says busy while any runs and shows the newest one
+        self.live_reqs: dict[int, tuple[dict, collections.deque]] = {}
         # #332: the API request monitor (/api-monitor) keeps the last 100 requests' prompts and answers in memory,
         # so it is off unless the config's "api_monitor" (or --api-monitor) turns it on
         self.api_monitor = False
@@ -1323,6 +1821,12 @@ class Service:
             raise ValueError("this engine cannot resize its VRAM use")
         if not self.fifo.acquire(timeout=self.vram_wait_s):
             raise ModelBusy("a request is still running; try again when it has finished")
+        # parallel requests do not hold the fifo: the engine's control lines (a prompt being admitted) are taken too,
+        # and the engine refuses a resize while a slot decodes
+        ctl = getattr(self.engine, "ctl", None) if getattr(self.engine, "batch", 0) else None
+        if ctl is not None and not ctl.acquire(timeout=self.vram_wait_s):
+            self.fifo.release()
+            raise ModelBusy("a request is still running; try again when it has finished")
         try:
             if not self.loaded():
                 self.vram_reserve = reserve_mib
@@ -1334,6 +1838,8 @@ class Service:
                   f"{out.get('expert_cache_full_mib')} MiB ({out.get('expert_slots')} experts)", flush=True)
             return {"status": "ok", **out}
         finally:
+            if ctl is not None:
+                ctl.release()
             self.fifo.release()
 
     def _say_died(self, e: Exception) -> None:
@@ -1449,10 +1955,19 @@ class Service:
                                        amd=getattr(self, "backend", None) == "hip")
 
     def _tok_s(self):
-        """tok/s over the last RATE_WINDOW_S seconds.  Returns 0.0 while nothing is generating."""
+        """tok/s over the last RATE_WINDOW_S seconds.  Returns 0.0 while nothing is generating.  With parallel requests:
+        the sum of theirs."""
+        with self.status_lock:
+            many = [(dict(st), list(rt)) for st, rt in self.live_reqs.values()]
+        if many:
+            return sum(self._rate_of(st, rt) for st, rt in many)
         with self.status_lock:
             s = dict(self.status)
             rate = list(self.rate)
+        return self._rate_of(s, rate)
+
+    @staticmethod
+    def _rate_of(s, rate):
         if not s.get("busy") or not s.get("first_token"):
             return 0.0
         now = time.time()
@@ -1465,10 +1980,9 @@ class Service:
     def _tok_s_mean(self):
         """The whole-request mean since the first token (the old formula), kept so the two can be compared."""
         with self.status_lock:
-            s = dict(self.status)
-        if not s.get("busy") or not s.get("first_token"):
-            return 0.0
-        return s["generated"] / max(1e-6, time.time() - s["first_token"])
+            many = [dict(st) for st, _ in self.live_reqs.values()] or [dict(self.status)]
+        return sum(s["generated"] / max(1e-6, time.time() - s["first_token"]) for s in many
+                   if s.get("busy") and s.get("first_token"))
 
     def _prefill_tok_s_mean(self):
         """Engine-reported mean over newly read tokens, excluding the cached prefix."""
@@ -1512,6 +2026,8 @@ class Service:
         the hardware (with a minute of history per series)."""
         with self.status_lock:
             s = dict(self.status)
+            if self.live_reqs:                          # parallel requests: the newest one's status
+                s = {**s, **dict(list(self.live_reqs.values())[-1][0]), "queued": s.get("queued", 0)}
             hist = list(self.history)
             totals = dict(self.totals)
         now = time.time()
@@ -1535,6 +2051,17 @@ class Service:
                 "tok_s_window_s": RATE_WINDOW_S if state == "generating" else None}
         if state == "reading" and progress:
             live["prompt_read"], live["prompt_total"] = progress
+        par = int(getattr(self.engine, "batch", 0) or 0)
+        if par:                                         # #465: the requests running together, slot by slot
+            with self.status_lock:
+                running = len(self.live_reqs)
+            slots = self.engine.slots_view() if hasattr(self.engine, "slots_view") else []
+            # the requests in flight that are not in a slot: the one alone on the solo path, or waiting their turn
+            in_slots = sum(1 for x in slots if x["state"] != "idle")
+            live.update(parallel=par, running=running, slots=slots, outside_slots=max(0, running - in_slots),
+                        waiting=int(getattr(self.engine, "waiting", 0) or 0))
+            if running and state == "idle":
+                live["state"] = "generating"
         engine = {"model": self.model, "max_context": self.engine.max_context, "images": self.vision is not None,
                   **dict(getattr(self.engine, "info", {}) or {})}
         tel = self.telemetry.snapshot() if getattr(self, "telemetry", None) else {"now": {}, "history": {}, "static": {}}
@@ -1569,7 +2096,9 @@ class Service:
             "started": int(self.started_at), "uptime_s": int(time.time() - self.started_at),
             "cache_max_tokens": ctx,
             "context": {"native": ctx, "max_positions": ctx},
-            "concurrency": {"serving": 1, "requested": 1},       # one request at a time; more wait their turn
+            # one request at a time (more wait their turn), or "parallel": N batch slots (#465)
+            "concurrency": {"serving": max(1, int(getattr(self.engine, "batch", 0) or 0)),
+                            "requested": max(1, int(getattr(self.engine, "batch", 0) or 0))},
             "dialects": ["/v1/chat/completions", "/v1/messages", "/v1/responses"],
             "vision": {"enabled": images, "available": images, "error": None},
             "activity": {"requests": totals["requests"] + int(busy), "in_flight": int(busy) + int(s.get("queued") or 0),
@@ -1676,13 +2205,13 @@ class Service:
             max_new = max(1, room)          # --fit-max-tokens: a shorter completion beats a 400
         return ids, kwargs.get("enable_thinking", True) is not False, max_new
 
-    def _note(self, n, evs):
+    def _note(self, n, evs, st=None, rate=None):
         with self.status_lock:
-            s = self.status
+            s = self.status if st is None else st
             s["generated"] = n
             if s.get("first_token") is None:
                 s["first_token"] = time.time()
-            self.rate.append((time.time(), n))          # the live rate's window over the last RATE_WINDOW_S
+            (self.rate if rate is None else rate).append((time.time(), n))   # the live rate's window (RATE_WINDOW_S)
             for ev in evs:
                 if ev.kind == "reasoning":
                     s["phase"] = "thinking"
@@ -1694,13 +2223,13 @@ class Service:
                     s["phase"] = "tool call complete"
                 s["tail"] = ((s.get("tail") or "") + (ev.text or ""))[-600:]
 
-    def _progress(self, last_print, every=1.0):
+    def _progress(self, last_print, every=1.0, st=None):
         """A progress line in the server window every `every` seconds while a request runs."""
         now = time.time()
         if now - last_print < every:
             return last_print
         with self.status_lock:
-            s = dict(self.status)
+            s = dict(self.status if st is None else st)
         el = now - s.get("started", now)
         if s.get("first_token") is None:
             pr = getattr(self.engine, "progress", None)   # (position reached, prompt tokens): a reused prefix counts
@@ -1731,10 +2260,16 @@ class Service:
         engine_last0 = getattr(self.engine, "last", None)
         trace = getattr(self.request_trace, "record", None)
         waiting = time.perf_counter()
+        # #465: with "parallel" (the engine's batch slots) requests run at once: each keeps its own status and rate
+        # window (self.status says busy while any runs); one at a time they are self.status / self.rate, as before
+        par = bool(getattr(self.engine, "batch", 0))
+        st = {} if par else self.status
+        rate = collections.deque(maxlen=32) if par else self.rate
         with self.status_lock:
             self.status["queued"] += 1
         try:
-            with self.fifo:
+            # --batch: the engine runs several requests at once (StrataEngine.generate_batched orders them)
+            with (contextlib.nullcontext() if getattr(self.engine, "batch", 0) else self.fifo):
                 try:
                     with self.status_lock:
                         if trace is not None:
@@ -1742,13 +2277,18 @@ class Service:
                             trace["state"] = "generating"
                         self.status["queued"] -= 1
                     # issue #27: it died in an earlier request (or was unloaded) - start it again instead of failing
-                    self.ensure_loaded()
+                    # (parallel requests do not hold the fifo: one of them starts it, the others wait for that)
+                    with (self.fifo if par else contextlib.nullcontext()):
+                        self.ensure_loaded()
                     with self.status_lock:
-                        self.status.update(busy=True, phase="reading the prompt", prompt_tokens=len(ids),
-                                           generated=0, started=time.time(), first_token=None, tool=None, tail="",
-                                           max_tokens=max_new)
+                        st.update(busy=True, phase="reading the prompt", prompt_tokens=len(ids),
+                                  generated=0, started=time.time(), first_token=None, tool=None, tail="",
+                                  max_tokens=max_new)
+                        if par:
+                            self.live_reqs[id(st)] = (st, rate)
+                            self.status.update(st)
                         self.last_request_at = time.time()
-                        self.rate.clear()               # the previous request's samples must not leak into this one
+                        rate.clear()                    # the previous request's samples must not leak into this one
                     before = getattr(self.engine, "last", None)
                     last_print = time.time()
                     prompt, thought = ids, 0            # thought: the reasoning tokens so far (the budget's count)
@@ -1759,7 +2299,7 @@ class Service:
                         try:
                             for t in gen:
                                 if t is None:               # heartbeat while the engine is quiet
-                                    last_print = self._progress(last_print)
+                                    last_print = self._progress(last_print, st=st)
                                     yield "ping", None
                                     continue
                                 n += 1
@@ -1779,8 +2319,8 @@ class Service:
                                     repeated = True     # #606: a degenerate output, not an answer: end it here
                                     break
                                 evs = parser.feed(detok.push(t))
-                                self._note(n, evs)
-                                last_print = self._progress(last_print)
+                                self._note(n, evs, st, rate)
+                                last_print = self._progress(last_print, st=st)
                                 for ev in evs:
                                     yield "event", ev
                                 if budget and parser.state == "reasoning":
@@ -1826,7 +2366,7 @@ class Service:
                             raw_ids.append(t)
                             thinking_n += parser.state == "reasoning"
                             evs = parser.feed(detok.push(t))
-                            self._note(n, evs)
+                            self._note(n, evs, st, rate)
                             for ev in evs:
                                 yield "event", ev
                         prompt = prompt + seg + extra
@@ -1844,12 +2384,12 @@ class Service:
                     # #266: settle this request's status, history and totals while still holding the fifo: once
                     # it is released the next request sets its own status, which this must not record or clear
                     with self.status_lock:
-                        if self.status.get("busy"):
+                        if st.get("busy"):
                             # only this request's DONE counts: same object means no DONE arrived (death, error,
                             # disconnect)
                             last = dict(getattr(self.engine, "last", {}) or {}) \
                                 if getattr(self.engine, "last", None) is not engine_last0 else {}
-                            started = self.status.get("started", time.time())
+                            started = st.get("started", time.time())
                             cvec = (getattr(self.engine, "info", {}) or {}).get("cvec", 0)
                             loaded = str(cvec) not in ("0", "", "None")
                             hit_rate = round(last["hits"] / last["lookups"], 3) if last.get("lookups") else None
@@ -1890,8 +2430,8 @@ class Service:
                                 self.last_timings = dict(timings, at=int(time.time())) if timings else None
                             self.last_request_at = time.time()
                             now = time.time()
-                            el = now - self.status.get("started", now)
-                            ft = self.status.get("first_token")
+                            el = now - st.get("started", now)
+                            ft = st.get("first_token")
                             rate = n / max(1e-6, now - ft) if ft else 0.0
                             hit_msg = f", expert cache {hit_rate*100:.1f}% hit" if hit_rate is not None else ""
                             if hit_msg and pcie_share:
@@ -1904,9 +2444,17 @@ class Service:
                                       "strata-<model>.json for every request) leaves room to answer", 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
-                        self.status.pop("tail", None)            # #212: the answer's end is not kept once it is done
-                        self.status.pop("tool", None)
+                        st["busy"] = False
+                        st.pop("tail", None)                     # #212: the answer's end is not kept once it is done
+                        st.pop("tool", None)
+                        if par:
+                            self.live_reqs.pop(id(st), None)
+                            if self.live_reqs:              # the newest request still running
+                                self.status.update(list(self.live_reqs.values())[-1][0])
+                            else:
+                                self.status.update(busy=False)
+                                self.status.pop("tail", None)
+                                self.status.pop("tool", None)
         finally:
             if emb:
                 Path(emb).unlink(missing_ok=True)

serve/test_parallel.py

@@ -0,0 +1,382 @@
+"""#465: "parallel": N - several requests decode together in the engine's batch slots (--batch N).
+
+A fake engine (a Python script speaking the engine's GEN / BGEN / BT / BDONE / BADM / BSTOP lines) runs behind the
+real StrataEngine and Service, so the server's side is tested without a GPU: requests up to N run at once, more
+wait for a slot, /metrics shows the slots, each request's history row is its own, and a conversation's next turn
+goes back to the slot that holds it."""
+import json
+import sys
+import tempfile
+import threading
+import time
+import unittest
+import urllib.request
+from pathlib import Path
+from unittest import mock
+
+from serve.frontend import ChatTemplate
+from serve.server import ByteTokenizer, Service, StrataEngine, engine_args, parallel_args, serve
+
+# The fake engine: one token every STEP seconds per active slot (a "window" serves every active slot at once);
+# GEN (solo) streams T lines; BGEN reads the prompt (one T line, DONE), answers BADM and continues in the slot.
+# Every reply is the bytes of "ok, done." then <|im_end|> (ByteTokenizer: 257).
+FAKE_BATCH = r'''import queue, sys, threading, time
+args = sys.argv[1:]
+slots = int(args[args.index("--batch") + 1]) if "--batch" in args else 0
+fit = int(args[args.index("--fit") + 1]) if "--fit" in args else slots
+STEP = 0.02
+CH = 32
+lines, stop = queue.Queue(), threading.Event()
+def reader():
+    for l in sys.stdin:
+        l = l.strip()
+        if l == "STOP":
+            stop.set()
+        else:
+            lines.put(l)
+    lines.put(None)
+threading.Thread(target=reader, daemon=True).start()
+print("INFO engine=0.1.39" + (f" batch_slots={fit}" if fit >= 2 else "") +
+      (" slot_cache=1" if "--slotcache" in args else ""), flush=True)
+print("READY 4096 stop", flush=True)
+LONG = list(b"LONGREPLY")
+def reply(ids):            # the rest of the reply after what the prompt already ends with (a request continued)
+    ids = [int(x) for x in ids]
+    long_one = any(ids[i:i + len(LONG)] == LONG for i in range(len(ids)))
+    R = list(b"ok, " + b"la " * 15 + b"done." if long_one else b"ok, done.") + [257]
+    k = max(k for k in range(len(R)) if k == 0 or ids[-k:] == R[:k])
+    return R[k:]
+active = {}          # slot -> [tokens left, max_new, produced]
+stopped = set()      # BSTOPped slots: they end "cancel"
+def window():        # one batch window: every active slot one token
+    for b in sorted(active):
+        left, max_new, produced = active[b]
+        t = left.pop(0) if left else 257
+        produced += 1
+        print(f"BT {b} {t}", flush=True)
+        if t == 257 or produced >= max_new:
+            fin = 'stop' if t == 257 else 'cancel' if b in stopped else 'length'
+            stopped.discard(b)
+            print(f"BDONE {b} {produced} {fin} 1.0", flush=True)
+            del active[b]
+        else:
+            active[b] = [left, max_new, produced]
+log = open(args[args.index("--log") + 1], "a") if "--log" in args else None
+while True:
+    try:
+        line = lines.get(timeout=STEP if active else None)
+    except queue.Empty:
+        line = ""
+    if line is None or line == "QUIT":
+        break
+    if line.startswith("BYIELD "):
+        continue                              # too late: the read it was for has ended
+    if line.startswith("BSTOP "):
+        b = int(line.split()[1])
+        if b in active:
+            active[b][1] = 0
+            stopped.add(b)
+        continue
+    if line.startswith(("GEN ", "BGEN ")):
+        f = line.split()
+        slot = int(f[1]) if f[0] == "BGEN" else None
+        max_new = int(f[2] if f[0] == "BGEN" else f[1])
+        ids = f[-1].split(",")
+        if log:
+            log.write(f"{f[0]} {slot} {len(ids)} {max(len(active), 0)}\n"); log.flush()
+        toks = reply(ids)
+        stop.clear()
+        # the prompt read: CH tokens a chunk, a PP line each; a BYIELD <s> gives way at a chunk boundary
+        given = None
+        for q in range(CH, len(ids) - 1, CH):
+            time.sleep(STEP / 2)
+            print(f"PP {q} {len(ids)} 1 1", flush=True)
+            held = []
+            while True:
+                try:
+                    l2 = lines.get_nowait()
+                except queue.Empty:
+                    break
+                if l2 and l2.startswith("BYIELD "):
+                    ys = int(l2.split()[1])
+                    if len(ids) - 1 - q > CH and (slot is None or ys == slot) and given is None:
+                        given = (ys, q)
+                else:
+                    held.append(l2)
+            for l2 in held:
+                lines.put(l2)
+            if given:
+                break
+            if active and (q // CH) % 4 == 0:     # the slots decode between the prompt's chunks (as the engine)
+                window()
+        if given:
+            if log:
+                log.write(f"YIELD {given[0]} {given[1]}\n"); log.flush()
+            print(f"YIELDED {given[0]} {given[1]}", flush=True)
+            print(f"DONE 0 {len(ids)} 5.0 0.0 cancel 0 0 0", flush=True)
+            if slot is not None:
+                print(f"BADM {slot} 0", flush=True)
+            continue
+        n = 1 if slot is not None else max_new
+        out = 0
+        for t in toks[:n]:
+            if stop.is_set():
+                break
+            print(f"T {t}", flush=True)
+            out += 1
+            time.sleep(STEP)
+        fin = "stop" if out and toks[out - 1] == 257 else ("cancel" if stop.is_set() else "length")
+        print(f"DONE {out} {len(ids)} 5.0 {out * STEP * 1000:.1f} {fin} 0 0 0", flush=True)
+        if slot is not None:
+            cont = fin == "length" and max_new > 1
+            if cont:
+                active[slot] = [toks[1:], max_new, 1]
+            print(f"BADM {slot} {1 if cont else 0}", flush=True)
+        continue
+    if line:
+        print("ERR unknown " + line[:20], flush=True)
+        continue
+    window()
+'''
+
+
+class ParallelArgs(unittest.TestCase):
+    def test_parallel_config_key(self):
+        self.assertEqual(parallel_args({"parallel": 4}, []), ["--batch", "4"])
+        self.assertEqual(parallel_args({"parallel": 1}, []), [])
+        self.assertEqual(parallel_args({}, []), [])
+        self.assertEqual(parallel_args({"parallel": 2}, ["--batch", "3"]), [])       # the args' own value wins
+        self.assertEqual(parallel_args({"parallel": "4"}, []), [])                   # said and ignored
+        self.assertEqual(parallel_args({"parallel": True}, []), [])
+        self.assertEqual(parallel_args({"parallel": 12}, []), ["--batch", "12"])     # the engine warns and caps
+        args = engine_args({"args": ["--pack", "p"], "parallel": 2})
+        self.assertEqual(args[-2:], ["--batch", "2"])
+
+
+class PickSlot(unittest.TestCase):
+    def engine(self, n):
+        e = StrataEngine("missing-executable", [], lazy=True)
+        e.batch = n
+        e.slot_order = list(range(n))
+        e.slot_busy = [False] * n
+        e.slot_held = [[] for _ in range(n)]
+        e.slot_used = [0.0] * n
+        e.slot_live = [None] * n
+        return e
+
+    def test_the_slot_that_holds_the_conversation(self):
+        e = self.engine(3)
+        e.slot_held = [[1, 2, 3], [], [1, 2, 3, 4, 5]]
+        e.slot_used = [5.0, 0.0, 9.0]
+        self.assertEqual(e.pick_slot([1, 2, 3, 4, 5, 6]), 2)       # the longest held start
+        self.assertEqual(e.pick_slot([1, 2, 3, 9]), 0)
+        self.assertEqual(e.pick_slot([7, 8]), 1)                    # no match: the empty slot first
+        e.slot_busy[1] = True
+        self.assertEqual(e.pick_slot([7, 8]), 0)                    # then the one used longest ago
+        e.slot_busy = [True] * 3
+        self.assertIsNone(e.pick_slot([7, 8]))
+        e.slot_busy = [False] * 3
+        self.assertEqual(e.pick_slot([1, 2, 3]), 1)                 # a held prompt is never the WHOLE prompt
+
+    def test_slots_view(self):
+        e = self.engine(2)
+        e.slot_held[1] = [1, 2]
+        e.slot_live[0] = {"slot": 0, "state": "decoding", "prompt_tokens": 9, "generated": 4,
+                          "started": time.time() - 2, "first_token": time.time() - 1}
+        v = e.slots_view()
+        self.assertEqual(v[0]["state"], "decoding")
+        self.assertEqual(v[0]["generated"], 4)
+        self.assertEqual(v[1], {"slot": 1, "state": "idle", "held_tokens": 2})
+
+
+class ParallelService(unittest.TestCase):
+    """The real StrataEngine and Service over HTTP, the fake engine behind them."""
+
+    def start(self, slots, fit=None, slot_cache=False):
+        import serve.server as server
+        self.tmp = tempfile.TemporaryDirectory()
+        script = Path(self.tmp.name) / "fake_strata.py"
+        script.write_text(FAKE_BATCH, encoding="utf-8")
+        self.log = Path(self.tmp.name) / "requests.log"
+        real = server.subprocess.Popen
+        extra = ["--batch", str(slots)] + (["--fit", str(fit)] if fit is not None else []) + ["--log", str(self.log)]
+        extra += ["--slotcache"] if slot_cache else []
+        with mock.patch.object(server.subprocess, "Popen",
+                               lambda cmd, **kw: real([sys.executable, str(script), *cmd[1:]], **kw)):
+            self.engine = StrataEngine("strata", extra)
+        tok = ByteTokenizer()
+        self.svc = Service(self.engine, tok, ChatTemplate(Path(__file__).parent / "chat_template.jinja"))
+        self.httpd = serve(self.svc, port=0)
+        self.base = f"http://127.0.0.1:{self.httpd.server_address[1]}"
+
+    def tearDown(self):
+        if getattr(self, "httpd", None):
+            self.httpd.shutdown()
+            self.httpd.server_close()
+        if getattr(self, "engine", None):
+            self.engine.unload()
+        if getattr(self, "tmp", None):
+            self.tmp.cleanup()
+
+    def chat(self, text, max_tokens=64):
+        body = {"messages": [{"role": "user", "content": text}], "max_tokens": max_tokens, "reasoning_effort": "none"}
+        req = urllib.request.Request(self.base + "/v1/chat/completions", data=json.dumps(body).encode(),
+                                     headers={"Content-Type": "application/json"})
+        with urllib.request.urlopen(req, timeout=60) as r:
+            return json.loads(r.read().decode())
+
+    def get(self, path):
+        with urllib.request.urlopen(self.base + path, timeout=10) as r:
+            return json.loads(r.read().decode())
+
+    def test_engine_reports_fewer_slots(self):
+        self.start(4, fit=2)
+        self.assertEqual(self.engine.batch, 2)
+        self.assertEqual(self.get("/v1/status")["concurrency"]["serving"], 2)
+
+    def test_engine_turns_batching_off(self):
+        self.start(4, fit=0)
+        self.assertEqual(self.engine.batch, 0)
+        self.assertEqual(self.chat("hi")["choices"][0]["message"]["content"], "ok, done.")
+
+    def test_concurrent_requests_share_the_slots(self):
+        self.start(2)
+        # the first one decodes long enough for the others to arrive while it runs (also on a busy PC)
+        texts = ["first question LONGREPLY", "second question, longer", "a third one waits for a slot"]
+        reply = {t: "ok, " + "la " * 15 + "done." if "LONGREPLY" in t else "ok, done." for t in texts}
+        out, errors, seen = {}, [], []
+
+        def go(t):
+            try:
+                out[t] = self.chat(t)
+            except Exception as e:   # noqa: BLE001 - reported below
+                errors.append(e)
+
+        threads = [threading.Thread(target=go, args=(t,)) for t in texts]
+        for th in threads:
+            th.start()
+            time.sleep(0.05)
+        deadline = time.time() + 30
+        while any(th.is_alive() for th in threads) and time.time() < deadline:
+            m = self.get("/metrics")["live"]
+            seen.append(m)
+            time.sleep(0.01)
+        for th in threads:
+            th.join(30)
+        self.assertEqual(errors, [])
+        for t in texts:
+            self.assertEqual(out[t]["choices"][0]["message"]["content"], reply[t])
+        # the slots were reported, and a request was admitted while another decoded in a slot (the fake logs how many
+        # slots were active at each admission)
+        self.assertTrue(any(m.get("parallel") == 2 and len(m.get("slots", [])) == 2 for m in seen))
+        admitted = [int(x.split()[3]) for x in self.log.read_text().splitlines() if x.startswith("BGEN")]
+        self.assertIn(1, admitted)
+        # every request recorded its own row; nothing is left running
+        with self.svc.status_lock:
+            self.assertEqual(len(self.svc.history), 3)
+            self.assertEqual(self.svc.live_reqs, {})
+            self.assertFalse(self.svc.status["busy"])
+        self.assertEqual(sorted(r["output_tokens"] for r in self.svc.history),
+                         sorted(len(reply[t]) + 1 for t in texts))
+        # never more than two in the slots at once (the fake logs how many were active at each admission)
+        active = [int(x.split()[3]) for x in self.log.read_text().splitlines() if x.startswith("BGEN")]
+        self.assertTrue(all(a <= 1 for a in active), active)
+
+    def race(self, first, second, delay=0.15, max_tokens=64):
+        """`first` sent, `second` after `delay` seconds; -> {text: (reply, seconds it took)}"""
+        res = {}
+
+        def go(t):
+            t0 = time.time()
+            r = self.chat(t, max_tokens=max_tokens)
+            res[t] = (r["choices"][0]["message"]["content"], time.time() - t0)
+        a = threading.Thread(target=go, args=(first,))
+        a.start()
+        time.sleep(delay)
+        go(second)
+        a.join(30)
+        return res
+
+    def test_a_long_prompt_gives_way_to_a_short_one(self):
+        """#656: a long solo prompt read gives way at a chunk boundary; the short request is answered first, then the
+        long one goes on (in the slot its part waited in) and is answered too."""
+        self.start(2)
+        long_q, short_q = "long " * 600, "short"
+        res = self.race(long_q, short_q)
+        self.assertEqual(res[long_q][0], "ok, done.")
+        self.assertEqual(res[short_q][0], "ok, done.")
+        # the long read takes ~1 s (3,000 tokens, 32 a chunk, 10 ms a chunk): the short one did not wait for it
+        self.assertLess(res[short_q][1], res[long_q][1] - 0.5, res)
+        log = self.log.read_text().splitlines()
+        self.assertTrue(any(x.startswith("YIELD ") for x in log), log)
+        with self.svc.status_lock:
+            self.assertEqual(len(self.svc.history), 2)
+        self.assertFalse(any(self.engine.slot_busy))
+
+    def test_an_admission_gives_way_too(self):
+        """The same while another request decodes in a slot: the long one is being admitted, a short one waits."""
+        self.start(3)
+        bg = threading.Thread(target=lambda: self.chat("background", max_tokens=60))
+        bg.start()
+        time.sleep(0.05)
+        long_q, short_q = "long " * 600, "short"
+        res = self.race(long_q, short_q)
+        bg.join(30)
+        self.assertEqual(res[long_q][0], "ok, done.")
+        self.assertLess(res[short_q][1], res[long_q][1] - 0.5, res)
+        log = self.log.read_text().splitlines()
+        self.assertTrue(any(x.startswith("YIELD ") for x in log), log)
+        self.assertFalse(any(self.engine.slot_busy))
+
+    def test_no_way_given_to_a_long_one(self):
+        """Two prompts of about the same length: no BYIELD (it only goes to a much shorter one)."""
+        self.start(2)
+        res = self.race("long " * 600, "lung " * 600)
+        self.assertEqual(sorted(r[0] for r in res.values()), ["ok, done."] * 2)
+        self.assertFalse(any(x.startswith("YIELD ") for x in self.log.read_text().splitlines()))
+
+    def test_left_alone_it_goes_back_to_the_solo_path(self):
+        """A request decoding in a slot whose neighbour finished goes back to the solo path (MTP drafts), continued
+        from its slot - only with an engine that keeps the slots' conversations (INFO slot_cache=1)."""
+        self.start(2, slot_cache=True)
+        res = {}
+        long = threading.Thread(target=lambda: res.setdefault("long", self.chat("LONGREPLY please", max_tokens=200)))
+        long.start()
+        time.sleep(0.05)
+        short = self.chat("short", max_tokens=200)
+        long.join(30)
+        self.assertEqual(short["choices"][0]["message"]["content"], "ok, done.")
+        self.assertEqual(res["long"]["choices"][0]["message"]["content"], "ok, " + "la " * 15 + "done.")
+        kinds = [x.split()[0] for x in self.log.read_text().splitlines()]
+        # solo, then both in slots, then the long one alone again: a GEN after the BGENs
+        self.assertIn("BGEN", kinds)
+        self.assertEqual(kinds[-1], "GEN", kinds)
+        self.assertFalse(any(self.engine.slot_busy))
+        with self.svc.status_lock:
+            self.assertEqual(len(self.svc.history), 2)
+
+    def test_next_turn_goes_back_to_its_slot(self):
+        self.start(2)
+        a = "conversation A, first turn"
+        r1 = {}
+        th = threading.Thread(target=lambda: r1.setdefault("x", self.chat("blocker LONGREPLY", max_tokens=200)))
+        th.start()
+        time.sleep(0.05)
+        first = self.chat(a)
+        th.join(30)
+        deadline = time.time() + 10                     # a slot left at its stop token is freed at its BDONE
+        while not [h for h in self.engine.slot_held if h] and time.time() < deadline:
+            time.sleep(0.02)
+        held = [h for h in self.engine.slot_held if h]
+        self.assertTrue(held, "a finished slot request keeps what its slot holds")
+        # the held tokens are the prompt and every token fed: all of the reply but its last token
+        reply = list(b"ok, done.") + [257]
+        mine = [h for h in held if h[-len(reply) + 1:] == reply[:-1] and bytes(x for x in h if x < 256).find(b"conversation A") >= 0]
+        self.assertEqual(len(mine), 1, held)
+        slot = self.engine.slot_held.index(mine[0])
+        self.assertEqual(self.engine.pick_slot(mine[0] + [10, 11]), slot)
+        self.assertEqual(first["choices"][0]["message"]["content"], "ok, done.")
+
+
+if __name__ == "__main__":
+    unittest.main()

setup.py

@@ -3276,6 +3276,60 @@ def small_card_note(ctx: int, draft_vocab: str | None) -> list[str]:
     return lines
 
 
+PARALLEL_MAX = 8               # #465: the engine's batch window holds at most 8 requests
+PARALLEL_SHARE = 0.2           # #465: the slots' sessions may take this share of the VRAM the expert cache would hold
+PARALLEL_HELD = 0.5            # #465: ... and only where the cache still holds this share of the experts beside them
+PARALLEL_COST_NOTE = ("parallel N reduces waiting for several users but costs about 10-25% speed per request on this "
+                      "card")
+
+
+def parallel_slot_gb(ctx: int, kv: str, streaming: bool) -> float:
+    """#465: the VRAM one batch slot's session takes: its KV cache (12 QSA layers; with KV streaming only the 32K
+    positions the attention reads stay in VRAM) and the DeltaNet state (~0.17 GB).  Measured: 0.56 GiB at 32K int8."""
+    kv_tok = 12 * (576 if kv == "q4_0" else 1056)
+    return (min(ctx, 32768) if streaming else ctx) * kv_tok / 1e9 + 0.17
+
+
+def parallel_recommend(vram_gbs, arena_gb: float, ctx: int, kv: str, streaming: bool) -> int:
+    """#465: how many requests at once ("parallel") to recommend: 0 = none (one at a time).  Only where the experts
+    mostly fit in VRAM - the expert cache (each card's VRAM less ~5 GB, every card of a layer split) still holds
+    PARALLEL_HELD of the model's experts beside the slots' sessions, which take at most PARALLEL_SHARE of it, up to 4.
+    Where the experts mostly run on the CPU a batch reads about as many experts as the requests one by one and every
+    slot's VRAM is expert cache lost: measured on a 12 GB RTX 5070 (Q2_0, 32K), a request alone 11-24% slower with 2-4
+    slots, 4 requests together 63 tok/s against 71 one after the other (docs/BATCHING.md)."""
+    if isinstance(vram_gbs, (int, float)):
+        vram_gbs = [vram_gbs]
+    cache_gb = sum(max(0.0, v - 5) for v in vram_gbs)
+    slot = parallel_slot_gb(ctx, kv, streaming)
+    best = 0
+    for n in (2, 3, 4):
+        if n * slot <= PARALLEL_SHARE * cache_gb and (cache_gb - n * slot) >= PARALLEL_HELD * arena_gb:
+            best = n
+    return best
+
+
+def parallel_note(asked: int | None, vram_gbs, arena_gb: float, ctx: int, kv: str, streaming: bool) -> list[str]:
+    """#465: what setup says about "parallel": the recommendation (or, where it would cost speed, why it is left at
+    one), or how the asked count compares with it (kept as asked: recommend, never force)."""
+    rec = parallel_recommend(vram_gbs, arena_gb, ctx, kv, streaming)
+    slot = parallel_slot_gb(ctx, kv, streaming)
+    if asked is None or asked <= 1:
+        if not rec:
+            return [f"Several requests at once: left at one at a time - {PARALLEL_COST_NOTE} (docs/BATCHING.md)."]
+        return [f"Several requests at once (opt-in): --parallel {rec} decodes up to {rec} together instead of one "
+                f"after the other (each takes ~{slot:.1f} GB of VRAM from the expert cache; docs/BATCHING.md)."]
+    lines = [f"parallel requests: {asked} at once (each takes ~{slot:.1f} GB of VRAM from the expert cache, "
+             f"{asked * slot:.1f} GB in all)"]
+    if asked > PARALLEL_MAX:
+        lines.append(f"the engine runs at most {PARALLEL_MAX} at once; it will use {PARALLEL_MAX}")
+    if not rec:
+        lines.append(f"recommended for this card: one at a time - {PARALLEL_COST_NOTE}; kept as you chose")
+    elif asked > rec:
+        lines.append(f"recommended for this card: {rec} - more slots leave fewer experts in VRAM, which can make every "
+                     "request slower; kept as you chose")
+    return lines
+
+
 DESKTOP_RESERVE_MIB = 3072     # #560 #516: what kept a KDE/Wayland desktop alive beside a full expert cache
 
 
@@ -3552,6 +3606,9 @@ def main() -> int:
     ap.add_argument("--vram-reserve-mib", type=int, metavar="N",
                     help="VRAM in MiB the engine leaves free for other programs (a game, another model; the engine's "
                          "default: 700); the expert cache takes that much less")
+    ap.add_argument("--parallel", type=int, metavar="N",
+                    help="up to N requests decode together (batch slots, opt-in; default: one at a time, the others "
+                         "wait). Each slot takes VRAM from the expert cache; setup says what it recommends")
     ap.add_argument("--kv-streaming", choices=["auto", "on", "off"], default="auto",
                     help="from a 64K context: keep the KV cache in RAM and only the attention's window in VRAM (more "
                          "experts fit on the GPU); auto: when the RAM has room for it")
@@ -4284,6 +4341,17 @@ def main() -> int:
         cfg["draft_vocab"] = draft_vocab
     if a.browser is not None:                          # #609: only when given (else an earlier choice is carried over)
         cfg["open_browser"] = a.browser
+    # #465: requests at once - written only when given (else an earlier "parallel" is carried over); a recommendation
+    streaming = "--kv-resident" in args
+    if a.parallel is not None:
+        if a.parallel >= 2:
+            cfg["parallel"] = a.parallel
+            for i, line in enumerate(parallel_note(a.parallel, [g.get("vram_gb", 0.0) for g in chosen],
+                                                   MODELS[model]["arena_gb"], ctx, kv, streaming)):
+                (ok if i == 0 else warn)(line)
+        else:
+            cfg["parallel"] = 1
+            ok("parallel requests: one at a time (--parallel 1)")
     if vision != "none":
         old_cfg = ROOT / f"strata-{tag.lower()}.json"
         vt = vision_tokens(a.vision_tokens, vision, old_cfg if old_cfg.is_file() else adopted)
@@ -4323,6 +4391,10 @@ def main() -> int:
     say(f"  Next time:        just run {'START-HERE.bat' if WIN else './setup.sh'} (or {script.name}) - it starts right away")
     if vision != "none":
         say("  Images:           send them in the chat page, in chat.py (/image <path>) or over the API")
+    if a.parallel is None:                             # #465: the opt-in, said once (nothing changes)
+        for line in parallel_note(None, [g.get("vram_gb", 0.0) for g in chosen], MODELS[model]["arena_gb"], ctx, kv,
+                                  "--kv-resident" in cfg["args"]):
+            say("  " + line)
     if tuned is False:                                 # #447: a failed tuning is repeated here, not only above
         say("  Tuning:           FAILED (the reason is above): the default settings stay - "
             f"{'START-HERE.bat' if WIN else './setup.sh'} --calibrate tries again")

src/core/conversation_cache_test.cpp

@@ -2,6 +2,7 @@
 
 #include <cstdio>
 #include <cstdlib>
+#include <vector>
 
 using namespace strata::core;
 
@@ -209,5 +210,96 @@ int main() {
         check(!disabled.put(image({1,2,3})) && !no_slots.put(image({1,2,3})), "disabled cache stores nothing");
         check(disabled.best(a,{},true).tokens == 0, "disabled cache has no matches");
     }
+    {
+        // layer-split parking: checkpoints moved apart into stage parts and put back, no running state copied
+        auto running = [](std::initializer_list<int32_t> ids, uint8_t salt, size_t parts) {
+            ConversationCheckpoint c;
+            c.ids = ids;
+            c.gdn.assign(256, salt);
+            c.tails.assign(32, uint8_t(salt + 1));
+            c.used = salt;
+            for (size_t k = 0; k < parts; ++k) {
+                ConversationCheckpoint p;
+                p.gdn.assign(128, uint8_t(salt + 10 + k));
+                c.stage_parts.push_back(std::move(p));
+            }
+            return c;
+        };
+        std::vector<ConversationCheckpoint> checks;
+        checks.push_back(running({1, 2}, 3, 2));
+        checks.push_back(running({1, 2, 3}, 5, 1));       // an incomplete set: kept aside, put back in place
+        checks.push_back(running({1, 2, 3, 4}, 7, 2));
+        const auto original = checks;
+        const uint8_t* gdn0 = checks[0].gdn.data();
+        const uint8_t* part1 = checks[2].stage_parts[1].gdn.data();
+        auto split = conversation_checkpoints_split(std::move(checks), 2);
+        check(checks.empty() && split.stage0.size() == 2 && split.rest.size() == 1 && split.parts.size() == 2 &&
+              split.parts[0].size() == 2 && split.parts[1].size() == 2, "split: complete checkpoints apart, the other aside");
+        check(split.stage0[0].stage_parts.empty() && split.parts[1][1].ids == original[2].ids &&
+              split.parts[1][1].used == original[2].used && split.parts[0][0].gdn == original[0].stage_parts[0].gdn,
+              "split: every stage part carries its checkpoint's ids and LRU stamp");
+        check(split.stage0[0].gdn.data() == gdn0 && split.parts[1][1].gdn.data() == part1,
+              "split moves the running state (same buffers, no copy)");
+        std::vector<ConversationCheckpoint> back;
+        check(conversation_checkpoints_merge(std::move(split), back) && back.size() == 3, "merge puts them back");
+        bool same = true;
+        for (size_t j = 0; j < back.size(); ++j) {
+            same = same && back[j].ids == original[j].ids && back[j].gdn == original[j].gdn &&
+                   back[j].tails == original[j].tails && back[j].used == original[j].used &&
+                   back[j].stage_parts.size() == original[j].stage_parts.size();
+            for (size_t k = 0; same && k < back[j].stage_parts.size(); ++k)
+                same = back[j].stage_parts[k].gdn == original[j].stage_parts[k].gdn;
+        }
+        check(same, "merge(split(x)) == x, in the original order");
+        check(back[0].gdn.data() == gdn0 && back[2].stage_parts[1].gdn.data() == part1, "merge moves them back too");
+
+        auto torn = conversation_checkpoints_split(std::vector<ConversationCheckpoint>(original), 2);
+        torn.parts[1].pop_back();
+        check(!conversation_checkpoints_merge(std::move(torn), back) && back.empty(), "a stage short of a part is refused");
+        auto mixed = conversation_checkpoints_split(std::vector<ConversationCheckpoint>(original), 2);
+        mixed.parts[0][0].ids = {9, 9};
+        check(!conversation_checkpoints_merge(std::move(mixed), back) && back.empty(), "a part of another checkpoint is refused");
+        // a restored image: stage 0's checkpoints and each stage's, with no record of an incomplete one
+        ConversationCheckpointSplit restored;
+        restored.stage0.push_back(running({4, 5}, 11, 0));
+        restored.parts.resize(2);
+        for (auto& p : restored.parts) { auto c = running({4, 5}, 12, 0); p.push_back(std::move(c)); }
+        check(conversation_checkpoints_merge(std::move(restored), back) && back.size() == 1 &&
+              back[0].stage_parts.size() == 2, "a restored image's parts are all whole");
+        auto none = conversation_checkpoints_split({}, 3);
+        check(conversation_checkpoints_merge(std::move(none), back) && back.empty(), "no checkpoints: nothing to do");
+    }
+    {
+        // the later stages' retained K/V, kept with the first stage's and limited together
+        auto kv = [](size_t n) {
+            std::vector<ConversationKv> v(2);
+            for (auto& l : v) l.k.resize(n, 1);
+            return v;
+        };
+        ConversationCache cache(1 << 20, 4);
+        std::vector<std::vector<ConversationKv>> stage_kv;
+        stage_kv.push_back(kv(1000));
+        stage_kv.push_back(kv(2000));
+        cache.retain(kv(500), 40, std::move(stage_kv));
+        check(cache.retained_bytes() >= 2 * (500 + 1000 + 2000), "retained bytes count every stage");
+        cache.limit_reuse(30);
+        auto r = cache.take_reuse();
+        check(r.kv.size() == 2 && r.stages.size() == 2 && r.unchanged_tokens == 30 &&
+              r.stages[0].unchanged_tokens == 30 && r.stages[1].unchanged_tokens == 30 &&
+              r.stages[1].captured_tokens == 40, "every stage's reuse is limited to the first rewrite");
+        check(cache.retained_bytes() == 0, "taken once");
+        std::vector<std::vector<ConversationKv>> big;
+        big.push_back(kv(1 << 20));
+        cache.retain(kv(500), 40, std::move(big));
+        check(cache.retained_bytes() == 0, "a stage's K/V over the budget drops the whole reuse");
+        std::vector<std::vector<ConversationKv>> one;
+        one.push_back(kv(100));
+        cache.retain(kv(500), 40, std::move(one));
+        cache.limit_reuse(0);
+        check(cache.retained_bytes() == 0, "a rewrite from the start drops every stage's reuse");
+        cache.retain(kv(500), 40);
+        r = cache.take_reuse();
+        check(r.kv.size() == 2 && r.stages.empty(), "no layer split: no stage reuse, as before");
+    }
     std::printf("conversation_cache_test: %d checks passed\n", checks);
 }

src/core/conversation_snapshot_test.cpp

@@ -187,6 +187,47 @@ void full_session(int fmt, int mode, int experts) {
         check(conversation_kv_verify(incremental.kv.back(),draft.state,g,70,false,fingerprint,err),"incremental draft authoritative and ring read-back");
         ids.resize(65);
     }
+    {
+        // a layer split's later stage: the same image WITHOUT the draft layer's K/V (draft == nullptr)
+        check(conversation_snapshot_restore(a,ss,g,draft.state,err)==ConversationRestore::restored,"restore A before stage images");
+        size_t with_draft=0,without=0;
+        check(conversation_snapshot_bytes(view,ss,g,draft.state,with_draft,err) &&
+              conversation_snapshot_bytes(view,ss,g,nullptr,without,err) && without<with_draft,
+              "a stage image's estimate leaves the draft out");
+        SavedConversation stage,stage_b,stage_back;
+        check(conversation_snapshot_save(stage,view,ss,g,nullptr,err),"capture a stage image");
+        check(stage.kv.size()==a.kv.size()-1 && equal(stage.kv[0],a.kv[0]) && stage.live.gdn==a.live.gdn,
+              "a stage image holds the session's own layers, no draft");
+        check(stage.bytes()<a.bytes(),"a stage image is smaller by the draft ring");
+        check(conversation_snapshot_validate(stage,ss,g,nullptr,err),"validate a stage image without a draft");
+        check(!conversation_snapshot_validate(stage,ss,g,draft.state,err),"a stage image is refused where a draft is expected");
+        check(!conversation_snapshot_validate(a,ss,g,nullptr,err),"an image with a draft is refused where none is expected");
+        fill(201);
+        check(conversation_snapshot_save(stage_b,view,ss,g,nullptr,err),"capture another stage state");
+        check(conversation_snapshot_restore(a,ss,g,nullptr,err)==ConversationRestore::invalid,
+              "restoring a draft image as a stage image is refused before any write");
+        check(conversation_snapshot_save(stage_back,view,ss,g,nullptr,err) && stage_back.live.gdn==stage_b.live.gdn &&
+              equal(stage_back.kv[0],stage_b.kv[0]),"the refusal left the stage untouched");
+        check(conversation_snapshot_restore(stage,ss,g,nullptr,err)==ConversationRestore::restored,"restore a stage image");
+        check(conversation_snapshot_save(stage_back,view,ss,g,nullptr,err) && stage_back.live.gdn==stage.live.gdn &&
+              stage_back.live.dead==stage.live.dead && equal(stage_back.kv[0],stage.kv[0]),"stage A/B/A exactness");
+        // its retained K/V: the next park copies only what changed, and equals a full capture
+        ConversationKvReuse reuse{stage.kv,65,65,{}};
+        SavedConversation fresh,incremental;
+        size_t peak=0,reused=0;
+        check(conversation_snapshot_save(fresh,view,ss,g,nullptr,err),"full stage capture reference");
+        check(conversation_snapshot_capture_bytes(reuse,view,ss,g,nullptr,peak,err),"admit an incremental stage capture");
+        check(conversation_snapshot_save(incremental,view,ss,g,nullptr,err,std::move(reuse),&reused),"incremental stage capture");
+        check(reused>0 && incremental.bytes()<=peak && equal(incremental.kv[0],fresh.kv[0]) &&
+              incremental.live.gdn==fresh.live.gdn,"an incremental stage capture reuses pages and equals a full one");
+        SavedConversation split_image=stage;
+        split_image.stage_images.push_back(stage);
+        check(!conversation_snapshot_validate(split_image,ss,g,draft.state,err),
+              "a layer split's image is refused by the whole-session form");
+        ConversationKvReuse wrong{a.kv,65,65,{}};
+        check(!conversation_snapshot_capture_bytes(wrong,view,ss,g,nullptr,peak,err),"a draft image's K/V is not a stage's reuse");
+    }
+    check(conversation_snapshot_restore(a,ss,g,draft.state,err)==ConversationRestore::restored,"restore A after stage images");
     check(conversation_checkpoint_restore(a.checkpoints[0],ss,g,err),"restore early running checkpoint");
     std::vector<uint8_t> spare(sizes.dead);
     cuda_check(cudaMemcpy(spare.data(),main.state.idx_pooled,sizes.dead,cudaMemcpyDeviceToHost));

src/core/conversation_state.cpp

@@ -193,7 +193,7 @@ bool conversation_checkpoint_restore(const ConversationCheckpoint& c, SessionSta
 }
 
 bool conversation_snapshot_bytes(const ConversationView& view, const SessionState& ss,
-                                 const ModelGeometry& g, const QsaState& draft, size_t& bytes, std::string& error) {
+                                 const ModelGeometry& g, const QsaState* draft, size_t& bytes, std::string& error) {
     bytes = 0;
     if (!view_validate(view, ss, g, error)) return false;
     ConversationStateSizes z;
@@ -203,7 +203,7 @@ bool conversation_snapshot_bytes(const ConversationView& view, const SessionStat
     if (!product(ids, {view.ids.size(), sizeof(int32_t)}) ||
         !product(images, {view.images.size(), sizeof(ConversationImageKey)}) ||
         !product(checkpoints, {view.checkpoints.size(), sizeof(ConversationCheckpoint)}) ||
-        !product(layers, {qsa + 1, sizeof(ConversationKv)}) ||
+        !product(layers, {qsa + (draft ? 1 : 0), sizeof(ConversationKv)}) ||
         !product(tails, {qsa, z.tail}) || !product(dead, {qsa, z.dead}) ||
         !product(positions, {qsa, z.block_pos})) return fail(error, "snapshot metadata byte count overflow");
     for (size_t n : {ids, images, checkpoints, layers, tails, dead, positions, z.gdn, ss.ple_hist ? z.ple : 0})
@@ -211,8 +211,8 @@ bool conversation_snapshot_bytes(const ConversationView& view, const SessionStat
     for (const auto& c : view.checkpoints)
         if (!metadata_bytes(c, bytes)) return fail(error, "checkpoint byte count overflow");
     const int64_t upto = (int64_t) view.ids.size(); // view_validate bounds this by signed max_cells
-    for (uint64_t i = 0; i <= qsa; ++i) {
-        const auto& st = i == qsa ? draft : owned(ss, (size_t) i);
+    for (uint64_t i = 0; i < qsa + (draft ? 1 : 0); ++i) {
+        const auto& st = i == qsa ? *draft : owned(ss, (size_t) i);
         const size_t n = conversation_kv_bytes(st, g, upto, i != qsa);
         if (!n || !add(bytes, n)) return fail(error, "invalid or overflowing K/V byte estimate");
     }
@@ -221,16 +221,16 @@ bool conversation_snapshot_bytes(const ConversationView& view, const SessionStat
 
 bool conversation_snapshot_capture_bytes(const ConversationKvReuse& reuse, const ConversationView& view,
                                          const SessionState& ss, const ModelGeometry& g,
-                                         const QsaState& draft, size_t& bytes, std::string& error) {
+                                         const QsaState* draft, size_t& bytes, std::string& error) {
     if (!conversation_snapshot_bytes(view, ss, g, draft, bytes, error)) return false;
     if (reuse.kv.empty()) return true;
-    const size_t layers = owned_qsa(ss) + 1;
+    const size_t layers = owned_qsa(ss) + (draft ? 1 : 0);
     if (reuse.kv.size() != layers || reuse.unchanged_tokens < 0 ||
         reuse.unchanged_tokens > reuse.captured_tokens || reuse.unchanged_tokens > int64_t(view.ids.size()))
         return fail(error, "invalid retained K/V prefix");
     for (size_t i = 0; i < layers; ++i) {
-        const bool index = i + 1 != layers;
-        const auto& st = index ? owned(ss, i) : draft;
+        const bool index = i < owned_qsa(ss);
+        const auto& st = index ? owned(ss, i) : *draft;
         if (!conversation_kv_validate(reuse.kv[i], st, g, reuse.captured_tokens, index, error)) return false;
         const size_t fresh = conversation_kv_bytes(st, g, int64_t(view.ids.size()), index);
         size_t retained = 0;
@@ -246,7 +246,7 @@ bool conversation_snapshot_capture_bytes(const ConversationKvReuse& reuse, const
 
 bool conversation_snapshot_save(SavedConversation& image, const ConversationView& view,
                                 const SessionState& ss, const ModelGeometry& g,
-                                const QsaState& draft, std::string& error,
+                                const QsaState* draft, std::string& error,
                                 ConversationKvReuse reuse, size_t* reused_bytes) {
     size_t estimate = 0;
     if (!conversation_snapshot_capture_bytes(reuse, view, ss, g, draft, estimate, error) || !sync(error)) return false;
@@ -259,7 +259,7 @@ bool conversation_snapshot_save(SavedConversation& image, const ConversationView
     const int64_t unchanged = reuse.kv.empty() ? 0 : reuse.unchanged_tokens;
     captured.kv = std::move(reuse.kv);
     const size_t layers = owned_qsa(ss);
-    captured.kv.resize(layers + 1);
+    captured.kv.resize(layers + (draft ? 1 : 0));
     if (!conversation_checkpoint_save(captured.live, ss, g, error)) return false;
     const int64_t upto = (int64_t) view.ids.size();
     for (size_t j = 0; j < layers; ++j)
@@ -267,15 +267,16 @@ bool conversation_snapshot_save(SavedConversation& image, const ConversationView
                                   unchanged, reused_bytes)) return false;
     // The draft's final cell may not have been computed when the output cap was
     // reached. Refresh that page even when the main prefix continued unchanged.
-    if (!conversation_kv_save(captured.kv.back(), draft, g, upto, false, error,
-                              std::max<int64_t>(0, unchanged - 1), reused_bytes)) return false;
+    if (draft && !conversation_kv_save(captured.kv.back(), *draft, g, upto, false, error,
+                                       std::max<int64_t>(0, unchanged - 1), reused_bytes)) return false;
     image = std::move(captured);
     return true;
 }
 
 bool conversation_snapshot_validate(const SavedConversation& image, const SessionState& ss,
-                                    const ModelGeometry& g, const QsaState& draft, std::string& error) {
-    if (!image.live.stage_parts.empty()) return fail(error, "layer-split parking is not supported");
+                                    const ModelGeometry& g, const QsaState* draft, std::string& error) {
+    if (!image.live.stage_parts.empty())
+        return fail(error, "a running checkpoint's stage parts belong in the stage images, not in live.stage_parts");
     if (image.geometry != geometry_key(g)) return fail(error, "incompatible runtime geometry");
     // an image holds exactly one carve's running state and K/V: same layer range or nothing
     if (image.layer_lo != ss.layer_lo || image.layer_hi != ss.layer_hi)
@@ -283,23 +284,49 @@ bool conversation_snapshot_validate(const SavedConversation& image, const Sessio
     const ConversationView view{image.live.ids, image.live.imgs, image.checkpoints, image.cvec};
     if (!view_validate(view, ss, g, error) || !conversation_checkpoint_validate(image.live, ss, g, error)) return false;
     const size_t layers = owned_qsa(ss);
-    if (image.kv.size() != layers + 1) return fail(error, "invalid K/V layer count");
+    if (image.kv.size() != layers + (draft ? 1 : 0)) return fail(error, "invalid K/V layer count");
     const int64_t upto = (int64_t) image.live.ids.size();
     for (size_t j = 0; j < layers; ++j)
         if (!conversation_kv_validate(image.kv[j], owned(ss, j), g, upto, true, error)) return false;
-    return conversation_kv_validate(image.kv.back(), draft, g, upto, false, error);
+    return !draft || conversation_kv_validate(image.kv.back(), *draft, g, upto, false, error);
 }
 
 ConversationRestore conversation_snapshot_restore(const SavedConversation& image, SessionState& ss,
-                                                   const ModelGeometry& g, const QsaState& draft, std::string& error) {
+                                                   const ModelGeometry& g, const QsaState* draft, std::string& error) {
     if (!conversation_snapshot_validate(image, ss, g, draft, error)) return ConversationRestore::invalid;
     if (!sync(error)) return ConversationRestore::transfer_failed;
     const int64_t upto = (int64_t) image.live.ids.size();
     for (size_t j = 0; j < owned_qsa(ss); ++j)
         if (!conversation_kv_restore(image.kv[j], owned(ss, j), g, upto, true, error))
             return ConversationRestore::transfer_failed;
-    if (!conversation_kv_restore(image.kv.back(), draft, g, upto, false, error) ||
+    if ((draft && !conversation_kv_restore(image.kv.back(), *draft, g, upto, false, error)) ||
         !conversation_checkpoint_restore(image.live, ss, g, error)) return ConversationRestore::transfer_failed;
     return ConversationRestore::restored;
 }
+
+// the draft layer's K/V included (the stage that owns the draft head, or no layer split)
+bool conversation_snapshot_bytes(const ConversationView& view, const SessionState& ss, const ModelGeometry& g,
+                                 const QsaState& draft, size_t& bytes, std::string& error) {
+    return conversation_snapshot_bytes(view, ss, g, &draft, bytes, error);
+}
+bool conversation_snapshot_capture_bytes(const ConversationKvReuse& reuse, const ConversationView& view,
+                                         const SessionState& ss, const ModelGeometry& g, const QsaState& draft,
+                                         size_t& bytes, std::string& error) {
+    return conversation_snapshot_capture_bytes(reuse, view, ss, g, &draft, bytes, error);
+}
+bool conversation_snapshot_save(SavedConversation& image, const ConversationView& view, const SessionState& ss,
+                                const ModelGeometry& g, const QsaState& draft, std::string& error,
+                                ConversationKvReuse reuse, size_t* reused_bytes) {
+    return conversation_snapshot_save(image, view, ss, g, &draft, error, std::move(reuse), reused_bytes);
+}
+bool conversation_snapshot_validate(const SavedConversation& image, const SessionState& ss, const ModelGeometry& g,
+                                    const QsaState& draft, std::string& error) {
+    // the whole-session form: a layer split's image (stage images) is restored stage by stage, never through here
+    if (!image.stage_images.empty()) return fail(error, "a layer split's image restored as a single session");
+    return conversation_snapshot_validate(image, ss, g, &draft, error);
+}
+ConversationRestore conversation_snapshot_restore(const SavedConversation& image, SessionState& ss,
+                                                   const ModelGeometry& g, const QsaState& draft, std::string& error) {
+    return conversation_snapshot_restore(image, ss, g, &draft, error);
+}
 } // namespace strata::core

src/core/native_dense.cpp

@@ -1,4 +1,5 @@
 #include "strata/core/native_dense.hpp"
+#include <cstdlib>
 #include "strata/core/weights.hpp"
 #include "strata/artifact/gguf_reader.hpp"
 #include "strata/kernels/native_mmvq.hpp"
@@ -77,7 +78,12 @@ NativeDense::~NativeDense() {
 }
 
 bool NativeDense::load(const std::vector<std::string>& shards, WeightTable& table, std::string& err,
-                       bool include_ple_key) {
+                       bool include_ple_key, int64_t layer_lo, int64_t layer_hi) {
+    auto outside = [&](const std::string& name) {   // a blk.<l>. tensor of another stage's layers
+        if (layer_hi < 0 || name.rfind("blk.", 0) != 0) return false;
+        const long l = std::strtol(name.c_str() + 4, nullptr, 10);
+        return l < layer_lo || l >= layer_hi;
+    };
     if (scratch_ || !weights_.empty()) { err = "native dense: already loaded"; return false; }
     if (shards.empty()) { err = "native dense: at least one GGUF shard is required"; return false; }
     try {
@@ -146,7 +152,7 @@ bool NativeDense::load(const std::vector<std::string>& shards, WeightTable& tabl
                 }
             }
             for (const auto& tensor : gguf.tensors()) {
-                if (!eligible(tensor, include_ple_key)) continue;
+                if (!eligible(tensor, include_ple_key) || outside(tensor.name)) continue;
                 if (!in_range(tensor.name) && tensor.name.find("ple") == std::string::npos) continue;
                 if (!seen.insert(tensor.name).second) {
                     err = "native dense: duplicate tensor " + tensor.name; return false;

src/core/verify.cpp

@@ -279,6 +279,12 @@ Verifier::~Verifier() {
     for (auto& e : exec_)
         if (e) cudaGraphExecDestroy(e);
     if (commit_exec_) cudaGraphExecDestroy(commit_exec_);
+    for (auto& kv : exec_bm_)
+        if (kv.second) cudaGraphExecDestroy(kv.second);
+    for (auto& kv : commit_bm_)
+        if (kv.second) cudaGraphExecDestroy(kv.second);
+    if (arena_b_) cudaFree(arena_b_);
+    if (h_commitb_) cudaFreeHost(h_commitb_);
     if (cs_) cudaStreamDestroy(cs_);
     if (sh_cs_) cudaStreamDestroy(sh_cs_);
     if (copy_) { cudaStreamSynchronize(copy_); cudaStreamDestroy(copy_); }
@@ -552,14 +558,18 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
     const int64_t TS = (s.idx_block - 1) * ID;
     const bool ple_on = ss.ple.ready() && ple_stage();
     auto Rt = [&](int t) { return R_ + (size_t) t * HC * N; };
-    const int G = (split_ && T >= 2) ? 2 : 1;
+    const int G = (split_ && T >= 2 && !batch_rec_) ? 2 : 1;   // a batch window is one group
     static const bool dec_batch = [] { const char* v = std::getenv("STRATA_DEC_BATCH"); return v == nullptr || std::atoi(v) != 0; }();
     auto stamp = [&](int64_t l, int i, int grp) {
         if (prof_on_ && grp == 0) gpu_stamp(prof_, (int) (l * kProfPer + i), cs);
         if (trace_m_ != nullptr) gpu_stamp(trace_m_, (int) ((l * kProfPer + i) * 2 + grp), cs);   // #649
     };
     const int tb_[2] = {0, (T + 1) / 2}, te_[2] = {G == 2 ? (T + 1) / 2 : T, T};
-    groups_[T] = G;
+    if (!batch_rec_) groups_[T] = G;
+    const int64_t nQall = g.n_qsa_layers();
+    // a batch window: row t is slot t, whose state lives in its own session
+    auto slot_ss = [&](int t) -> SessionState& { return batch_rec_ ? *slots_[(size_t) brow_[t]] : ss; };
+    const int hrow0 = batch_rec_ ? row_base_ : 0;   // a slot group's own hand-off rows
 
     // ---- the window's inputs, from mapped staging
     copy_i32_from_mapped(tok_, m_tok_, T, cs);
@@ -573,9 +583,10 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
     // residual, pending write and inject (see set_stage)
     const int64_t HB = Verifier::handoff_floats(g);
     if (lb_ > 0) {
-        copy_from_mapped(R_, hand_in_, (int64_t) T * HC * N, cs);
-        copy_from_mapped(bo_, hand_in_ + (size_t) T * HC * N, (int64_t) T * N, cs);
-        copy_from_mapped(inj2_, hand_in_ + (size_t) T * (HC + 1) * N, (int64_t) T * HC, cs);
+        const float* hin = hand_in_ + (size_t) hrow0 * HB;   // the group's rows: [R][bo][inj] contiguous
+        copy_from_mapped(R_, hin, (int64_t) T * HC * N, cs);
+        copy_from_mapped(bo_, hin + (size_t) T * HC * N, (int64_t) T * N, cs);
+        copy_from_mapped(inj2_, hin + (size_t) T * (HC + 1) * N, (int64_t) T * HC, cs);
     } else if (const NativeEmbed* ne = native_embed()) {       // plan v0.3 P6: the GGUF-form table
         ne->gather_dev(tok_, T, emb_, cs);
         broadcast_streams(emb_, R_, N, (int) HC, T, cs);
@@ -651,19 +662,20 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 PleOut po;
                 po.normalized = normalized;
                 po.result = Rt(t);
+                float* hist = batch_rec_ ? slot_ss(t).ple_hist : ss.ple.hist;   // the row's own history
                 try {
                     if (ple_batch_kv) {
-                        ple_block_projected(xn_ + (size_t) t * HC * N, z_ + (size_t) t * N, Rt(t), ss.ple.hist,
+                        ple_block_projected(xn_ + (size_t) t * HC * N, z_ + (size_t) t * N, Rt(t), hist,
                                             ss.ple.w, po, ss.ple.scratch, cs);
                     } else {
-                        ple_block(ple_ + t * N, Rt(t), ss.ple.hist, ss.ple.w, po, ss.ple.scratch, cs);
+                        ple_block(ple_ + t * N, Rt(t), hist, ss.ple.w, po, ss.ple.scratch, cs);
                     }
-                    ple_history_advance(ss.ple.hist, normalized, cs);
+                    ple_history_advance(hist, normalized, cs);
                 } catch (const std::exception& e) {
                     err = std::string("verify PLE: ") + e.what();
                     return false;
                 }
-                copy_from_mapped(hist_snap_ + (size_t) t * HS, ss.ple.hist, HS, cs);
+                copy_from_mapped(hist_snap_ + (size_t) t * HS, hist, HS, cs);
             }
             pending = false;
         }
@@ -706,6 +718,15 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 native_quantize_q8_1(xm, xq_, (int) N, n, cs);
                 native_mmvq(wqkv->native_type, wqkv->native_data, xq_, qkv + (size_t) tb * C, (int) N, (int) C, n, cs);
                 stamp(l, 2, grp);
+                if (batch_rec_) {   // each row from its own slot's conv history, one row each
+                    for (int t = tb; t < te; ++t) {
+                        SessionState& sx = slot_ss(t);
+                        const float* cx = sx.gdn_state + (size_t) (gi - sx.gdn_ord0) * gdn_floats +
+                                          (uint64_t) g.ssm_state_size * g.ssm_v_heads * g.ssm_state_size;
+                        gdn_conv_l2_multi(cx, qkv + (size_t) t * C, (const float*) wc->data, hb + (size_t) t * C, (int) C,
+                                          (int) (2 * HK), EPS, 1, cs, 0);
+                    }
+                } else
                 gdn_conv_l2_multi(conv, qkv, (const float*) wc->data, hb, (int) C, (int) (2 * HK), EPS, n, cs, tb);
                 stamp(l, 3, grp);
                 gdn_ab_multi(xm, (const uint16_t*) wa->data, (const uint16_t*) wb->data, (const float*) wdt->data,
@@ -715,6 +736,15 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 native_mmvq(wg->native_type, wg->native_data, xq_, z_ + (size_t) tb * ZV, (int) N, (int) ZV, n, cs);
                 stamp(l, 5, grp);
                 // the recurrence from the untouched state over tokens [0, te); outputs only for this group's
+                if (batch_rec_) {   // each row's recurrence from its own slot's state, one token
+                    for (int t = tb; t < te; ++t) {
+                        SessionState& sx = slot_ss(t);
+                        float* stx = sx.gdn_state + (size_t) (gi - sx.gdn_ord0) * gdn_floats;
+                        gdn_step_norm_multi(stx, hb + (size_t) t * C, (int) C, gate + (size_t) t * HV, beta + (size_t) t * HV,
+                                            z_ + (size_t) t * ZV, (const float*) wnm->data, EPS, y_ + (size_t) t * ZV,
+                                            (int) HK, (int) HV, 1, nullptr, cs, 0);
+                    }
+                } else
                 gdn_step_norm_multi(state, hb, (int) C, gate, beta, z_, (const float*) wnm->data, EPS, y_, (int) HK,
                                     (int) HV, te, nullptr, cs, tb);
                 stamp(l, 6, grp);
@@ -759,8 +789,14 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                     fwht256_inplace_cuda(vcur_ + tb * NKV * HD, (int64_t) n * NKV, cs);
                 }
                 stamp(l, 8, grp);
+                if (batch_rec_) {   // every slot's indexer tail, restored by its commit
+                    for (int t = tb; t < te; ++t)
+                        copy_from_mapped(tail_snap_b_ + ((size_t) brow_[t] * nQall + qi) * TS, slot_ss(t).qsa_states[qi].idx_tail,
+                                         TS, cs);
+                } else
                 if (grp == 0) copy_from_mapped(tail_snap_ + (size_t) qi * TS, st.idx_tail, TS, cs);
                 for (int t = tb; t < te; ++t) {
+                    const QsaState& st = slot_ss(t).qsa_states[qi];   // the row's own K/V (ss's outside a batch)
                     const int32_t* step_t = step_ + t * kStepCount;
                     if (st.kv_hybrid) {   // K8V4: the unused half's lanes folded onto the used pool (layer.cpp)
                         kv_append_q8_step(st.k_q, st.k_q, st.k_scale, st.k_scale, st.page_table, step_t,
@@ -777,11 +813,13 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                         kv_append_step(st.k_pool, st.v_pool, st.page_table, step_t, kcur_ + t * NKV * HD,
                                        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)
+                for (int t = tb; t < te; ++t) {
+                    const QsaState& sx = slot_ss(t).qsa_states[qi];
+                    const QsaIndexerBuffers ib{sx.idx_tail, sx.idx_dead, sx.idx_pooled, sx.idx_block_pos};
                     native_qsa_indexer_append(idx_raw + t * ID, step_ + t * kStepCount + kStepPos, 0,
-                                              (const float*) wikn->data, EPS, ib, s, st.max_cells,
+                                              (const float*) wikn->data, EPS, ib, s, sx.max_cells,
                                               rope_scaling(), cs);
+                }
                 stamp(l, 9, grp);
                 native_mmvq(wq->native_type, wq->native_data, xq_, qfull_ + tb * NH * 2 * HD, (int) N, (int) (NH * 2 * HD),
                             n, cs);
@@ -814,6 +852,20 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 }
                 }
                 stamp(l, 10, grp);
+                if (batch_rec_) {   // each row selects and attends over its own slot's K/V
+                    for (int t = tb; t < te; ++t) {
+                        const QsaState& sx = slot_ss(t).qsa_states[qi];
+                        qsa_block_scores(sx.idx_pooled, sx.idx_dead, qidx_ + t * IQ * ID, step_ + t * kStepCount, 1,
+                                         max_blocks_, s, scores_ + (size_t) t * max_blocks_, cs);
+                        qsa_block_topk(scores_ + (size_t) t * max_blocks_, step_ + t * kStepCount, 1, max_blocks_, cap_, s,
+                                       sel_ + (size_t) t * cap_, cs);
+                        qsa_kv_resolve(sx, *g_, sel_ + (size_t) t * cap_, step_ + t * kStepCount, 1, cap_, cs);
+                        const QsaAttnPools px = qsa_attn_pools(sx);
+                        qsa_decode_attn_batch(qcur_ + t * NH * HD, px, sel_ + (size_t) t * cap_, step_ + t * kStepCount,
+                                              cap_, s, attn_scratch_ + (size_t) t * attn_scratch_floats_,
+                                              attn_ + t * NH * HD, 1, cs);
+                    }
+                } else {
                 qsa_block_scores(st.idx_pooled, st.idx_dead, qidx_ + tb * IQ * ID, step_ + tb * kStepCount, n, max_blocks_,
                                  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,
@@ -825,6 +877,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 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);
+                }
                 stamp(l, 13, grp);
                 if (st.kv_rot || st.kv_hybrid) fwht256_inplace_cuda(attn_ + tb * NH * HD, (int64_t) n * NH, cs);   // back: H^-1 = H
                 if (qb) native_qsa_gate_apply(attn_ + tb * NH * HD, qfull_ + tb * NH * 2 * HD, attn32_ + tb * NH * HD,
@@ -1057,9 +1110,10 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             if (l + 1 < le_ && !pre(l + 1, grp)) return false;
         }
     if (le_ < g.n_layers) {   // a layer split's earlier stage: hand the residual on, no head
-        copy_from_mapped(hand_out_, R_, (int64_t) T * HC * N, cs);
-        copy_from_mapped(hand_out_ + (size_t) T * HC * N, bo_, (int64_t) T * N, cs);
-        copy_from_mapped(hand_out_ + (size_t) T * (HC + 1) * N, inj2_, (int64_t) T * HC, cs);
+        float* hout = hand_out_ + (size_t) hrow0 * HB;
+        copy_from_mapped(hout, R_, (int64_t) T * HC * N, cs);
+        copy_from_mapped(hout + (size_t) T * HC * N, bo_, (int64_t) T * N, cs);
+        copy_from_mapped(hout + (size_t) T * (HC + 1) * N, inj2_, (int64_t) T * HC, cs);
         return true;
     }
 
@@ -1299,10 +1353,45 @@ void Verifier::stage_inputs(int T, const int32_t* tokens, int64_t pos0) {
     staged_ = true;
 }
 
+void Verifier::collect_profile() {
+    const ModelGeometry& g = *g_;
+    cudaMemcpy(prof_h_.data(), prof_, prof_h_.size() * 8, cudaMemcpyDeviceToHost);
+    const int64_t L = g.n_layers;
+    auto at = [&](int64_t l, int i) { return prof_h_[(size_t) (l * kProfPer + i)]; };
+    // D8: the derived columns below read stamps the hc-read kernels write themselves or the next
+    // layer's first.  A slot no kernel stamped is 0 and its unsigned difference wrapped to ~1e19 ns -
+    // which is why (gap)/head/hc0 printed ~1e14 ms per window.  A missing or out-of-order stamp now
+    // contributes nothing.
+    const auto gap = [](unsigned long long to, unsigned long long from) {
+        return (from != 0 && to != 0 && to >= from) ? (double) (to - from) : 0.0;
+    };
+    for (int64_t l = 0; l < L; ++l) {
+        const int kind = is_qsa_layer(g, l) ? 1 : 0;
+        unsigned long long prev = at(l, 0);
+        for (int i = 1; i <= 24; ++i) {
+            const unsigned long long x = at(l, i);
+            if (x == 0 || x < prev) continue;
+            prof_sum_[kind][i] += (double) (x - prev);
+            prev = x;
+        }
+        if (l + 1 < L) prof_sum_[kind][25] += gap(at(l + 1, 0), at(l, 24));
+        const double dn = gap(at(l, 27), at(l, 0)), dd = gap(at(l, 28), at(l, 27)), du = gap(at(l, 1), at(l, 28));
+        if (dn > 0 && dd > 0 && du > 0) {   // the split exists: show it split, not twice
+            prof_sum_[kind][27] += dn;      // hc-read0: norm
+            prof_sum_[kind][28] += dd;      //           down
+            prof_sum_[kind][29] += du;      //           up (through both halves, as before)
+            prof_sum_[kind][1] -= gap(at(l, 1), at(l, 0));   // (hc-read0 shown split)
+        }
+    }
+    prof_sum_[0][26] += gap(at(L, 1), at(L, 0));
+    ++prof_windows_;
+}
+
 bool Verifier::run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool, void* user, int32_t* out,
                    std::string& err) {
     using namespace strata::kernels;
     const OnDevice on_device(device_);
+    last_batch_ = false;
     if (T < 1 || T > max_t_) { err = "verify: window size out of range"; return false; }
     if (released_.load()) { err = "verify: an earlier window never finished on the GPU (#267); restart the engine"; return false; }
     const ModelGeometry& g = *g_;
@@ -1449,38 +1538,7 @@ bool Verifier::run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool,
         cudaStreamSynchronize(copy_);   // no host function of this window may raise flag B in the next one
         copy_used_ = false;
     }
-    if (prof_on_ && G == 1) {       // the window's GPU stage stamps
-        cudaMemcpy(prof_h_.data(), prof_, prof_h_.size() * 8, cudaMemcpyDeviceToHost);
-        const int64_t L = g.n_layers;
-        auto at = [&](int64_t l, int i) { return prof_h_[(size_t) (l * kProfPer + i)]; };
-        // D8: the derived columns below read stamps the hc-read kernels write themselves or the next
-        // layer's first.  A slot no kernel stamped is 0 and its unsigned difference wrapped to ~1e19 ns -
-        // which is why (gap)/head/hc0 printed ~1e14 ms per window.  A missing or out-of-order stamp now
-        // contributes nothing.
-        const auto gap = [](unsigned long long to, unsigned long long from) {
-            return (from != 0 && to != 0 && to >= from) ? (double) (to - from) : 0.0;
-        };
-        for (int64_t l = 0; l < L; ++l) {
-            const int kind = is_qsa_layer(g, l) ? 1 : 0;
-            unsigned long long prev = at(l, 0);
-            for (int i = 1; i <= 24; ++i) {
-                const unsigned long long x = at(l, i);
-                if (x == 0 || x < prev) continue;
-                prof_sum_[kind][i] += (double) (x - prev);
-                prev = x;
-            }
-            if (l + 1 < L) prof_sum_[kind][25] += gap(at(l + 1, 0), at(l, 24));
-            const double dn = gap(at(l, 27), at(l, 0)), dd = gap(at(l, 28), at(l, 27)), du = gap(at(l, 1), at(l, 28));
-            if (dn > 0 && dd > 0 && du > 0) {   // the split exists: show it split, not twice
-                prof_sum_[kind][27] += dn;      // hc-read0: norm
-                prof_sum_[kind][28] += dd;      //           down
-                prof_sum_[kind][29] += du;      //           up (through both halves, as before)
-                prof_sum_[kind][1] -= gap(at(l, 1), at(l, 0));   // (hc-read0 shown split)
-            }
-        }
-        prof_sum_[0][26] += gap(at(L, 1), at(L, 0));
-        ++prof_windows_;
-    }
+    if (prof_on_ && G == 1) collect_profile();   // the window's GPU stage stamps
     // ---- 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
@@ -1535,7 +1593,7 @@ void Verifier::set_plan_slot(int grp) {
     sink_.ptr = (unsigned long long*) (base + ptr_off);
     sink_.ptr2 = sink_.ptr + cap;
     sink_.start2 = base + ptr_off + 4 * cap;
-    const int G = groups_[last_t_] > 0 ? groups_[last_t_] : 1;
+    const int G = last_batch_ ? 1 : (groups_[last_t_] > 0 ? groups_[last_t_] : 1);
     const int64_t per = G == 2 ? kStagingBlobs / 2 : kStagingBlobs;
     sink_.staging = (unsigned long long) (staging_ + (size_t) (grp * per) * strata::kernels::cpu::expert_layout().max_blob);
     sink_.staging_cap = per;
@@ -1679,6 +1737,417 @@ bool Verifier::wait_commit(std::string& err) {
     return next_ == nullptr || next_->wait_commit(err);
 }
 
+
+// ================================ BATCH WINDOWS (see init_slots) ================================
+
+bool Verifier::init_slots(const std::vector<SessionState*>& slots, std::string& err) {
+    const OnDevice on_device(device_);
+    if (g_ == nullptr || ss_ == nullptr) { err = "verify: init_slots before init"; return false; }
+    if (slots.empty() || (int) slots.size() > max_t_) {
+        err = "verify: init_slots needs 1.." + std::to_string(max_t_) + " sessions";
+        return false;
+    }
+    for (SessionState* x : slots) {
+        if (x == nullptr || x->layer_lo != ss_->layer_lo || x->layer_hi != ss_->layer_hi ||
+            x->max_cells != ss_->max_cells || x->gdn_ord0 != ss_->gdn_ord0 || x->qsa_ord0 != ss_->qsa_ord0) {
+            err = "verify: a slot's session is not carved like the stage's own (layer range, context)";
+            return false;
+        }
+    }
+    const strata::kernels::QsaShapes s = shapes_of(*g_);
+    const int64_t S = (int64_t) slots.size(), CB = 2 + max_t_;
+    const int64_t TS = (s.idx_block - 1) * g_->idx_key_dim, nQ = g_->n_qsa_layers();
+    void* d = nullptr;
+    if (!mapped((size_t) (S * CB * 4 + 16), (void**) &h_commitb_, (void**) &m_commitb_)) {
+        err = "verify: the batch commit staging failed";
+        return false;
+    }
+    const uint64_t a = ((uint64_t) S * CB * 4 + 255) & ~255ull;
+    if (cudaMalloc(&d, a + (uint64_t) S * std::max<int64_t>(nQ, 1) * TS * 4) != cudaSuccess) {
+        err = "verify: the batch buffers do not fit";
+        return false;
+    }
+    arena_b_ = d;
+    commitb_ = (int32_t*) d;
+    tail_snap_b_ = (float*) ((uint8_t*) d + a);
+    slots_ = slots;
+    slot_sp_.assign(slots.size(), sampling_);   // greedy until set_slot_sampling
+    std::fprintf(stderr, "strata verify: batch windows of up to %lld sequences (layers [%lld, %lld))\n", (long long) S,
+                 (long long) lb_, (long long) le_);
+    return true;
+}
+
+bool Verifier::capture_batch(const int* rows, int S, int hbase, std::string& err) {
+    cudaGraphExec_t& ex = exec_bm_[batch_key(rows, S, hbase)];
+    if (ex != nullptr) return true;
+    if (cudaStreamBeginCapture(cs_, cudaStreamCaptureModeThreadLocal) != cudaSuccess) {
+        err = "verify: begin batch capture failed";
+        return false;
+    }
+    batch_rec_ = true;
+    row_base_ = hbase;
+    for (int t = 0; t < S; ++t) brow_[t] = rows[t];
+    std::string rerr;
+    const bool ok = record_window(S, cs_, rerr);
+    batch_rec_ = false;
+    row_base_ = 0;
+    cudaGraph_t graph = nullptr;
+    const cudaError_t ce = cudaStreamEndCapture(cs_, &graph);
+    if (!ok || ce != cudaSuccess) {
+        if (graph) cudaGraphDestroy(graph);
+        err = !ok ? rerr : std::string("verify: end batch capture: ") + cudaGetErrorString(ce);
+        return false;
+    }
+    const cudaError_t ie = cudaGraphInstantiate(&ex, graph, 0);
+    cudaGraphDestroy(graph);
+    if (ie != cudaSuccess) { err = std::string("verify: batch instantiate: ") + cudaGetErrorString(ie); return false; }
+    cudaGraphUpload(ex, cs_);
+    cudaStreamSynchronize(cs_);
+    std::string list;
+    for (int t = 0; t < S; ++t) list += (t ? "," : "") + std::to_string(rows[t]);
+    std::fprintf(stderr, "strata verify: captured the batch window over slots %s\n", list.c_str());
+    return true;
+}
+
+bool Verifier::capture_commit_batch(const int* rows, int S, int hbase, std::string& err) {
+    cudaGraphExec_t& cex = commit_bm_[batch_key(rows, S, hbase)];
+    if (cex != nullptr) return true;
+    using namespace strata::kernels;
+    const ModelGeometry& g = *g_;
+    const QsaShapes s = shapes_of(g);
+    const int64_t C = g.ssm_conv_channels, HV = g.ssm_v_heads, ZV = g.ssm_value_dim, ID = g.idx_key_dim, MT = max_t_;
+    const int64_t CB = 2 + MT, nQ = g.n_qsa_layers();
+    const uint64_t gdn_floats = (uint64_t) g.ssm_state_size * g.ssm_v_heads * g.ssm_state_size +
+                                (uint64_t) g.ssm_conv_channels * (g.ssm_d_conv - 1);
+    const int64_t TS = (s.idx_block - 1) * ID;
+    const int64_t HS = (int64_t) NG_HIST * NG_HC_DIM;
+    if (cudaStreamBeginCapture(cs_, cudaStreamCaptureModeThreadLocal) != cudaSuccess) {
+        err = "verify: begin batch commit capture failed";
+        return false;
+    }
+    bool ok = true;
+    try {
+        for (int t = 0; t < S; ++t)
+            copy_i32_from_mapped(commitb_ + (size_t) rows[t] * CB, m_commitb_ + (size_t) rows[t] * CB, CB, cs_);
+        int64_t qsa_index = 0, gdn_index = 0;
+        for (int64_t l = 0; l < lb_; ++l) (is_qsa_layer(g, l) ? qsa_index : gdn_index) += 1;
+        for (int64_t l = lb_; l < le_ && ok; ++l) {
+            const LayerView v(*wt_, l);
+            if (!is_qsa_layer(g, l)) {
+                const WeightRef* wnm = need(v, "ssm_norm.weight", err);
+                if (!wnm) { ok = false; break; }
+                for (int t = 0; t < S; ++t) {
+                    SessionState& sx = *slots_[(size_t) rows[t]];
+                    const int32_t* keep = commitb_ + (size_t) rows[t] * CB;
+                    float* state = sx.gdn_state + (size_t) (gdn_index - sx.gdn_ord0) * gdn_floats;
+                    float* conv = state + (uint64_t) g.ssm_state_size * g.ssm_v_heads * g.ssm_state_size;
+                    gdn_conv_commit(conv, qkv_L_ + (size_t) gdn_index * MT * C + (size_t) t * C, (int) C, keep, cs_);
+                    gdn_step_norm_multi(state, h_L_ + (size_t) gdn_index * MT * C + (size_t) t * C, (int) C,
+                                        gate_L_ + (size_t) gdn_index * MT * HV + (size_t) t * HV,
+                                        beta_L_ + (size_t) gdn_index * MT * HV + (size_t) t * HV, z_ + (size_t) t * ZV,
+                                        (const float*) wnm->data, EPS, y_dummy_ + (size_t) t * ZV, (int) g.ssm_k_heads,
+                                        (int) HV, 1, keep, cs_);
+                }
+                ++gdn_index;
+            } else {
+                const WeightRef* wikn = need(v, "indexer.k_norm.weight", err);
+                if (!wikn) { ok = false; break; }
+                for (int t = 0; t < S; ++t) {
+                    const QsaState& st = slots_[(size_t) rows[t]]->qsa_states[qsa_index];
+                    copy_from_mapped(st.idx_tail, tail_snap_b_ + ((size_t) rows[t] * nQ + qsa_index) * TS, TS, cs_);
+                    const QsaIndexerBuffers ib{st.idx_tail, st.idx_dead, st.idx_pooled, st.idx_block_pos};
+                    native_qsa_indexer_append(idx_raw_L_ + (size_t) (qsa_index * MT + t) * ID, commitb_ + (size_t) rows[t] * CB + 2,
+                                              0, (const float*) wikn->data, EPS, ib, s, st.max_cells, rope_scaling(), cs_);
+                }
+                ++qsa_index;
+            }
+        }
+        if (ok && ss_->ple.ready() && ple_stage())
+            for (int t = 0; t < S; ++t)
+                copy_indexed(slots_[(size_t) rows[t]]->ple_hist, hist_snap_ + (size_t) t * HS, HS,
+                             commitb_ + (size_t) rows[t] * CB + 1, HS, cs_);
+    } catch (const std::exception& e) {
+        err = std::string("verify batch commit: ") + e.what();
+        ok = false;
+    }
+    cudaGraph_t graph = nullptr;
+    const cudaError_t ce = cudaStreamEndCapture(cs_, &graph);
+    if (!ok) {
+        if (graph) cudaGraphDestroy(graph);
+        return false;
+    }
+    if (ce != cudaSuccess || cudaGraphInstantiate(&cex, graph, 0) != cudaSuccess) {
+        if (graph) cudaGraphDestroy(graph);
+        err = std::string("verify: batch commit capture: ") + cudaGetErrorString(ce);
+        return false;
+    }
+    cudaGraphDestroy(graph);
+    return true;
+}
+
+bool Verifier::stage_batch(const int* rows, int S, int hbase, const int32_t* tokens, const int64_t* pos,
+                           std::string& err) {
+    using namespace strata::kernels;
+    if (S < 1 || S > max_t_ || hbase < 0 || hbase + S > (int) slots_.size()) {
+        err = "verify: batch rows out of range (init_slots)";
+        return false;
+    }
+    for (int t = 0; t < S; ++t) {
+        bool dup = false;
+        for (int u = 0; u < t; ++u) dup = dup || rows[u] == rows[t];
+        if (rows[t] < 0 || rows[t] >= (int) slots_.size() || dup) { err = "verify: a batch row's slot is out of range or twice"; return false; }
+    }
+    if (released_.load()) { err = "verify: an earlier window never finished on the GPU (#267); restart the engine"; return false; }
+    const ModelGeometry& g = *g_;
+    for (int t = 0; t < S; ++t)
+        if (pos[t] < 0 || pos[t] + 1 > slots_[(size_t) rows[t]]->max_cells) {
+            err = "verify: slot " + std::to_string(rows[t]) + " runs past its context";
+            return false;
+        }
+    if (!capture_batch(rows, S, hbase, err) || !capture_commit_batch(rows, S, hbase, err)) return false;
+    const Clock::time_point t0 = Clock::now();
+    const QsaShapes s = shapes_of(g);
+    for (int t = 0; t < S; ++t) {
+        h_tok_[t] = tokens[t];
+        qsa_step_fill(h_step_ + t * kStepCount, pos[t], s);
+        for (int64_t h = 0; h < g.n_head; ++h) h_pos_[t * g.n_head + h] = (int32_t) pos[t];
+        int32_t* pk = h_pos_ + (size_t) max_t_ * g.n_head;
+        int32_t* pi = pk + (size_t) max_t_ * g.n_head_kv;
+        for (int64_t h = 0; h < g.n_head_kv; ++h) pk[t * g.n_head_kv + h] = (int32_t) pos[t];
+        for (int64_t h = 0; h < g.idx_q_heads; ++h) pi[t * g.idx_q_heads + h] = (int32_t) pos[t];
+    }
+    if (ss_->ple.ready() && ple_stage()) {
+        uint32_t ple_rows[kVerifyMaxT * PLE_N_HEADS];   // (not `rows`: that is the slots of the window's rows)
+        for (int t = 0; t < S; ++t) {
+            const SessionState& sx = *slots_[(size_t) rows[t]];
+            int32_t prev[2] = {sx.ple_prev[0], sx.ple_prev[1]};
+            ngram_rows(&tokens[t], prev, 1, ss_->ple.consts, ple_rows + t * PLE_N_HEADS);
+        }
+        if (!ss_->ple.table->gather_batch(ple_rows, (size_t) S, h_ple_, err)) return false;
+    }
+    // the commit's rows (a batch window keeps every row): [1, 0, position, -1 ..] per slot
+    const int64_t CB = 2 + max_t_;
+    for (int t = 0; t < S; ++t) {
+        int32_t* c = h_commitb_ + (size_t) rows[t] * CB;
+        c[0] = 1;
+        c[1] = 0;
+        c[2] = (int32_t) pos[t];
+        for (int64_t j = 3; j < CB; ++j) c[j] = -1;
+    }
+    *(volatile uint32_t*) h_seq_ = 0;
+    *(volatile uint32_t*) h_flag_ = 0;
+    *(volatile uint32_t*) h_flagA_ = 0;
+    *(volatile uint32_t*) h_flagB_ = 0;
+    std::atomic_thread_fence(std::memory_order_seq_cst);
+    last_t_ = S;
+    for (int t = 0; t < S; ++t) last_rows_[t] = rows[t];
+    row_base_ = hbase;   // the graphs' key (and nothing else) reads it until the next stage_batch
+    last_batch_ = true;
+    for (int t = 0; t < S; ++t) { last_tokens_[t] = tokens[t]; last_pos_b_[t] = pos[t]; }
+    ms_host += ms_since(t0);
+    return true;
+}
+
+bool Verifier::run_slots(int S, const int32_t* tokens, const int64_t* pos, PoolMultiFn pool, void* user, int32_t* out,
+                         std::string& err) {
+    int rows[8] = {};
+    for (int t = 0; t < S && t < 8; ++t) rows[t] = t;
+    return run_slot_rows(rows, S, tokens, pos, pool, user, out, err);
+}
+
+bool Verifier::run_slot_rows(const int* rows, int S, const int32_t* tokens, const int64_t* pos, PoolMultiFn pool,
+                             void* user, int32_t* out, std::string& err) {
+    using namespace strata::kernels;
+    const OnDevice on_device(device_);
+    const ModelGeometry& g = *g_;
+    if (!stage_batch(rows, S, 0, tokens, pos, err)) return false;
+    const cudaError_t le = cudaGraphLaunch(exec_bm_[batch_key(rows, S, 0)], cs_);
+    if (le != cudaSuccess) { err = std::string("verify: batch launch: ") + cudaGetErrorString(le); return false; }
+    (void) cudaStreamQuery(cs_);
+    volatile uint32_t* const seq = h_seq_;
+    volatile uint32_t* const flag = h_flag_;
+    const int64_t steps = le_ - lb_;
+    for (int64_t k = 0; k < steps; ++k) {
+        const int64_t l = lb_ + k;
+        const uint32_t want = (uint32_t) (k + 1);
+        const Clock::time_point a = Clock::now();
+        auto last_flush = a;
+        uint32_t spins = 0;
+        progress_at("verify batch: waiting for the GPU to reach layer", l);
+        while (*seq < want) {
+            _mm_pause();
+            if ((++spins & 1023u) != 0) continue;
+            const auto now = Clock::now();
+            if (now - last_flush > std::chrono::microseconds(2000)) {
+                last_flush = now;
+                const cudaError_t q = cudaStreamQuery(cs_);
+                if (q != cudaErrorNotReady && *seq < want) {
+                    err = "verify batch: layer " + std::to_string(l) + " never rang (" +
+                          (q == cudaSuccess ? std::string("graph finished") : std::string(cudaGetErrorString(q))) + ")";
+                    return false;
+                }
+            }
+            if (now - a > std::chrono::seconds(20)) {
+                err = "verify batch: timed out at layer " + std::to_string(l) + released_note(release_gpu_waits(5000));
+                return false;
+            }
+        }
+        const Clock::time_point b = Clock::now();
+        cur_layer_ = want - 1;
+        set_plan_slot(0);
+        progress_at("verify batch: the CPU experts of layer", l);
+        if (pool != nullptr) pool(user, h_x_, h_ids_, S, ss_->k, h_ymiss_, l);
+        progress_tick();
+        std::atomic_thread_fence(std::memory_order_seq_cst);
+        _mm_sfence();
+        if (*(volatile uint32_t*) h_flagA_ != want) {        // the pool did not publish a plan: an empty one
+            sink_.counts[0] = 0;
+            sink_.counts[1] = 0;
+            sink_.counts[2] = 0;
+            sink_.start[0] = 0;
+            sink_.start2[0] = 0;
+            std::atomic_thread_fence(std::memory_order_seq_cst);
+            *(volatile uint32_t*) h_flagA_ = want;
+            raise_flag(h_flagB_, want);
+        }
+        *flag = want;
+        ms_wait += std::chrono::duration<double, std::milli>(b - a).count();
+        ms_pool += ms_since(b);
+    }
+    const cudaError_t se = cudaStreamSynchronize(cs_);
+    if (se != cudaSuccess) { err = std::string("verify batch: ") + cudaGetErrorString(se); return false; }
+    cudaStreamSynchronize(copy_);
+    if (prof_on_) collect_profile();
+    ++windows;
+    if (le_ < g.n_layers) return next_ == nullptr || next_->run_slot_rows(rows, S, tokens, pos, pool, next_user_, out, err);
+    if (!sample_rows(S, err)) return false;
+    for (int t = 0; t < S; ++t) out[t] = ((volatile int32_t*) h_out_)[t];
+    progress_at("decode");
+    progress_beat();
+    return true;
+}
+
+bool Verifier::commit_slots(std::string& err) {
+    const OnDevice on_device(device_);
+    if (!last_batch_ || last_t_ < 1) { err = "verify: commit_slots without a batch window"; return false; }
+    const int S = last_t_;
+    const Clock::time_point t0 = Clock::now();
+    std::atomic_thread_fence(std::memory_order_seq_cst);
+    const cudaError_t le = cudaGraphLaunch(commit_bm_[batch_key(last_rows_, S, row_base_)], cs_);
+    if (le != cudaSuccess) { err = std::string("verify: batch commit launch: ") + cudaGetErrorString(le); return false; }
+    const cudaError_t se = cudaStreamSynchronize(cs_);
+    if (se != cudaSuccess) { err = std::string("verify: batch commit: ") + cudaGetErrorString(se); return false; }
+    if (ple_stage())
+        for (int t = 0; t < S; ++t) {
+            SessionState& sx = *slots_[(size_t) last_rows_[t]];
+            sx.ple_prev[0] = sx.ple_prev[1];
+            sx.ple_prev[1] = last_tokens_[t];
+        }
+    ms_commit += ms_since(t0);
+    return next_ == nullptr || next_->commit_slots(err);
+}
+
+bool Verifier::sample_rows(int S, std::string& err) {
+    bool any = false;
+    for (int t = 0; t < S; ++t) {
+        strata::kernels::SamplerParams sp = slot_sp_[(size_t) last_rows_[t]];
+        if (sp.greedy || sp.temperature <= 0.0f) continue;
+        sp.counter = (uint64_t) last_pos_b_[t];   // Philox(seed, position): the solo window's draw for this position
+        sp.penalty_last_n = 0;
+        strata::kernels::sample_tokens(head_logits_ + (size_t) t * (size_t) n_vocab_, 1, (int) n_vocab_, nullptr, 0, sp,
+                                       m_out_ + t, cs_);
+        any = true;
+    }
+    if (any && cudaStreamSynchronize(cs_) != cudaSuccess) { err = "verify batch: the row sampling failed"; return false; }
+    return true;
+}
+
+bool Verifier::batch_launch(int base, int S, const int32_t* tokens, const int64_t* pos, std::string& err) {
+    const OnDevice on_device(device_);
+    if (b_running_) { err = "verify: batch_launch while this stage is busy"; return false; }
+    int rows[8] = {};
+    for (int t = 0; t < S && t < 8; ++t) rows[t] = base + t;
+    if (!stage_batch(rows, S, base, tokens, pos, err)) return false;
+    cudaError_t le = cudaGraphLaunch(exec_bm_[batch_key(rows, S, base)], cs_);
+    if (le == cudaSuccess) le = cudaGraphLaunch(commit_bm_[batch_key(rows, S, base)], cs_);   // right behind it: every row is kept
+    if (le != cudaSuccess) { err = std::string("verify: batch launch: ") + cudaGetErrorString(le); return false; }
+    (void) cudaStreamQuery(cs_);
+    if (ple_stage())   // the host's side of the commit (the hash's last two tokens)
+        for (int t = 0; t < S; ++t) {
+            SessionState& sx = *slots_[(size_t) rows[t]];
+            sx.ple_prev[0] = sx.ple_prev[1];
+            sx.ple_prev[1] = tokens[t];
+        }
+    b_running_ = true;
+    b_k_ = 0;
+    b_steps_ = le_ - lb_;
+    b_last_ = Clock::now();
+    return true;
+}
+
+int Verifier::batch_poll(PoolMultiFn pool, void* user, std::string& err) {
+    if (!b_running_) return 1;
+    const OnDevice on_device(device_);
+    volatile uint32_t* const seq = h_seq_;
+    const int S = last_t_;
+    while (b_k_ < b_steps_) {
+        const uint32_t want = (uint32_t) (b_k_ + 1);
+        if (*seq < want) {
+            const auto now = Clock::now();
+            if (now - b_last_ > std::chrono::milliseconds(2)) {
+                const cudaError_t q = cudaStreamQuery(cs_);
+                if (q != cudaErrorNotReady && *seq < want) {
+                    err = "verify batch: layer " + std::to_string(lb_ + b_k_) + " never rang (" +
+                          (q == cudaSuccess ? std::string("graph finished") : std::string(cudaGetErrorString(q))) + ")";
+                    b_running_ = false;
+                    return -1;
+                }
+                if (now - b_last_ > std::chrono::seconds(20)) {
+                    err = "verify batch: timed out at layer " + std::to_string(lb_ + b_k_) + released_note(release_gpu_waits(5000));
+                    b_running_ = false;
+                    return -1;
+                }
+            }
+            return 0;
+        }
+        const Clock::time_point b = Clock::now();
+        cur_layer_ = want - 1;
+        set_plan_slot(0);
+        if (pool != nullptr) pool(user, h_x_, h_ids_, S, ss_->k, h_ymiss_, lb_ + b_k_);
+        progress_tick();
+        std::atomic_thread_fence(std::memory_order_seq_cst);
+        _mm_sfence();
+        if (*(volatile uint32_t*) h_flagA_ != want) {        // the pool did not publish a plan: an empty one
+            sink_.counts[0] = 0;
+            sink_.counts[1] = 0;
+            sink_.counts[2] = 0;
+            sink_.start[0] = 0;
+            sink_.start2[0] = 0;
+            std::atomic_thread_fence(std::memory_order_seq_cst);
+            *(volatile uint32_t*) h_flagA_ = want;
+            raise_flag(h_flagB_, want);
+        }
+        *(volatile uint32_t*) h_flag_ = want;
+        ms_pool += ms_since(b);
+        b_last_ = Clock::now();
+        ++b_k_;
+    }
+    const cudaError_t q = cudaStreamQuery(cs_);
+    if (q == cudaErrorNotReady) return 0;
+    if (q != cudaSuccess) { err = std::string("verify batch: ") + cudaGetErrorString(q); b_running_ = false; return -1; }
+    const cudaError_t qc = cudaStreamQuery(copy_);   // no host function of this window may raise flag B in the next
+    if (qc == cudaErrorNotReady) return 0;
+    if (prof_on_) collect_profile();
+    if (last_stage()) {
+        if (!sample_rows(S, err)) { b_running_ = false; return -1; }
+        for (int t = 0; t < S; ++t) b_out_[t] = ((volatile int32_t*) h_out_)[t];
+    }
+    ++windows;
+    b_running_ = false;
+    progress_beat();
+    return 1;
+}
+
 bool Verifier::copy_logits(int t, float* host) const {
     if (next_ != nullptr) return next_->copy_logits(t, host);   // a layer split: the head is on the last stage
     if (head_logits_ == nullptr || host == nullptr || t < 0 || n_vocab_ <= 0) return false;

tools/batch_interleave_test.py

@@ -0,0 +1,237 @@
+#!/usr/bin/env python3
+"""Batch slots, the parts batch_test.py does not reach (#465): a long prompt read WHILE other slots decode (its
+chunks interleaved with their windows), and a conversation's next turn continued from the slot that holds it.
+
+  1. solo references (GEN): A, then A's next turn (from the live session), B, C (a long prompt)
+  2. the same in the slots: A in slot 0, B in slot 1 next to it, then C in slot 2 while A and B decode - every slot's
+     tokens must equal its solo tokens
+  3. A's next turn (BGEN into slot 0, which still holds A): read from the slot, its tokens must equal the solo next turn
+  4. a long prompt D gives way (BYIELD) at its first chunk boundary, a short E is admitted, D goes on from its slot
+     while E decodes: both equal their solo tokens
+  6. A in a slot, stopped after 50 tokens and continued on the solo path (GEN of A + its tokens, from the slot): the
+     whole equals solo A (what the server does with a request left alone in a slot)
+  5. (a measurement) a solo next turn continued from a slot: the drafts accepted
+  7. F, A's history as a client sends it back without the reply's thinking, from slot 1's turn checkpoint: equal to
+     F solo (which continues from the same checkpoint of the live chain)
+
+Exact comparisons need the same settings as batch_test.py (this script sets STRATA_IQ_MT_MIN=1):
+  python tools/batch_interleave_test.py --exe build/strata --config strata-<model>.json \\
+      --extra "--pcie-frac 0 --adapt-every 1000000 --no-prefill-borrow"
+Without --no-prefill-borrow the slots decoding during C's read see the expert cache without the slots C's prompt
+borrowed (those experts run on the CPU, which rounds differently): A's and B's text may then drift from solo.
+"""
+import argparse, json, os, sys, time
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parent))
+from batch_test import Engine, tokenizer  # noqa: E402
+
+LONG = ("The history of mathematics is long and full of surprising turns. Early civilizations counted with tally marks, "
+        "then with symbols, and later with place-value systems that made arithmetic far easier. ")
+
+
+def run(eng, out, line, slot=None):
+    """Send one GEN/BGEN; returns (tokens of the request line, the BADM flag) - BT/BDONE of other slots are kept."""
+    eng.send(line)
+    got = []
+    for l in out:
+        if l.startswith("T "):
+            got.append(int(l.split()[1]))
+        elif l.startswith(("BT ", "BDONE ")):
+            eng.pending.append(l)
+        elif l.startswith("ERR"):
+            raise SystemExit("engine: " + l)
+        elif slot is None and l.startswith("DONE"):
+            return got, None
+        elif slot is not None and l.startswith("BADM "):
+            return got, l.split()[2] == "1"
+    raise SystemExit(f"the engine ended during: {line[:40]} (see {eng.log_path})")
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--exe", required=True)
+    ap.add_argument("--config", required=True)
+    ap.add_argument("--max-new", type=int, default=200)
+    ap.add_argument("--long", type=int, default=5000, help="tokens of the long prompt C (several prompt chunks)")
+    ap.add_argument("--extra", default="")
+    a = ap.parse_args()
+    cfg = json.loads(Path(a.config).read_text())
+    tok = tokenizer(cfg["tokenizer"])
+
+    def chat(q):
+        return tok.encode(f"<|im_start|>user\n{q}<|im_end|>\n<|im_start|>assistant\n<think>\n\n</think>\n\n",
+                          parse_special=True)
+    body = LONG * (a.long // max(1, len(tok.encode(LONG))) + 1)
+    A = chat("Write a Python function that merges overlapping intervals, then explain it.")
+    B = chat("Compare TCP and QUIC: handshake, congestion control, multiplexing.")
+    C = chat(body + "\nSummarize the text above in three sentences.")
+    D = chat(body + "\nWhat are the three most important ideas in the text above?")
+    E = chat("Describe the life cycle of a star like the Sun.")
+    print(f"prompts: A {len(A)}, B {len(B)}, C {len(C)} tokens", flush=True)
+    eng = Engine(a.exe, cfg, 3, {"STRATA_IQ_MT_MIN": "1"}, a.extra.split())
+    eng.pending = []
+    out = eng.lines()
+    ids = lambda v: ",".join(map(str, v))
+    M = a.max_new
+
+    # 1. solo
+    sA, _ = run(eng, out, f"GEN {M} {ids(A)}")
+    A2 = A + sA + chat("Now do the same in Rust.")
+    sA2, _ = run(eng, out, f"GEN {M} {ids(A2)}")
+    # F: A's history as a client sends it back WITHOUT the reply's thinking - it shares A only up to A's last turn
+    # boundary (the checkpoint there); solo it continues from that checkpoint of the live chain
+    turn_at = max(i for i, t in enumerate(A) if t == A[0])          # A[0] is <|im_start|>
+    F = A[:turn_at] + tok.encode("<|im_start|>assistant\nA short answer.<|im_end|>\n", parse_special=True) + \
+        chat("Now in Go.")
+    sF, _ = run(eng, out, f"GEN {M} {ids(F)}")
+    sB, _ = run(eng, out, f"GEN {M} {ids(B)}")
+    sC, _ = run(eng, out, f"GEN {M} {ids(C)}")
+    sD, _ = run(eng, out, f"GEN {M} {ids(D)}")
+    sE, _ = run(eng, out, f"GEN {M} {ids(E)}")
+    print(f"solo: A {len(sA)}, A2 {len(sA2)}, B {len(sB)}, C {len(sC)}, D {len(sD)}, E {len(sE)} tokens", flush=True)
+
+    # 2. the slots: A, B, then the long C while they decode
+    got = {0: [], 1: [], 2: []}
+    done = {}
+    t0 = time.time()
+    for s, P in ((0, A), (1, B), (2, C)):
+        first, cont = run(eng, out, f"BGEN {s} {M} {ids(P)}", slot=s)
+        got[s] += first
+        if not cont:
+            done[s] = True
+    t_admit = time.time() - t0
+    during_c = sum(1 for l in eng.pending if l.startswith("BT "))   # windows the slots ran while C was admitted
+
+    def drain():
+        while eng.pending or len(done) < 3:
+            l = eng.pending.pop(0) if eng.pending else next(out)
+            if l.startswith("BT "):
+                _, s, y = l.split()
+                got[int(s)].append(int(y))
+            elif l.startswith("BDONE "):
+                done[int(l.split()[1])] = True
+            elif l.startswith("ERR"):
+                raise SystemExit("engine: " + l)
+    drain()
+    print(f"slots: admissions {t_admit:.1f}s; {during_c} slot tokens arrived while the admissions read their prompts",
+          flush=True)
+    ok = True
+    for s, ref, name in ((0, sA, "A"), (1, sB, "B"), (2, sC, "C")):
+        same = got[s] == ref
+        ok &= same
+        d = next((k for k in range(min(len(got[s]), len(ref))) if got[s][k] != ref[k]), None)
+        print(f"slot {s} ({name}): {len(got[s])} tokens, solo {len(ref)}: "
+              f"{'IDENTICAL' if same else f'DIFFERS at {d}'}", flush=True)
+
+    # 3. A's next turn from slot 0 (the engine's log says "slot 0 gave back ...")
+    first, cont = run(eng, out, f"BGEN 0 {M} {ids(A2)}", slot=0)
+    got2, done = first, {}
+    if cont:
+        while True:
+            l = eng.pending.pop(0) if eng.pending else next(out)
+            if l.startswith("BT 0 "):
+                got2.append(int(l.split()[2]))
+            elif l.startswith("BDONE 0 "):
+                break
+    same = got2 == sA2
+    ok &= same
+    d = next((k for k in range(min(len(got2), len(sA2))) if got2[k] != sA2[k]), None)
+    print(f"A's next turn from its slot: {len(got2)} tokens, solo {len(sA2)}: "
+          f"{'IDENTICAL' if same else f'DIFFERS at {d}'}", flush=True)
+    # 4. a long prompt D gives way (BYIELD) to a short E at its first chunk boundary, then goes on from its slot
+    eng.send(f"BGEN 2 {M} {ids(D)}")
+    eng.send("BYIELD 2")
+    yielded = None
+    for l in out:
+        if l.startswith("YIELDED "):
+            yielded = l
+        elif l.startswith(("BT ", "BDONE ")):
+            eng.pending.append(l)
+        elif l.startswith("ERR"):
+            raise SystemExit("engine: " + l)
+        elif l.startswith("BADM 2 "):
+            break
+    print(f"D gave way: {yielded}", flush=True)
+    ok &= yielded is not None
+    got4 = {1: [], 2: []}
+    first, cont1 = run(eng, out, f"BGEN 1 {M} {ids(E)}", slot=1)
+    got4[1] += first
+    first, cont2 = run(eng, out, f"BGEN 2 {M} {ids(D)}", slot=2)
+    got4[2] += first
+    done4 = {s for s, c in ((1, cont1), (2, cont2)) if not c}
+    while len(done4) < 2:
+        l = eng.pending.pop(0) if eng.pending else next(out)
+        if l.startswith("BT "):
+            _, s, y = l.split()
+            if int(s) in got4:
+                got4[int(s)].append(int(y))
+        elif l.startswith("BDONE "):
+            done4.add(int(l.split()[1]))
+    for s, ref, name in ((1, sE, "E (short, admitted while D waited)"), (2, sD, "D (gave way, then went on)")):
+        same = got4[s] == ref
+        ok &= same
+        d = next((k for k in range(min(len(got4[s]), len(ref))) if got4[s][k] != ref[k]), None)
+        print(f"slot {s} {name}: {len(got4[s])} tokens, solo {len(ref)}: "
+              f"{'IDENTICAL' if same else f'DIFFERS at {d}'}", flush=True)
+    # 6. back to the solo path (what the server does with a request left alone in a slot): A in slot 1, BSTOP after
+    # 50 tokens, then GEN of A + what it produced - the engine continues from the slot with MTP drafts again
+    first, cont = run(eng, out, f"BGEN 1 {M} {ids(A)}", slot=1)
+    got6, sent = list(first), False
+    while cont:
+        l = eng.pending.pop(0) if eng.pending else next(out)
+        if l.startswith("BT 1 "):
+            got6.append(int(l.split()[2]))
+            if len(got6) >= 50 and not sent:
+                eng.send("BSTOP 1")
+                sent = True
+        elif l.startswith("BDONE 1 "):
+            break
+    tail6, _ = run(eng, out, f"GEN {M - len(got6)} {ids(A + got6)}") if len(got6) < M else ([], None)
+    same = got6 + tail6 == sA
+    ok &= same
+    both = got6 + tail6
+    d = next((k for k in range(min(len(both), len(sA))) if both[k] != sA[k]), None)
+    print(f"A in a slot, then solo again after {len(got6)} tokens: {len(both)} tokens, solo {len(sA)}: "
+          f"{'IDENTICAL' if same else f'DIFFERS at {d}'}", flush=True)
+
+    # 5. (a measurement, not a check) a SOLO next turn continued from slot 0: its tokens are exact either way, but the
+    # draft layer's own K/V was built for another conversation, so fewer drafts may be accepted than in a solo next
+    # turn whose drafter read the conversation (A2 in step 1)
+    A3 = A2 + got2 + chat("Now in Go.")
+    eng.send(f"GEN {M} {ids(A3)}")
+    for l in out:
+        if l.startswith(("BT ", "BDONE ")):
+            eng.pending.append(l)
+        elif l.startswith("DONE"):
+            f = l.split()
+            print(f"solo next turn from a slot: drafts accepted {f[6]} of {f[7]} ({f[1]} tokens in {f[4]} ms)",
+                  flush=True)
+            break
+    # 7. F from slot 1's turn checkpoint (slot 1 holds A since step 6): the engine log says "(its turn checkpoint)"
+    first, cont = run(eng, out, f"BGEN 2 {M} {ids(F)}", slot=2)
+    got7 = list(first)
+    while cont:
+        l = eng.pending.pop(0) if eng.pending else next(out)
+        if l.startswith("BT 2 "):
+            got7.append(int(l.split()[2]))
+        elif l.startswith("BDONE 2 "):
+            break
+    same = got7 == sF
+    ok &= same
+    d = next((k for k in range(min(len(got7), len(sF))) if got7[k] != sF[k]), None)
+    print(f"F (A's history without the reply's thinking) from slot 1's turn checkpoint: {len(got7)} tokens, solo "
+          f"{len(sF)}: {'IDENTICAL' if same else f'DIFFERS at {d}'}", flush=True)
+    eng.send("QUIT")
+    eng.p.wait(timeout=180)
+    log = Path(eng.log_path).read_text(errors="replace")
+    for l in log.splitlines():
+        if "drafts accepted" in l:
+            print("  log:", l.split("strata serve: ")[-1][:160], flush=True)
+    for key in ("gave back", "its turn checkpoint", "takes", "the prompt was read in", "gives way"):
+        print(f"engine log '{key}': {log.count(key)} lines", flush=True)
+    return 0 if ok else 2
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/batch_test.py

@@ -0,0 +1,161 @@
+#!/usr/bin/env python3
+"""Batching test: the same prompts decoded alone (GEN, the usual path with drafts) and together in the batch
+windows (BGEN), greedy.  Every slot's tokens must equal its solo tokens; prints the aggregate decode rate.
+
+  python3 tools/batch_test.py --exe engine/strata --config strata-<model>.json --batch 4 --n 4 \
+      --extra "--layer-split 12,24,36 --trim-stage-weights --pcie-frac 0 --adapt-every 1000000"
+"""
+import argparse, json, os, subprocess, sys, tempfile, threading, time
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parent))
+import strata_tokenizer as ST  # noqa: E402
+
+QUESTIONS = [
+    "Explique en detail le fonctionnement d'un B-tree.",
+    "Write a Python function that merges overlapping intervals, then explain it.",
+    "Compare TCP and QUIC: handshake, congestion control, multiplexing.",
+    "Raconte l'histoire du calcul de pi, d'Archimede a nos jours.",
+    "What are the main causes of the French Revolution?",
+    "Explain how a transformer's attention works, step by step.",
+    "Donne une recette de tarte aux pommes, etape par etape.",
+    "Describe the life cycle of a star like the Sun.",
+]
+
+
+def tokenizer(path):
+    t = Path(path)
+    vocab = json.loads((t / "vocab.json").read_text(encoding="utf-8"))
+    tokens = [None] * len(vocab)
+    for s, i in vocab.items():
+        tokens[i] = s
+    merges = (t / "merges.txt").read_text(encoding="utf-8").split("\n")
+    types = json.loads((t / "token_type.json").read_text())
+    return ST.Tokenizer(tokens, merges, types)
+
+
+class Engine:
+    def __init__(self, exe, cfg, batch, extra_env, extra_args):
+        args = list(cfg["args"])
+        if len(cfg.get("gpu") or []) > 1:
+            args += ["--layer-split", str(cfg.get("layer_split") or "auto")]
+        if batch:
+            args += ["--batch", str(batch)]
+        args += extra_args
+        env = dict(os.environ, **extra_env)
+        env["LD_LIBRARY_PATH"] = ":".join(cfg.get("lib_dirs", []) + [env.get("LD_LIBRARY_PATH", "")])
+        if os.name == "nt":
+            env["PATH"] = os.pathsep.join(cfg.get("lib_dirs", []) + [env.get("PATH", "")])
+        self.log_path = os.environ.get("BATCH_TEST_LOG") or os.path.join(tempfile.gettempdir(), "batch_test_engine.log")
+        self.log = open(self.log_path, "w")
+        self.p = subprocess.Popen([exe, "--serve", *args], cwd=cfg.get("cwd"), stdin=subprocess.PIPE,
+                                  stdout=subprocess.PIPE, stderr=self.log, text=True, bufsize=1, env=env)
+        for line in self.p.stdout:
+            if line.startswith("READY"):
+                return
+        raise SystemExit(f"the engine ended before READY - see {self.log_path}")
+
+    def send(self, line):
+        self.p.stdin.write(line + "\n")
+        self.p.stdin.flush()
+
+    def lines(self):
+        for line in self.p.stdout:
+            yield line.rstrip("\n")
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--exe", required=True)
+    ap.add_argument("--config", required=True)
+    ap.add_argument("--batch", type=int, default=4)
+    ap.add_argument("--n", type=int, default=4, help="prompts (<= --batch)")
+    ap.add_argument("--max-new", type=int, default=64)
+    ap.add_argument("--skip-solo", action="store_true")
+    ap.add_argument("--keys", default="", help='sampling keys for every request, e.g. "temperature=0.7 top_k=20"')
+    ap.add_argument("--extra", default="", help='more engine arguments in one string, e.g. "--adapt-every 1000000"')
+    ap.add_argument("--mt-min", default="1", help="STRATA_IQ_MT_MIN for the engine (1: exact; empty: the default)")
+    a = ap.parse_args()
+    cfg = json.loads(Path(a.config).read_text())
+    tok = tokenizer(cfg["tokenizer"])
+    prompts = []
+    for q in QUESTIONS[: a.n]:
+        text = f"<|im_start|>user\n{q}<|im_end|>\n<|im_start|>assistant\n<think>\n\n</think>\n\n"
+        prompts.append(tok.encode(text, parse_special=True))
+    env = {"STRATA_DECODE_TIMING": "1", **({"STRATA_VERIFY_PROFILE": "1"} if os.environ.get("PROF") else {})}
+    if a.mt_min:
+        env["STRATA_IQ_MT_MIN"] = a.mt_min
+    eng = Engine(a.exe, cfg, a.batch, env, a.extra.split())
+    out = eng.lines()
+
+    solo = []
+    if not a.skip_solo:
+        for i, ids in enumerate(prompts):
+            eng.send(" ".join(x for x in ("GEN", str(a.max_new), a.keys, ",".join(map(str, ids))) if x))
+            got, t0 = [], time.time()
+            for line in out:
+                if line.startswith("T "):
+                    got.append(int(line.split()[1]))
+                elif line.startswith("DONE") or line.startswith("ERR"):
+                    print(f"solo {i}: {len(got)} tokens in {time.time() - t0:.1f}s  {line[:60]}", flush=True)
+                    break
+            solo.append(got)
+
+    # batch: admit every prompt, then the slots decode together
+    got = {i: [] for i in range(len(prompts))}
+    done = {}
+    t_admit = time.time()
+    first_bt = None
+    for i, ids in enumerate(prompts):
+        eng.send(" ".join(x for x in ("BGEN", str(i), str(a.max_new), a.keys, ",".join(map(str, ids))) if x))
+        for line in out:
+            if line.startswith("T "):
+                got[i].append(int(line.split()[1]))
+            elif line.startswith("BT "):
+                _, s, y = line.split()
+                got[int(s)].append(int(y))
+                first_bt = first_bt or time.time()
+            elif line.startswith("BDONE "):
+                done[int(line.split()[1])] = line
+            elif line.startswith("ERR"):
+                print("batch:", line)
+                return 1
+            elif line.startswith("BADM "):
+                if line.split()[2] == "0":
+                    done[i] = line
+                break
+    t_admitted = time.time()
+    for line in out:
+        if line.startswith("BT "):
+            _, s, y = line.split()
+            got[int(s)].append(int(y))
+            first_bt = first_bt or time.time()
+        elif line.startswith("BDONE "):
+            done[int(line.split()[1])] = line
+            if len(done) == len(prompts):
+                break
+        elif line.startswith("ERR"):
+            print("batch:", line)
+            return 1
+    t_end = time.time()
+    total = sum(len(v) for v in got.values())
+    print(f"batch: {len(prompts)} slots, {total} tokens; admissions {t_admitted - t_admit:.1f}s; "
+          f"then {t_end - t_admitted:.1f}s -> aggregate {total / max(t_end - t_admit, 1e-9):.1f} tok/s overall, "
+          f"{sum(len(v) for v in got.values()) / max(t_end - (first_bt or t_admit), 1e-9):.1f} tok/s from the first batch token",
+          flush=True)
+    ok = True
+    for i in range(len(prompts)):
+        b = got[i]
+        if solo:
+            same = b == solo[i]
+            ok &= same
+            first_diff = next((k for k in range(min(len(b), len(solo[i]))) if b[k] != solo[i][k]), None)
+            print(f"slot {i}: {len(b)} tokens, solo {len(solo[i])}: {'IDENTICAL' if same else f'DIFFERS at {first_diff}'}")
+        print("   ", ascii(tok.decode(b)[:160]))
+    eng.send("QUIT")
+    eng.p.wait(timeout=180)   # the next run needs the GPUs back
+    return 0 if ok else 2
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/early_close_test.py

@@ -0,0 +1,47 @@
+#!/usr/bin/env python3
+"""A client that stops reading a streamed answer early must not leave its tokens to the next request.
+
+  STRATA_KEY=... python3 tools/early_close_test.py http://127.0.0.1:8090
+"""
+import json, os, sys, threading, urllib.request
+
+URL = sys.argv[1].rstrip("/") + "/v1/chat/completions"
+KEY = os.environ.get("STRATA_KEY", "")
+
+
+def req(content, stream, max_tokens=300):
+    body = json.dumps({"model": "m", "max_tokens": max_tokens, "temperature": 0, "stream": stream,
+                       "messages": [{"role": "user", "content": content}],
+                       "chat_template_kwargs": {"enable_thinking": False}}).encode()
+    return urllib.request.Request(URL, data=body, headers={"Content-Type": "application/json",
+                                                           "Authorization": f"Bearer {KEY}"})
+
+
+def early_close(content, keep_lines=5):
+    r = urllib.request.urlopen(req(content, True), timeout=600)
+    for _ in range(keep_lines):
+        r.readline()
+    r.close()                                   # the client goes away mid-answer
+
+
+def answer(content):
+    with urllib.request.urlopen(req(content, False, 60), timeout=600) as r:
+        m = json.load(r)["choices"][0]["message"]
+        return (m.get("content") or "") + " " + (m.get("reasoning_content") or m.get("reasoning") or "")
+
+
+ok = True
+# alone (solo path), then the same with a second request running (batch slots)
+for label, background in (("solo", False), ("batch", True)):
+    bg = None
+    if background:
+        bg = threading.Thread(target=answer, args=("Count slowly from 1 to 60, one number per line.",))
+        bg.start()
+    early_close("Describe the weather in Montpellier in great detail, at length.")
+    a = answer("Reply with exactly one word: the color of the sky on a clear day.")
+    good = "blue" in a.lower() and "montpellier" not in a.lower()
+    ok &= good
+    print(f"{label}: {'OK' if good else 'STALE'} -> {a[:120]!r}", flush=True)
+    if bg:
+        bg.join()
+sys.exit(0 if ok else 1)

tools/parking_test.py

@@ -0,0 +1,89 @@
+#!/usr/bin/env python3
+"""Conversation parking test: a follow-up to conversation A must decode the same tokens whether A's state was still
+live (reference engine: A, then the follow-up) or came back from the parking cache (second engine: A, then B, which
+parks A, then the follow-up).  Greedy; run with --extra "--pcie-frac 0 --adapt-every 1000000" for exactness.
+
+  python3 tools/parking_test.py --exe engine/strata --config strata-<model>.json \
+      --extra "--layer-split 12,24,36 --conversation-cache-mib 8192 --conversation-cache-slots 4 --pcie-frac 0"
+"""
+import argparse, json, re, sys, time
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parent))
+from batch_test import Engine, tokenizer  # noqa: E402
+
+DOC = ("The history of the printing press begins with movable type in East Asia and continues with Gutenberg's "
+       "press in Mainz around 1450, which spread across Europe within decades, lowered the cost of books, changed "
+       "how scholars, merchants and churches shared knowledge, and fed the Reformation and the scientific "
+       "revolution. ")
+
+
+def chat(user):
+    return f"<|im_start|>user\n{user}<|im_end|>\n<|im_start|>assistant\n<think>\n\n</think>\n\n"
+
+
+def gen(eng, out, ids, n):
+    eng.send(f"GEN {n} " + ",".join(map(str, ids)))
+    got = []
+    for line in out:
+        if line.startswith("T "):
+            got.append(int(line.split()[1]))
+        elif line.startswith("DONE") or line.startswith("ERR"):
+            return got, line
+    raise SystemExit("the engine ended")
+
+
+def run(a, cfg, tok, with_b):
+    # STRATA_SNAPSHOT_VERIFY: the engine reads the restored draft ring back (on the GPU that holds it) after a restore
+    eng = Engine(a.exe, cfg, 0, {"STRATA_IQ_MT_MIN": "1", "STRATA_SNAPSHOT_VERIFY": "1"}, a.extra.split())
+    out = eng.lines()
+    pa = tok.encode(chat(DOC * a.repeat + "\nSummarize this text in five sentences."), parse_special=True)
+    ans, _ = gen(eng, out, pa, a.max_new)
+    if with_b:
+        pb = tok.encode(chat("Write a short poem about the sea, then explain its metaphors." * 20), parse_special=True)
+        gen(eng, out, pb, a.max_new)
+    end = tok.encode("<|im_end|>\n", parse_special=True)
+    follow = pa + ans + ([] if ans and ans[-1] in end else end) + \
+        tok.encode(chat("Now give three keywords for that text, with one sentence each.")[0:], parse_special=True)
+    t0 = time.time()
+    got, done = gen(eng, out, follow, a.max_new)
+    dt = time.time() - t0
+    if with_b:   # B again: A is parked a second time, from retained K/V (every stage's) - only its new tail copied
+        gen(eng, out, pb, 8)
+    eng.send("QUIT")
+    eng.p.wait(timeout=180)
+    log = open("/tmp/batch_test_engine.log").read()
+    restored = re.findall(r"restored \d+ tokens[^\n]*", log)
+    if with_b and "SNAPSHOT_VERIFY draft=" not in log:
+        print("no SNAPSHOT_VERIFY line: the draft ring was not read back", flush=True)
+    reparked = [int(x) for x in re.findall(r"parked \d+ tokens .*reused_kv_bytes=(\d+)", log)]
+    return len(pa), len(follow), got, done, dt, restored, reparked
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--exe", required=True)
+    ap.add_argument("--config", required=True)
+    ap.add_argument("--repeat", type=int, default=60, help="copies of the paragraph in conversation A")
+    ap.add_argument("--max-new", type=int, default=120)
+    ap.add_argument("--extra", default="")
+    a = ap.parse_args()
+    cfg = json.loads(Path(a.config).read_text())
+    tok = tokenizer(cfg["tokenizer"])
+    ref = run(a, cfg, tok, with_b=False)
+    park = run(a, cfg, tok, with_b=True)
+    for name, r in (("live (reference)", ref), ("parked (A, B, A)", park)):
+        f = r[3].split()
+        print(f"{name}: prompt A {r[0]} tokens, follow-up {r[1]} tokens, reused {f[8] if len(f) > 8 else '?'}, "
+              f"follow-up in {r[4]:.2f} s; restore lines: {r[5][-1:] or 'none'}")
+    same = ref[2] == park[2]
+    first = next((k for k in range(min(len(ref[2]), len(park[2]))) if ref[2][k] != park[2][k]), None)
+    print("follow-up tokens:", "IDENTICAL" if same else f"DIFFER at {first}", f"({len(ref[2])} / {len(park[2])})")
+    print("   ", repr(tok.decode(park[2])[:200]))
+    reused = park[6][-1] if park[6] else 0
+    print(f"second park of A: {reused} bytes of K/V reused (retained from its restore)")
+    return 0 if same and park[5] and reused > 0 else 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/test_setup_parallel.py

@@ -0,0 +1,53 @@
+"""#465: setup's recommendation for "parallel" (several requests decoded together in the engine's batch slots).
+No GPU or download needed: python tools/test_setup_parallel.py"""
+import sys
+import unittest
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parent.parent))
+import setup as S  # noqa: E402
+
+
+class Parallel(unittest.TestCase):
+    def test_slot_size(self):
+        # measured: a slot's session at 32K, int8 KV: 0.56 GiB = 0.60 GB
+        self.assertAlmostEqual(S.parallel_slot_gb(32768, "int8", False), 0.59, delta=0.03)
+        # KV streaming keeps only the attention's 32K positions in VRAM, whatever the context
+        self.assertEqual(S.parallel_slot_gb(262144, "int8", True), S.parallel_slot_gb(32768, "int8", False))
+        self.assertLess(S.parallel_slot_gb(32768, "q4_0", False), S.parallel_slot_gb(32768, "int8", False))
+
+    def test_recommendation_only_where_the_experts_mostly_fit(self):
+        q2, iq3s = S.MODELS["Q2_0"]["arena_gb"], S.MODELS["IQ3_S"]["arena_gb"]
+        # most experts on the CPU: slots cost speed per request (measured on the 12 GB 5070) - none recommended
+        self.assertEqual(S.parallel_recommend(12, q2, 32768, "int8", False), 0)
+        self.assertEqual(S.parallel_recommend(16, q2, 32768, "int8", False), 0)
+        self.assertEqual(S.parallel_recommend(24, iq3s, 32768, "int8", False), 0)
+        # the cache holds most of the experts beside the slots: recommended
+        self.assertEqual(S.parallel_recommend(24, q2, 32768, "int8", False), 3)
+        self.assertEqual(S.parallel_recommend(48, q2, 32768, "int8", False), 4)       # capped at 4
+        self.assertEqual(S.parallel_recommend([24, 24], iq3s, 32768, "int8", False), 4)   # a layer split: every card
+        self.assertEqual(S.parallel_recommend([16, 16], q2, 32768, "int8", False), 4)
+        self.assertEqual(S.parallel_recommend(32, q2, 262144, "int8", False), 0)      # a 262K KV per slot: no
+        self.assertEqual(S.parallel_recommend(32, q2, 262144, "int8", True), 4)       # ... unless it streams
+
+    def test_notes_recommend_never_force(self):
+        q2 = S.MODELS["Q2_0"]["arena_gb"]
+        # 12 GB: left at one, with the reason
+        note = S.parallel_note(None, 12, q2, 32768, "int8", False)
+        self.assertEqual(len(note), 1)
+        self.assertIn("one at a time", note[0])
+        self.assertIn("costs about 10-25% speed per request on this card", note[0])
+        self.assertIn("--parallel 3", S.parallel_note(None, 24, q2, 32768, "int8", False)[0])
+        # asked on 12 GB: kept, with the cost said
+        lines = S.parallel_note(4, 12, q2, 32768, "int8", False)
+        self.assertIn("4 at once", lines[0])
+        self.assertTrue(any("one at a time" in x and "kept as you chose" in x and "10-25%" in x for x in lines))
+        # asked above the recommendation where it pays
+        lines = S.parallel_note(4, 24, q2, 32768, "int8", False)
+        self.assertTrue(any("recommended for this card: 3" in x for x in lines))
+        self.assertTrue(any("at most 8" in x for x in S.parallel_note(12, 48, q2, 32768, "int8", False)))
+        self.assertEqual(len(S.parallel_note(3, 24, q2, 32768, "int8", False)), 1)   # as recommended: no warning
+
+
+if __name__ == "__main__":
+    unittest.main()