#3584 [Engine] suffix-reuse: conversation KV prefix mount, quant-safe splices, state checkpoint ladder

Description

## 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.
- This replaces **#2351**, which was closed in the 2026-10-19 tracker reset for 0.2.0-beta, not on
  content. The repo's `reopen` check passes on this body, but GitHub declined to reopen the old PR
  ("Validation Failed"), so this is a new pull request from the same branch (`sfx-group`, now
  `cf6c6984`) with the rebase applied and the template filled in.

## 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.