Skip to content

serve: tokenise the prompt from the last shared prefix, not from the start - #567

Closed
gputier wants to merge 848 commits into
Niko1221:mainfrom
gputier:pr/prompt-encoder-0.1.37
Closed

gputier wants to merge 848 commits into
Niko1221:mainfrom
gputier:pr/prompt-encoder-0.1.37

Conversation

@gputier

@gputier gputier commented Oct 3, 2026 •

Copy link
Copy Markdown
Contributor

Based on v0.1.39 since 5 October (it was first written on v0.1.37). The rebase takes blange48's port: on 0.1.39 the plain spans of #537 (a </think> quoted in a message is text) also decide which special matches are taken, so the shared prefix also ends where the two prompts' plain spans first differ, and the rest is encoded with its spans shifted by the resume point.

A chat client sends the whole conversation every turn, and Service.prepare runs BPE over all of it each time (serve/server.py, self.tok.encode(prompt, parse_special=True)), even when the engine's conversation cache already holds the whole prefix. /v1/messages/count_tokens does the same.

This PR adds PromptEncoder in tools/strata_tokenizer.py. It keeps the last few rendered prompts with their ids. For a new prompt it finds the last special-token boundary shared with one of them, reuses the ids up to that boundary, and sends only the text after it through BPE.

The ids match a full encode by construction. Encoding splits the text at special-token matches and encodes each stretch between two matches on its own, so nothing crosses a match. Whether a match is taken at position q depends only on text[q : q + max_special_len] and on whether q lies inside a plain span, so the cut sits at least max_special_len - 1 characters before the end of the common prefix. The class docstring has the reasoning.

The tokenizer gains max_special_len and encode_marked() (encode() plus the offset and id count after each special-token match). ByteTokenizer gets the same two, so the mock-engine tests exercise the encoder. 0.1.39's Service.encode_prompt(), which prepare and count_tokens already use, now calls the encoder. encode() itself is untouched.

Diff: 5 files, 3 of them tests.

Measurement

A conversation of 100,893 tokens (433 K characters) that grows by one turn, encoding time for the new prompt, 5 repetitions:

  • full encode: 235 ms median (231 to 241)
  • PromptEncoder: 1.5 ms median (1.05 to 2.07)

Container with 2 cores, Python 3.12, llama.cpp's ggml-vocab-qwen35.gguf (with the template's added tokens), synthetic text mixing English, code and CJK. This is not the model pack's tokenizer and not the target machine, so the absolute numbers will differ there.

Evidence that the ids are identical

On v0.1.39, in a Linux container (Python 3.13, llama.cpp's qwen35 vocabulary): serve.test_prompt_encoder and serve.test_server pass 146 tests, tools/test_strata_tokenizer.py passes 11 with 1 skipped (the pack tokenizer). The rest of this section was measured on the v0.1.37 version.

  • serve/test_prompt_encoder.py (blange48's, added with the rebase): a fuzz against the full encode with random plain spans, spans that change inside the shared part, and edited histories, plus quoted <think> / </think> across turns through Service.encode_prompt.

  • tools/test_strata_tokenizer.py (new, 11 tests): llama.cpp's qwen35 vectors, encode_marked against encode, chat_golden.json, and conversations compared token for token with tok.encode. Those cover several turns with unicode and literal special tokens in the text, tool definitions, calls and responses, images, a changed reasoning effort or system message, an edited earlier turn, a conversation that gets shorter, an empty message, and interleaved conversations, with 1, 2 and 4 kept prompts. Separate tests use byte-level vocabularies whose special tokens overlap each other, which is where a cut without the look-ahead margin breaks.

  • serve/test_server.py, IncrementalPrompts: six turns over HTTP against the mock engine, the same conversations through Service.encode_prompt with both the byte tokenizer and qwen35, and a tokenizer without encode_marked.

  • Mutations: with the margin set to 0, 5 tests fail. With the cut one id too late, 305 fail.

  • Fuzz run outside the test suite: 40,000 cases, no difference in ids. 8,000 on qwen35 (seeds 1000 to 1099) and 32,000 on byte vocabularies with overlapping literals.

  • Full Python suite (tools/test_*.py, serve/test_*.py) in a Linux container: 454 passed, 9 skipped, 0 failed on this branch, against 440 passed, 9 skipped on v0.1.37. That run uses scratch copies of both trees where setup.EXE is forced to strata.exe. Without it, 47 tests fail on Linux, the same 47 on a clean v0.1.37 (test_setup_golden and one test_setup_amd test), because they expect the Windows executable name. The change is not in this PR. The skips are the pack tokenizer, jsonschema and the Windows job objects. No existing test is modified.

Limits

  • No fallback switch. The default path changes, with the same ids, and there is no environment variable to go back to the full encode. I can add one if you want it.
  • Service now imports strata_tokenizer, and so regex, whenever the tokenizer has encode_marked. That includes the byte tokenizer used in tests without a pack. regex is already pinned in requirements.txt.
  • The server keeps up to 4 rendered prompts in memory. Measured at 239,670 tokens (572,670 characters): one entry holds 3.6 MB of ids on top of its text, so the cap of 4 entries is about 15 MB of ids.
  • /v1/messages/count_tokens shares the same encoder, so a count request can evict another client's entry. The cost is speed only, the ids stay the same.
  • Not tested against the real engine or the model pack's vocabulary, only the mock engine and llama.cpp's qwen35 vocabulary.

Niko1221 and others added 30 commits October 2, 2026 09:35
… a failed tuning says why (Niko1221#447)

tools/calibrate.py read any "gpu" list as a layer split and added --layer-split auto, so a one-card config with a
helper card for the expert tier ("gpu": [0] + --expert-cache-device1) was refused by the engine before it read a
weight, a "0,2" string never got its split, and split_skip_if_fits was ignored. The tuner now builds the arguments
with serve.server.engine_args, the same as a normal start.

When the tuning fails, setup prints the engine's own last error line from the log (only lines written during the
tuning), and says the failure again where it is not lost: before --calibrate starts the model, and in the
"All set." summary of an install that tuned.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…etup recommends the big card alone (Niko1221#448)

- engine (layer splits only): when a stage's cache cannot fund the chunk CUDA0 alone would choose, a WARNING
  names it (CUDA index, name, cache slots, the slots the chunk borrows) and points at the helper cache
  (--expert-cache-device1, docs/SECOND_GPU.md).  The chunk choice itself is unchanged: fits() is split into
  fits_one() with the same tests; one GPU never reaches the new log.
- setup: a pair whose second card is under 12 GB beside a first card of 12 GB or more is still offered, but
  the first card alone is recommended (the default with --yes and Enter; offer_together asks "n" by default),
  with a note.  A named --gpus pair is kept.  Equal pairs (2x 8/16/24 GB) are unchanged.
- The automatic drop of the small card from an auto split is not done: stages load their weights before the
  boundary search and the exact prompt-buffer bytes need the sessions made after it.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…the totals (Niko1221#457)

The server already read drafts_offered / drafts_accepted from the engine's DONE line but only returned them in a
response's timings (draft_n / draft_n_accepted). Each /metrics request entry now carries both (null when the
engine did not report them) and the totals keep their running sums since the server started, so a dashboard such
as llm-visuals can show the acceptance. Additive: no existing field changes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…he --family/--model to use (Niko1221#444)

With Unsloth's UD-IQ3_XXS or UD-Q2_K_XL in the folder, setup either said only "has no IQ3_XXS model file"
(--family unsloth --model IQ3_XXS) or took the UD-IQ3_XXS file for the GSQ-RCO IQ3_XXS by the size in its name.
Now:
- a chosen shard whose name is a quantization setup has no size for (UD-IQ3_XXS, UD-Q2_K_XL, a K-quant) stops
  with "<file> is UD-IQ3_XXS, a GGUF Strata cannot run" and what Strata runs: ISTA-DASLab's GSQ-RCO files (and
  Swift 1.5's and the Coder's) and Unsloth's UD-Q4_K_XL only;
- a folder without the chosen file but with other GGUFs lists the ones Strata cannot use and the --family/--model
  of the ones it can (an empty folder keeps check_shards' "missing");
- a size the family does not have names the families that do (--family qwen --model IQ3_XXS), and with
  --gguf-dir also what Strata runs.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e unpruned sizes for that (Niko1221#438)

README and docs/MODELS.md said only "Weaker outside coding". Niko1221#438 measured Chinese answers wrong or looping on the
Coder where English was fine; both now say plainly that this includes Chinese and other CJK text and point to the
sizes that keep every expert (Q2_0, IQ2_XS, IQ3_S).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#461), the low-RAM resident mode on Windows 32 GB PCs (Niko1221#467), fixes Niko1221#460 Niko1221#459 Niko1221#446 Niko1221#447 Niko1221#457 Niko1221#448 Niko1221#444; setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The fill from the GGUF (--stream-experts) read each expert's three slices,
copied the blob and waited for the copy before the next read: ~12,000-18,500
experts with neither the SSD nor PCIe kept busy. Now the slots are admitted
in profile order first (same placement), reads go in file order by up to 8
threads into page-locked batches of 64, and a batch's copies overlap the next
batch's reads. STRATA_FILL_SERIAL=1 keeps the serial fill.

Tested on the B70: cold page cache, Coder fill 23.4 GiB 76.2 s -> 18.8 s;
launch to first token Coder 82 s -> 26 s, IQ2_XS 120 s -> 41 s; outputs
identical to the serial fill.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… read (Niko1221#471)

The engine's `strata serve: prompt ...` line now says "= 0 reused + 12288 of 98179 read in ... (cancelled)" with the
rate from the tokens read, and DONE gets a 15th field: the prompt tokens actually read. The server's /metrics history
and totals count reused + read for a cancelled read (prompt_total keeps the whole prompt); an older engine's DONE line
(no such field) is read as before.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Engine: when the MTP draft head's subset does not fit, it prints what the head needs, the VRAM free, and the smaller
shipped subsets (--draft-vocab cyrillic / en, their MiB from the head's real row size) with how to pick one
(START-HERE.bat / ./setup.sh --draft-vocab en, kept in the model's config). Text only, on the failure path.
Server: a start that ended that way carries the engine's hint (or the same advice for an older engine) in its
"the engine exited before it was ready" error. Setup: a card under 14 GB with the default subset gets a note
recommending --draft-vocab en; nothing is changed, an explicit or saved choice gets no note, and the golden
configs are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…o1221#475)

In a git clone they run `git pull --ff-only` (a zip copy is told to download the new zip), then `setup.py --update`:
what a plain start does to an install before it starts the model - the Python packages, the ready-made engine when
MIN_ENGINE is newer (a compiled one when its source changed), each model's config upgrades and its draft subset, the
Windows AMD runtime - and nothing else: no question, no model files, no start. Both scripts run from one parsed
block, since the pull can rewrite the running script. Documented in docs/INSTALL.md (Updating) and the README.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…profile-save, Niko1221#477)

`--expert-profile-save P` (--serve) writes the adaptive tier's learning as a profile in tools/make_profile.py's
format, byte for byte (read_expert_profile reads it back): every pair ranked - the experts resident in VRAM now
first, then the routing counted since the start (the decayed usage summed at each adapt step, before its decay),
then the profile the engine started from. On QUIT / stdin EOF and every --expert-profile-save-every minutes
(default 10, 0 = exit only) between requests, through <P>.tmp and a rename. Without the flag nothing is counted or
written: `heat` stays empty and every added loop is over it.

Server: "expert_profile_save" (and "expert_profile_save_every") in the config pass the flags, and the next start
begins from the saved file instead of the config's --expert-profile when its header matches that profile's
layers x experts. tests/core/expert_profile_save_test.cpp (CPU only) checks the order and the bytes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…, the draft-head hint (Niko1221#474), UPDATE.bat (Niko1221#475), --expert-profile-save (Niko1221#477); setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)

Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32
values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync),
then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in
registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue)
and one down launch into the per-slot rows the existing combine reads.  Experts are read where they are (cache
slot or streamed ring slot, a pointer per expert): no gather.

Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with
STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU).  Everything else - the default, native/IQ packs, small
chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch).

tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS
1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5.  One layer at 2048 tokens x top 10
over 512 experts: MMQ path 11.12 ms, fused 3.13 ms.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e layer's experts) by default

The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch
waited for slots its own first batch frees.  With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM
ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides).

RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 ->
2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%.  1024 slots: 4K +24-27%, but the 32K chunk falls
from 8192 to 6144 (+6%).  The default path's output is identical to 0.1.35's on both prompts.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… hold only what the fused kernels need

In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the
grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H
and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the
GDN set does).  A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for
stream_all_min() - 1 tokens.  The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default
stays 512).  Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events.

The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert
of the layer, routed or not - at a streamed chunk's size nearly all are.

RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots):
  4K:   MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046)
  32K:  MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153)
  128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346
        (8192), + 1024-slot ring 2,392 (6144)
A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per
8K tokens.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ault (two layers' experts)

With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots,
under the 90% cap).  RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating
pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to
6144 and is still faster.  Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656
(+21.5%), 128K 2,128 -> 2,454 (+15.3%).  The ring does not change the numbers (the same output as at 512 slots).
Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…y argmax, the same ids and tokens

- qsa_block_topk: calls of up to 16 queries (decode windows) run on a cluster of 8 CTAs per query
  (block_topk_cluster_kernel): keys in shared memory, digit histograms pushed through distributed shared memory,
  one cluster barrier per radix pass, a parallel digit search, cluster-wide offsets for the ordered emit. The
  threshold and the cut at it are the one-CTA kernels' (same tie rule, NaN keys 0). STRATA_QSA_CLUSTER=0: off.
  RTX 5070 per call: 21.9 -> 15.6 us at 32K, 58 -> 18 at 128K, 200 -> 22 at 262K.
- sample_tokens greedy without penalties: a cluster of 8 CTAs per row (sampler_greedy_cluster_kernel), four loads
  in flight per thread, results merged in CTA 0 (lowest index on ties, NaN never taken, 0 when none).
  STRATA_ARGMAX_MULTI=0: off. 39.6 -> 5.9 us per call at 248,320 logits.
- Both launch through cudaLaunchKernelEx with the cluster attribute (captured in the decode graphs) only where
  they run: CUDA, compute capability 9+ (STRATA_EMULATE_CC respected), the kernel's own PTX/SASS at sm_90+ (an
  older build's JIT copy is a trap), room for one cluster - checked once per device. HIP and sm_75/86/89 keep the
  one-CTA kernels.
- decode_cluster_parity (ctest; skips below sm_90): bitwise vs the one-CTA dispatch and the reference kernel over
  434 top-k cases (contexts 7..262,144, every tail weight, 1-16 queries, ties / NaN / +-inf / -0 / denormal / all
  NaN) and 216 argmax cases, plus a CUDA graph capturing both (first calls inside the capture); --bench timings.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)

src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert -
cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of
the native packs.  The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign
tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit
loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which
multiply at m16n8k16).  The activations are int8 per 32 values in natural order (quantize_act_native); fragments by
ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add.  A work item is 256
weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128).

Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the
FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers.  sm_75, HIP and the
default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy
tokens, 4K and 32K).

tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision
reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 /
2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5),
IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 /
2.72e-2 (4.4e-5); both tile shapes.  One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192
tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6,
12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5.  MMQ's products alone (no gathers) at
3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them.
(The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be
queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0.
prefill_fused_moe_test's timing part has the same pattern.)

End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt
tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K
1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%).  The IQ3
packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s
of 3.4 s), so faster kernels change little.  Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at
162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89,
IQ3_S 31 / 256.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…(STRATA_PF_FUSED=0 off), the native IQ packs' fused kernels opt-in (=1), cluster decode kernels; the fused test waits for its uploads

The ring rule: 1024 slots for the Q2_0 pack's fused experts (P3), 512 for native packs with fused layers.
Quality (teacher-forced vs the FP16 path, Q2_0): fused as close as MMQ at 8K, closer at 32K (phaseA-tests s18-tf).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
device, expert_source, native_head, prefill gemm, prefill.cpp, generate.cpp
merged from dpct of aeb35be and d9ab843 with sycl/tools/merge_upstream.py (the
3-way merge restricted to files upstream changed, dpct kernel-name hashes
canonicalized to the port's). Version 0.1.35.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ls) and their switches

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The PCIe probe sets the share of the missed experts the GPU reads over
the link (--pcie-frac): the native default 0.55 from 20 GB/s up, less
on a slower link (#44). On Windows HIP it timed its copies with events
that do not bracket them there: an RX 6800 read between 3,348 and
25,999 GB/s in 24 engine starts, so every link kept the x16 share. A
standalone copy of the probe on that card: its events measure 0.13 ms
for the 1 GiB, the host clock 43 ms (25.1 GB/s).

On Windows HIP the probe now times the same copies on the host, between
two synchronizes (1 GiB, tens of milliseconds): the same RX 6800 on
PCIe 4.0 x16 reads 27.2 GB/s in every run, which keeps 0.55 there. On
a PCIe 4.0 x8 or 3.0 x16 link the share now shrinks as #44 intended.
CUDA and Linux keep the events.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…DGET)

On Windows, hipMemGetInfo (ROCclr's PAL backend) reports the card's size
minus this process's own allocations; what the desktop and other
programs hold on the card is not subtracted. On a card that also drives
the desktop the engine counts the ~930 MiB those hold as free, and
--expert-cache auto fills the card past what Windows lets the process
keep in VRAM: WDDM moves part of it out to system RAM. On an RX 6800
(HIP SDK 7.2) with setup's configuration: a 9.83 GiB cache, the card at
16132 of its 16368 MiB, 929 MiB more of the engine in system RAM than
with this change, 30.1-31.0 tok/s decode. On a 0.1.30 build,
--vram-reserve-mib 1500 left the card as full (16037 MiB), since the
reserve is taken from the same figure.

Windows gives every process a video memory budget that does account
for the others (DXGI QueryVideoMemoryInfo). On that card it is 799 MiB
below the card's size while the process is small and about 1.8 GiB
below once it holds 11 GiB; allocations past it still succeed.
cudaMemGetInfo, which the HIP shim maps to hipMemGetInfo, now lowers
hipMemGetInfo's free figure by what the budget withholds (the card's
size minus the budget), on Windows only. What is sized from it then
stays within the budget: --expert-cache auto sizes the cache from it,
and the existing check after the slots are written shrinks it ("only
8 MiB free once the slots are written ... shrinking"): an 8.05 GiB
cache, the card at 15230 MiB, 40.6-41.5 tok/s.
Lowering HIP's own figure, instead of taking the budget minus DXGI's
usage, keeps memory the HIP runtime has freed but still caches counted
as free, as HIP counts it.

dxgi.dll is loaded at run time, so nothing new is linked; the adapter
is found by the device's LUID. Without dxgi.dll, the adapter or a
budget below the card's size, hipMemGetInfo's figure stands. The first
query logs the budget once. STRATA_WDDM_BUDGET=0 turns it off. CUDA
and Linux builds are unchanged.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…g forever (Niko1221#481)

A request whose engine prints nothing for engine_silence_s (default 300 s; while a prompt is read, a chunk may take
3x the previous one's time, the first up to its tokens at 50 tok/s more) ends with EngineSilent (an EngineDied): the
engine is killed and the next request starts it again. The STOP drain, under the request FIFO, is no longer an
untimed wait. "engine_silence_s": 0 keeps the old behaviour. Tests: unit (silence, PP chunks, first chunk, STOP
never acknowledged, 0, config) and over HTTP with a fake engine process (non-stream, stream, stop token).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…0 log lines (Niko1221#496)

Every start failure, not only the draft-head case (Niko1221#474), now carries the engine's own reason: the last 20 lines
this start wrote to the engine log (never an earlier start's). The failed start also closes its pipes and its
handle on the log. Tests: the tail itself, and a fake engine that exits with a reason in its log.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221 and others added 4 commits October 4, 2026 10:14
# Conflicts:
#	src/core/verify.cpp
#	src/program/generate.cpp
…l); UD-Q4_K_XL stays experimental

- "experimental" is per model now: MODELS["UD-Q4_K_XL"] keeps it, the unsloth family and UD-IQ4_XS drop it. The
  first menu lists the family without [experimental], with a one-line about (UD-IQ4_XS: a 94 GB download; under
  ~80 GB of RAM part of its experts are read from the SSD). Its size menu lists UD-IQ4_XS first (the default) and
  UD-Q4_K_XL second; an experimental size always sorts last. --model names are unchanged.
- UD-IQ4_XS: no EXPERIMENTAL warning, --check verdict without it. Images are asked (off by default), per model: the
  engine's image path has no restriction for this pack, which uses the original model's image encoder; not yet run
  with images. UD-Q4_K_XL keeps images off.
- MCP: the fallback tables mirror setup's (per-model experimental/vision), sizes in setup's order, per-size
  "experimental" and "images", the vision check per model.
- Docs: README and its six translations, MODELS, UNSLOTH_Q4, AI_SETUP, MCP_SERVER.
- Tests: the menus, the default, --model by name, images, --check, the pack/tokenizer paths; the golden configs
  are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
…e-mib 1000 if a request stalls

With the image encoder on the GPU, the 700 MiB reserve can leave ~200 MiB free on a 12 GB card (UD-IQ4_XS
end-to-end check on an RTX 5070). Setup prints a one-line recommendation; the written config is unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
@Niko1221

Niko1221 commented Oct 4, 2026

Copy link
Copy Markdown
Owner

Thank you, and thanks @sergqwer for the test with the pack's own tokenizer. This is independent of the work in 0.1.39, so it goes to a later release. We will review it then together with the server's other prompt-cache changes.

zuiyi233 added a commit to zuiyi233/Strata that referenced this pull request Oct 4, 2026
@blange48

blange48 commented Oct 4, 2026

Copy link
Copy Markdown
Contributor

In case it helps the review: 0.1.39's plain spans (#537, a </think> quoted in a message is text) also decide which special matches are taken, so the resume point has to respect them. I rebased this on 0.1.39 with that one addition: the shared prefix ends where the two prompts' plain spans first differ, and the tail is encoded with its spans shifted by the resume point. Branch: https://github.com/blange48/Strata/tree/port-567 (commit d07043e, your authorship kept as co-author).

Checks:

  • serve/test_prompt_encoder.py fuzzes it against the full encode: random plain spans, spans that change inside the shared part, and edited histories. It also covers quoted <think> / </think> across turns through Service.encode_prompt.
  • With the Qwen tokenizer, a 12-turn agent conversation up to 145K tokens gives identical ids on every turn. Encoding takes ~83 ms per turn instead of 485 ms at the end.

blange48 added a commit to blange48/Strata that referenced this pull request Oct 4, 2026
…ecent one (Niko1221#567 on 0.1.39)

Niko1221#567 (by gputier) reuses the ids of a recent prompt up to a special-token boundary both share and
encodes only the rest: the same ids as a full encode, because nothing crosses a special match. 0.1.39's
plain spans (Niko1221#537: a </think> quoted in a message is text) also decide which matches are taken, so the
shared prefix now ends where the two prompts' plain spans first differ, and the rest is encoded with its
spans shifted. With the Qwen tokenizer, an agent conversation of 145K tokens: 485 ms -> 83 ms per turn,
ids identical on every turn. serve/test_prompt_encoder.py: a fuzz against the full encode with random
plain spans (changing in the shared part too, and edited histories), a turn's reuse, and quoted
<think> / </think> across turns through Service.encode_prompt.

Co-Authored-By: gputier <gputier@users.noreply.github.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…the start

Every request re-encoded the whole conversation: 235 ms for a 100K-token
prompt on a two-core container, each turn. PromptEncoder keeps the last
prompts with their ids and resume points, and encodes only what follows
the longest shared prefix: 1.5 ms for the same turn.

On 0.1.39 the plain spans of Niko1221#537 (a </think> quoted in a message is
text) also decide which special matches are taken. The shared prefix now
ends where the two prompts' plain spans first differ, and the rest is
encoded with its spans shifted by the resume point (blange48's port).

The ids are the ones tok.encode gives, token for token. The tests compare
both paths on growing, edited and shortened conversations, with tools,
images and special-token literals in the text, on the byte tokenizer and
on the qwen35 vocabulary, and fuzz random plain spans against the full
encode. A tokenizer without resume points keeps the full encode.

Co-Authored-By: benoit lange <benoit@lange.xyz>
@gputier
gputier force-pushed the pr/prompt-encoder-0.1.37 branch from f29e527 to bfe4bcf Compare October 5, 2026 15:20
@gputier

gputier commented Oct 5, 2026

Copy link
Copy Markdown
Contributor Author

Thanks @sergqwer for the run with the pack's own tokenizer, and @blange48 for the port. You're right about #537: on 0.1.39 the resume point has to respect the plain spans, and my branch didn't.

The PR branch is now on v0.1.39 (bfe4bcf), with your change and your fuzz test (serve/test_prompt_encoder.py), and you as co-author. The original tests stay: tools/test_strata_tokenizer.py and IncrementalPrompts in serve/test_server.py, adapted to 0.1.39's encode_prompt(messages, tools, kwargs).

On v0.1.39, Linux container, Python 3.13, llama.cpp's qwen35 vocabulary: serve.test_prompt_encoder and serve.test_server pass 146 tests, and tools/test_strata_tokenizer.py passes 11 with 1 skipped (the pack tokenizer, which I don't have).

@Niko1221 nothing more to do on my side until your prompt-cache review.

@blange48

blange48 commented Oct 5, 2026

Copy link
Copy Markdown
Contributor

Thanks @gputier, great that it's on 0.1.39 with both test sets.

One heads-up for the review: #931 (control-token text inside a message stays text) changes the same spot in encode_prompt. It replaces the <think>-only plain spans with spans for every control token (mark_literals / unmark_literals), so the two PRs conflict textually. They do combine cleanly, because the encoder already takes arbitrary plain spans. The merge just passes #931's spans to PromptEncoder.encode:

plain = ()
if changed:
    prompt, plain = unmark_literals(prompt, self.literals)
if self.prompts is not None:
    return self.prompts.encode(prompt, plain)
return self.tok.encode(prompt, parse_special=True, plain=plain) if plain else \
    self.tok.encode(prompt, parse_special=True)

We run exactly that in production since this afternoon: 0.1.39 with both PRs on a 4-GPU split, agents (OpenCode, Claude Code) and chat. The server tests pass (193, including #931's and the fuzz test). Follow-up turns of a long conversation now reach their first token in 0.22-0.24 s.

…ncode_prompt

test_turns compared the ids with tok.encode of the rendered prompt, which
reads <|im_end|> in a message as a control token. With Niko1221#931 that text
stays text, so the reference is now encode_prompt on a Service without
the encoder: the same check, whatever encode_prompt does with literals.
The tokenizer without resume points forwards everything else to
ByteTokenizer, so it also has Niko1221#931's control_tokens.

Both pass on this branch (146 tests) and merged with Niko1221#931 (152 tests).
@gputier

gputier commented Oct 5, 2026

Copy link
Copy Markdown
Contributor Author

Thanks @blange48, I merged #931 into this branch locally with your resolution and it holds for the code. Two of this PR's own tests did not, though: IncrementalPrompts.test_turns and test_a_tokenizer_without_resume_points_encodes_in_full in serve/test_server.py. Your port (d07043e) doesn't carry them, only test_prompt_encoder.py, so your 193 wouldn't have run them.

Pushed d55679c: test_turns now takes its reference from encode_prompt on a Service without the encoder, and the stub forwards everything but encode_marked / max_special_len to ByteTokenizer. serve.test_server and serve.test_prompt_encoder in a Linux container, Python 3.13: 146 tests on this branch alone, 152 with #931 merged, all passing. So the two PRs can go in either order; whichever is second only needs your encode_prompt resolution.

@sergqwer

sergqwer commented Oct 5, 2026

Copy link
Copy Markdown

Ran d55679c on Windows 11 (Python 3.13) with the tokenizer you were missing: the pack's own (Qwen3.8-Flash-Next, 248,320 tokens, STRATA_TOKENIZER=<pack>/tokenizer).

  • tools/test_strata_tokenizer.py: 10 of 11 pass with the pack tokenizer. The one failure is the test, not the encoder. test_llama_cpp_vectors gets its tokenizer from load_tokenizer(), which prefers $STRATA_TOKENIZER, and then compares it with llama.cpp's qwen35 vectors. With both set, it encodes those vectors with another vocabulary. The first one differs at once: [1086, 220, 19, 220, 25229, 3818] against [1122, 220, 19, 220, 26062, 3951]. With only STRATA_GGUF_PY all 11 pass. The oracle could build its tokenizer from vocab_gguf() directly.
  • The serve suite with the pack tokenizer: 274 of 275. The error is test_responses.OverHttp.test_json_schema_text_format, which 0.1.39 has on its own wherever jsonschema is not installed (serve tests: test_json_schema_text_format's schema failure only where jsonschema is installed #785).

test_llama_cpp_vectors took its tokenizer from load_tokenizer(), which
prefers $STRATA_TOKENIZER: with a pack set, it encoded qwen35's vectors
with the pack's vocabulary and failed on the first one (sergqwer's run).
It now builds the tokenizer from vocab_gguf() itself. With a pack whose
merges differ, the test fails before the change and passes after.
@gputier

gputier commented Oct 5, 2026

Copy link
Copy Markdown
Contributor Author

Thanks @sergqwer, you're right: the oracle took load_tokenizer(), so with STRATA_TOKENIZER set it checked qwen35's vectors against the pack's vocabulary. Fixed in 76b916e: test_llama_cpp_vectors builds its tokenizer from vocab_gguf() (a new gguf_tokenizer(path), which load_tokenizer() also uses), and the other tests still prefer the pack.

Checked with a stand-in pack, since I don't have the real one: the qwen35 vocabulary with the template's added tokens and half its merges. With STRATA_TOKENIZER pointing at it, test_llama_cpp_vectors fails before the change and passes after; the whole file passes 11 of 11 with and without the pack, the vectors included.

@Niko1221

Niko1221 commented Oct 6, 2026

Copy link
Copy Markdown
Owner

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

fragtion added a commit to fragtion/StrataPool that referenced this pull request Oct 6, 2026
Niko1221#567, gputier's port)

From gputier/StrataGP pr/prompt-encoder-0.1.37 (bfe4bcf, d55679c, 76b916e), merged onto 0.1.40:
PromptEncoder keeps the last prompts with their ids and resume points and encodes only what follows
the longest shared prefix.  Kept with this tree's Niko1221#931 (a control token's text in a message is text):
the plain spans go to PromptEncoder as they went to the full encode.
Checked here on llama.cpp's ggml-vocab-qwen35.gguf: the tokenizer tests (llama.cpp's vectors, chat
goldens, growing/edited/shortened conversations) pass, and a 60-turn conversation encodes identically
both ways; at 110K tokens a turn took 601 ms with the full encode and 12.7 ms with PromptEncoder
(this container's CPU).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0157fjixQD2nRxvCvtXq66DD
@gputier
gputier deleted the pr/prompt-encoder-0.1.37 branch October 10, 2026 21:07
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 1138 commit(s) 61b3fb5..ee95207 on https://github.com/gputier/StrataGP, restricted to the
paths this fork still mirrors: 2 file(s) under tools -- 2 clean,
0 with conflict markers.

Original authors: Guillaume PUTIER

Original commits touching these paths (5 subject(s)):
  - serve, tools: Niko1221#567's tests and ByteTokenizer without repeated code
  - serve: Niko1221#567's kept prompts, one per conversation, none refused for the context
  - serve: count tokens without storing the prompt in the encoder
  - serve: the prompt is tokenised from the last shared prefix, not from the start
  - tools: llama.cpp's qwen35 vectors are checked with the qwen35 vocabulary

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/gputier/StrataGP
Fork-Replay-Branch: main
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: ee95207
Fork-Replay-Group: tools
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 1138 commit(s) 61b3fb5..ee95207 on https://github.com/gputier/StrataGP, restricted to the
paths this fork still mirrors: 2 file(s) under tools -- 2 clean,
0 with conflict markers.

Original authors: Guillaume PUTIER

Original commits touching these paths (5 subject(s)):
  - serve, tools: Niko1221#567's tests and ByteTokenizer without repeated code
  - serve: Niko1221#567's kept prompts, one per conversation, none refused for the context
  - serve: count tokens without storing the prompt in the encoder
  - serve: the prompt is tokenised from the last shared prefix, not from the start
  - tools: llama.cpp's qwen35 vectors are checked with the qwen35 vocabulary

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/gputier/StrataGP
Fork-Replay-Branch: main
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: ee95207
Fork-Replay-Group: tools
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 1138 commit(s) 61b3fb5..ee95207 on https://github.com/gputier/StrataGP, restricted to the
paths this fork still mirrors: 2 file(s) under tools -- 2 clean,
0 with conflict markers.

Original authors: Guillaume PUTIER

Original commits touching these paths (5 subject(s)):
  - serve, tools: Niko1221#567's tests and ByteTokenizer without repeated code
  - serve: Niko1221#567's kept prompts, one per conversation, none refused for the context
  - serve: count tokens without storing the prompt in the encoder
  - serve: the prompt is tokenised from the last shared prefix, not from the start
  - tools: llama.cpp's qwen35 vectors are checked with the qwen35 vocabulary

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/gputier/StrataGP
Fork-Replay-Branch: main
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: ee95207
Fork-Replay-Group: tools
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 1138 commit(s) 61b3fb5..ee95207 on https://github.com/gputier/StrataGP, restricted to the
paths this fork still mirrors: 2 file(s) under tools -- 2 clean,
0 with conflict markers.

Original authors: Guillaume PUTIER

Original commits touching these paths (5 subject(s)):
  - serve, tools: Niko1221#567's tests and ByteTokenizer without repeated code
  - serve: Niko1221#567's kept prompts, one per conversation, none refused for the context
  - serve: count tokens without storing the prompt in the encoder
  - serve: the prompt is tokenised from the last shared prefix, not from the start
  - tools: llama.cpp's qwen35 vectors are checked with the qwen35 vocabulary

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/gputier/StrataGP
Fork-Replay-Branch: main
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: ee95207
Fork-Replay-Group: tools
Fork-Replay-Automated: true
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.