[justindressler]

Problem and result

Problem. Strata re-reads the whole prompt from scratch on every request, and again after every
splice or mid-history edit. On a conversation with a long shared prefix that is repeated work: a
16-turn agent session re-prefilled the identical leading context sixteen times, and one mid-history
edit threw away the entire cache and re-read the whole 9.4k-token prefix.

Result. Three independent, default-off flags let the engine mount a parked conversation's KV
state when the incoming prompt shares a verified prefix, and bound the re-read after an edit. With
the flags off, answers are unchanged; with them on, prefill on cache hits drops from ~20 s to ~10 s
and exact replays from ~20 s to 0.2 s. Measured on the RTX 3090 that this repo's test rig uses, on a
164-row ledger (9.2k-9.6k prompt tokens): a 16-turn edit chain went from 234.2 s to 162.4 s
(-31%) with identical answers turn for turn, and an exact replay costs 0.2 s at 99.9% reuse.

This revision also rebases the branch from v0.1.45 (477ac13c) onto current main
(9de04cd4, v0.2.1.4), which is what the earlier revision could not be reviewed against.

Related discussion

• Tail-fork probe requested by @MrKop-git in this thread (a coding agent's sibling sub-agent shape:
  two requests sharing the whole prefix, differing only in the final instruction). Ran on four flag
  arms; the B-row reuse and wall times are in that comment.
• Prior close: this PR was closed in the 2026-10-19 tracker reset for 0.2.0-beta, not on content.

Approach

Three orthogonal layers, each usable alone:

1. --cache-suffix-reuse <N> — mount a parked previous conversation whose token prefix matches
   the incoming prompt (LCP-verified), narrowing the read window. Works across all five KV formats
   (fp16, int8, int8+KV_ROT, k8v4, q4_0), splicing quantized formats through their own quant/dequant
   paths so reused blocks stay byte-identical.
2. --prompt-cache-dense-tail <N> (FW-1) — dense checkpointing of the tail region so late-turn
   conversations restore cheaply.
3. --prompt-cache-state-every <N> (FW-2b) — state checkpoint ladder at a token cadence,
   capping the re-read gap after mid-history edits and compaction-style rewrites.

Correctness posture: mounts only from LCP-verified prefixes; the recurrence-state restore refuses any
rung above the divergence point (fail-closed), so a stale-state probe from development never reaches
main. Compatibility: additive flags, defaults off, no change to the wire format, model format or
existing cache behaviour. Design and invariants: docs/SUFFIX_REUSE.md in this PR.

What changed in the rebase, for review

The diff against main is 23 files, +6,272 / -19. The rebase itself touches one file for real (a
five-line CMakeLists.txt guard) plus the 19 rewritten upstream lines in generate.cpp; the rest is
the feature exactly as it was before, which is what makes this gateable in one sitting.

• One conflict, not many: CMakeLists.txt. Upstream added a Win32 HIP delay-load block
  (delayimp, /DELAYLOAD:hipblas.dll, libhipblaslt.dll, rocblas.dll); this branch adds a
  compile definition. Both kept, upstream's first. Our STRATA_HAVE_CUDA_BUILD=1 define is scoped
  to the GPU build so a CPU translation unit never reaches the fail-closed #error in
  kv_splice_transport.hpp:92. The wrapper sits inside the outer GPU block and is therefore
  redundant in that position — kept deliberately as defence in depth.
• generate.cpp auto-merged, so a clean merge proves nothing by itself. All 19 upstream lines
  this branch rewrites (!gpu_sync_ok() -> sync_ok, checkpoint_save -> save_ok, short-circuit
  and fail-closed return false preserved) were audited line by line, and the whole merged
  translation unit was compiled in both trees.
• One deliberate delta to state plainly: lap() now accumulates unconditionally, where upstream
  returns early unless ckpt_trace. That is this branch's own change (the FW-2a probe reads those
  timings without STRATA_TRACE), not a rebase artifact: two extra Clock::now() calls per
  checkpoint boundary. "Byte-identical with the flags off" therefore means identical output and KV
  behaviour, with that one timing accumulator added.

Validation

Revision: branch rebased onto main at 9de04cd4 (v0.2.1.4); new head cf6c6984, 36 commits
replayed, nothing squashed. Environment: WSL2, 32 cores, RTX 3090 24 GB, CUDA 13.0 (sm_86).
Model: Qwen3.8-Flash-Next abliterated GSQ-RCO IQ3_XXS, int8 KV, MTP spec-4, temperature 0.

Passed:

• Host ctest (CPU build, -j2): 49/54 passed, 1 skipped, 4 failed, 21.5 s. The four failures
  (native_hadamard_test, tq_cpu_parity, tq_cpu_parity_avx2, isa_guard_check) fail
  identically on an unmodified main worktree built the same way — this machine's fetched
  llama.cpp lacks the trellis types and has a pre-existing AVX-512 invariant violation. None of
  those four touch a file in this PR.

• Partial-prefix suite: 82/82 (the PR-2351 baseline), 0 failed.

• Off-parity check: PASS — with the flags off, engine output is unchanged against base.

• CPU build: exit 0, 0 warnings, 57 targets, 53 executables, 65 s.

• generate.cpp merged-TU compile via the repo's own metal_compat shim: 0 errors, 0 warnings
  in both trees (16,611 lines this branch vs 16,573 upstream).

• CUDA build (-DCMAKE_CUDA_ARCHITECTURES=86, CUDA 13.0): exit 0, strata binary 83,388,904 B.
  Built single-threaded; two concurrent nvcc front-ends crash on this host, which is a machine
  constraint, not a source finding. 21 nvcc warnings, all in files this PR does not touch.

• Needle-in-a-haystack on the rebased engine: 5 of 5 depths found (10/30/50/70/90%, 23,470
  prompt tokens).

• 16-turn edit chain (cold, appends, inserts at rows 40/120/5, delete at 60, four exact replays;
  each turn asks for the amount on a named row id, an amount that occurs nowhere else in the prompt):

  flags	correct	wall, 16 turns	reuse on append / mid-insert	reuse on exact replay	cold read
  --cache-suffix-reuse 64 only	14/16	234.2 s	0% (19.9 s per turn)	99.9% (0.2 s)	14.8 s
  + --prompt-cache-dense-tail 4096 --prompt-cache-state-every 256	14/16	162.4 s (-31%)	55-57% (10.4 s per turn)	99.9% (0.2 s)	20.0 s

• Control: stock upstream 9de04cd4 with none of this branch's commits ran the identical chain
  and produced the same 16 answers, turn for turn, 14/16 — zero disagreements against either arm
  above. The two failures are the model's, not the cache's: turn 5 returns a neighbouring row's
  amount, and turn 6 reproduces that answer from cache. A single cold request for turn 5's prompt,
  in a fresh engine process, returns the same wrong answer — so it is the model reading a 9.4k-row
  ledger, independent of any caching.

• Deep edit (insert at row 5, below every checkpoint) answered correctly in all three arms via the
  fail-closed restore-and-re-read path, at full read cost. The re-read is bounded by the deepest
  anchor: designed behaviour, not a miss.

• Independent review of this revision, scoped to the rebase delta rather than a re-read of the feature:
  APPROVE-WITH-FINDINGS, 0 blocker, 0 major, 4 minor from two reviewers. The two findings they
  shared are addressed here: the evidence snippet for the import/guard blocks is corrected, and the
  redundant CMakeLists.txt guard is documented above instead of left to confuse a reader.

Failed:

• None observed on this revision.

Not run:

• HIP and SYCL builds. Only CUDA was built and executed. The shared code paths are covered by the
  host tests above, but those backends are unverified on this revision. Per AGENTS.md this PR does
  touch backend-shared files, so this is a stated gap, not an oversight.
• Metal, multi-GPU, and any architecture other than sm_86: no hardware here.
• A determinism study. At temperature 0 the full-prefill path is not reliably reproducible run to run
  on knife-edge questions (observed 162 then 161 on identical prompts, same binary and config, with
  spec-4 decoding and an 84-99% expert-cache hit). That is why the chain scores an unambiguous row-id
  lookup, and why no claim above rests on a single run's answer.

Limits

• Reuse is bounded by what the parked pool holds and by the deepest checkpoint below the edit; a miss
  falls back to a full read. Edits deeper than every anchor still pay full prefill cost.
• The win is workload-shaped: near-full on append and replay, partial on mid-history edits.
• The dense-tail/state-cadence stack costs +5.2 s on the cold read (14.8 s -> 20.0 s). That is the
  trade, and it is why the flags are opt-in rather than default.
• FW-2c (session-local partial K/V splice, state-advance kernel) is deliberately not in this PR —
  parked with a plan in docs/SUFFIX_REUSE.md §10.
• Disk persistence of the conversation pool is a separate effort.
• Numbers are single-machine (one RTX 3090, WSL2); no multi-GPU or other-architecture claims.

[MrKop-git]

Tail-fork probe for the sibling sub-agent shape

One request shape worth measuring with this PR's flags is what a coding agent's sibling sub-agent (or a retry) hits: two requests share the entire leading context and differ only in the final instruction. The first one pays the cold prefill; the second one forks at the tail.

tools/tail_fork_probe.py (script below) sends exactly that pair — request A cold, then request B with A's body byte-identical and only the last instruction replaced — and prints the engine's own prompt_n / cached_tokens (reused) / read plus wall time for both arms. The B row is the one to quote: its reused is how much of the shared prefix survived, its wall is what the sibling actually waits for.

python3 tools/tail_fork_probe.py --words 5000 --url http://127.0.0.1:8080

Determinism: fixed synthetic body, temperature 0, thinking off, standard library only — the two arms are byte-identical up to the tail by construction.

How it maps to the three flags

Flag	What the probe's B row shows
--cache-suffix-reuse <N>	A's parked conversation mounts on B by the LCP (the shared body); reused should cover it and read should shrink to the tail
--prompt-cache-dense-tail <N>	how far back B re-reads once the tail has diverged — checkpoint density around the fork point
--prompt-cache-state-every <N>	the re-read gap after a fork/rewrite, at a token cadence

Suggested arms, same box and session, base first: flags off / --cache-suffix-reuse alone / the checkpoint ladder alone / all three. Quote B's reused, read, wall and the B reuse: …% line the script prints, and note the link tier the boxes in the PR's tables use, so the B rows stay comparable.

One caveat the script handles

prompt_n is not self-describing: on a checkpoint hit some builds report the tokens computed that call rather than the prompt's full length. The script dumps raw usage/timings for both arms, so a small prompt_n can be told apart from a truncated prompt instead of guessed.

tools/tail_fork_probe.py

#!/usr/bin/env python3
"""Tail-fork (sibling sub-agent) probe: same leading context, replaced tail.

Sends request A cold, then request B that shares A's body and replaces only the final
instruction - the shape a coding agent's sibling sub-agent or a retry hits. Prints the
engine's own prompt_n / reused / read and the wall time for both, so the reuse can be
compared across PCIe link tiers. Standard library only.

    python3 tools/tail_fork_probe.py --words 5000 [--url http://127.0.0.1:8080]

The B row is the one to quote: its `reused` is how much of the shared prefix survived,
its `wall` is what the sibling actually waits for.
"""
from __future__ import annotations

import argparse
import json
import time
import urllib.request

VOCAB = ("river stone market lantern copper winter garden signal harbor needle orbit velvet canyon ember "
         "ledger quartz meadow anchor violet thunder").split()

TAIL_A = "\n\nSummarize what kinds of words appear in the list, in five sentences."
TAIL_B = "\n\nList the ten most frequent words and say what each of them means."


def post(url: str, prompt: str, max_tokens: int) -> tuple[dict, float]:
    req = urllib.request.Request(
        f"{url}/v1/chat/completions",
        data=json.dumps({"messages": [{"role": "user", "content": prompt}], "max_tokens": max_tokens,
                         "temperature": 0, "chat_template_kwargs": {"enable_thinking": False}}).encode(),
        headers={"Content-Type": "application/json"})
    t0 = time.time()
    with urllib.request.urlopen(req, timeout=3600) as r:
        body = json.loads(r.read())
    return body, time.time() - t0


def report(label: str, body: dict, wall: float, raw: bool = False) -> dict:
    usage = body.get("usage", {}) or {}
    timings = body.get("timings", {}) or {}
    prompt_n = timings.get("prompt_n") or usage.get("prompt_tokens")
    reused = usage.get("cached_tokens", 0) or 0
    read = (prompt_n - reused) if prompt_n is not None else None
    out = {"label": label, "prompt_n": prompt_n, "reused": reused, "read": read,
           "wall_s": round(wall, 2), "prompt_ms": timings.get("prompt_ms"),
           "prompt_tps": timings.get("prompt_per_second")}
    print(f"{label}: prompt {prompt_n} = {reused} reused + {read} read | "
          f"wall {wall:.2f}s | prompt {out['prompt_ms']} ms ({out['prompt_tps']} tok/s)")
    if raw:
        # prompt_n is not self-describing: on a checkpoint hit some builds report the
        # tokens computed this call rather than the prompt's full length. Dump the raw
        # fields so a "281" can be told apart from a truncated prompt instead of guessed.
        print(f"    raw usage   = {json.dumps(usage, sort_keys=True)}")
        print(f"    raw timings = {json.dumps(timings, sort_keys=True)}")
    return out


def main() -> None:
    ap = argparse.ArgumentParser()
    ap.add_argument("--url", default="http://127.0.0.1:8080")
    ap.add_argument("--words", type=int, default=5000, help="word count of the shared body")
    ap.add_argument("--max-tokens", type=int, default=24)
    a = ap.parse_args()

    # deterministic body: the two arms must be byte-identical up to the tail
    body = " ".join(VOCAB[i % len(VOCAB)] for i in range(a.words))
    head = f"Here is a list of words:\n{body}"
    prompt_a, prompt_b = head + TAIL_A, head + TAIL_B

    print(f"tail-fork probe: body {a.words} words")
    print(f"  A chars={len(prompt_a)}  B chars={len(prompt_b)}  shared prefix chars={len(head)}"
          f"  tails {len(TAIL_A)}/{len(TAIL_B)}")
    print(f"  url={a.url}  max_tokens={a.max_tokens}\n")

    ra, wa = post(a.url, prompt_a, a.max_tokens)
    a_row = report("A (cold, full prompt)      ", ra, wa, raw=True)

    rb, wb = post(a.url, prompt_b, a.max_tokens)
    b_row = report("B (tail fork, shared prefix)", rb, wb, raw=True)

    if a_row["prompt_n"] and b_row["reused"]:
        pct = 100.0 * b_row["reused"] / b_row["prompt_n"]
        print(f"\nB reuse: {pct:.1f}% of {b_row['prompt_n']} prompt tokens")
    print("B is the row to compare across link tiers.")


if __name__ == "__main__":
    main()

Happy to adjust the shape — e.g. a fan-out variant (one body, N different tails) — if that fits how you want to present suffix-reuse results.

🤖 Generated with Claude Code

[justindressler]

Tail-fork probe results — RTX 3090, IQ3_XXS, int8 KV

Ran your probe shape (A cold, then B = A byte-identical with only the final instruction replaced) against four flag arms. B rows below.

Parameters

Build	this PR's branch (v0.1.43 + PR patches), engine sha256 edcd90045016ed7dd119ac6ec729dd3ea8e7acd78ad81cf44360898c0f9b5570, verified against the deployed binary before the run
Model	Qwen3.8-Flash-Next (125B MoE), GSQ-RCO IQ3_XXS pack
KV / decode	int8 KV cache, MTP speculative (--spec 4 --spec-min-p 0.5), resident experts, --prefill auto
Hardware	single RTX 3090 24 GB (Ampere), consumer desktop, PCIe via OcuLink, driver 616.92, WSL2 — single GPU, no tensor parallelism
Requests	POST /v1/chat/completions, temperature 0, enable_thinking false, max_tokens 64; served through llama-swap, one preset per arm (flag sets below)
Probe	your tail_fork_probe.py prompt shape verbatim — same cyclic VOCAB, same TAIL_A/TAIL_B — driven by an equivalent stdlib-only runner

Environment caveat: this rig is not a clean benchmark host. The engine runs inside a container on WSL2, the model weights are read from a Windows NTFS drive through the WSL file bridge (container bind mount), and the GPU is reached through the container device bridge. Every storage-bound path is penalized relative to a native Linux install on local NVMe — engine cold start here is 100–200 s where a local filesystem would load a fraction of that time. The B rows are warm and compute-bound (reuse ratio and re-read tokens are unaffected by this), but the A-row cold-read numbers — including the visible part of the ladder's checkpoint-save tax — should be read as an upper bound set by this storage stack, not the mechanism's intrinsic cost on normal hardware.

Arms (identical except cache flags):

arm	flags
off	none (upstream defaults: --prompt-cache 6, --prompt-cache-every 16384)
sfx	--cache-suffix-reuse 64 --conversation-cache-mib 4096
ladder	--conversation-cache-mib 4096 --prompt-cache-every 1024 --prompt-cache-dense-tail 4096 --prompt-cache-state-every 256
all	sfx + ladder combined

Method

Per arm: cold engine start, one-token warm trigger, then A → B at each body size sequentially in one session (conversation parking stays warm between requests), unload + verified GPU drain before the next arm. reused is the engine's own usage.prompt_tokens_details.cached_tokens, cross-checked against the engine log's read-path: resume=N debug lines. On your prompt_n caveat: this build reported the full prompt length in both arms (5285 / 21035), so there was no computed-vs-full ambiguity to disambiguate. One pass per cell, no seeds — determinism is by construction (fixed cyclic body, temperature 0), but the wall numbers are single-shot.

Why 5000 and 20000 words

• w5000 (5,285 prompt tokens) sits entirely below the stock --prompt-cache-every 16384 cadence. It isolates the case where default checkpointing gives a fork nothing (0% reuse) and shows what the dense tail buys on its own.
• w20000 (21,035 prompt tokens) is a realistic multi-file coding-agent context and deliberately crosses the stock 16384 cadence, so the ladder's extra gain is attributable to tail density (checkpoints every 1024 within the last 4096, state every 256) rather than merely to having checkpoints at all.
• Together they show rung-nearest-fork behavior at two scales: with the ladder, B re-reads 1022 tokens at both sizes (the dense-tail granularity); on stock defaults the w20000 fork re-reads 3629 (resume at 16384, fork at ~20013).

B rows

arm	w5000 reused / read	wall	w20000 reused / read	wall
off	0 / 5285 (0%)	11.2 s	16384 / 21035 (77.9%)	10.6 s
sfx	0 / 5285 (0%)	11.7 s	16384 / 21035 (77.9%)	11.2 s
ladder	4263 / 5285 (80.6%)	5.4 s	20013 / 21035 (95.1%)	5.1 s
all	4263 / 5285 (80.6%)	5.1 s	20013 / 21035 (95.1%)	5.1 s

Mechanism (engine log, ladder arm, w20000 B): read-path: resume=20013 read_from=20013 n=21035 → prompt 21035 tokens = 20013 reused + 1022 read in 4142 ms (246.7 tok/s). splice_narrowed=0 throughout — expected for a tail fork: the strict checkpoint resume at the rung nearest the fork serves it, so this shape exercises the checkpoint ladder, not the splice path.

Cost side (A rows)

arm	w5000 A (cold read)	w20000 A
off	11.9 s, 0 cached	30.9 s, 0 cached
ladder	42.8 s, 0 cached	54.0 s, 4263 cached*

The ladder taxes the cold read (checkpoint + state saves during prefill). At w5000 that's 3.6× the stock cold read in this run — noticeably larger than the steady-state tax we measured on a live client (~+2 s per 3.9k prefill tokens with state-every 256), so the cold-read component isn't fully decomposed yet (state cadence vs full-checkpoint copies vs first-request expert paging). Break-even is ≈2 forks per cold read at these sizes; a sibling fan-out (one body, N tails) amortizes immediately.

(*) One artifact to disclose: the probe's cyclic 20-word vocabulary makes the w5000 body a strict prefix of the w20000 body, so within the ladder/all arms the parked w5000 conversation prefix-mounted onto A@w20000 (4263 reused). It slightly flatters A@w20000's wall; it cannot serve B (the fork at ~20013 is beyond the 5.3k-token park), so the B rows are unaffected.

Read

• sfx-only == off on this shape, as designed: A's end is not a prefix of B, so suffix-reuse alone cannot mount the fork. The parked conversation + checkpoint ladder is what serves it.
• The stock 16384 cadence already catches 77.9% at w20000; the PR's tail density closes the remaining gap (95.1%) and extends the win to bodies under the stock cadence (80.6% at w5000 vs 0%).
• B wall halves at both sizes (10.6 → 5.1 s at w20000; 11.2 → 5.4 s at w5000) — the sibling sub-agent waits about half as long.

Happy to run the fan-out variant (one body, N different tails) next — that's the shape where the cold-read tax amortizes fastest.

[Niko1221]

Thank you. This is a large engine change and we could not review and gate it in this pass. If you rebase it on 0.1.44.1 we will schedule it for its own round.

[justindressler]

Rebased onto v0.1.45 and re-verified end-to-end (head now 2ea89b6, force-pushed with lease from 24de45cc).

Rebase: 6 conflict groups resolved (struct-field union for Match incl. upstream's new serial; both-keep ordering in generate.cpp). Two rebase artifacts were caught by the test suite and fixed in e2dae10c — an orphan scope closer in the restore search, and Match initializer field alignment (16 partial-prefix failures → 82/82).

Full re-verification on the rebased head (RTX 3090, IQ3_XXS, int8 KV):

• Conversation suite 10/10, partial-prefix 82/82 (PR-2351 baseline), coverage-measured 95% host lines (269/269 cases)
• Live 4-turn probe: cold 26.9s → append 98.4% reuse (1.3s) → mid-edit splice fired at the edit token (gap re-read correct) → identical replay 99.8% (0.7s); all answers correct
• NIAH needle 21.9k-token haystack: 5/5 depths (10–90%), exact passphrase recall, unique filler per depth so no cache shortcuts
• Post-rebase independent review (6.7k-line diff): APPROVE-WITH-FINDINGS, 11 MINOR / 0 MAJOR / 0 BLOCKER — 9 confirmations of correctness, 2 optional items addressed in 2ea89b6

CUDA build: sm_86, toolkit 13.0, -j2. HIP remains compile-unverified (no AMD hardware here) — unchanged from the original submission.

[Niko1221]

Thanks for the pull request.

We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.

To reopen it takes about two minutes:

1. Rebase on main if it no longer applies, and re-run your checks on 0.2.0-beta.
2. Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other].
3. Fill in the pull request template: the problem and result, and the validation (what you ran, on which revision, what you saw).
4. Comment /reopen. A check runs and either reopens the pull request or lists what is still missing.

The rules are in CONTRIBUTING.md.

[justindressler]

/reopen

[github-actions]

The check passed, but GitHub would not reopen it (Validation Failed: ). If the branch was deleted, restore it or open a new pull request and link this one.