feat(v4): prefix checkpoints (memory + disk), resumable 4096-token prefill segments, cancellable prefill, gateway prefix hint

[dcutugno]

Stacked on #988 (dual-SSD mirror); the diff against dev includes it until #988 merges.

Problem

An agent client (opencode, aider…) sends an 8k-token system prompt and then a new
request per turn/session. V4 keeps one KV slot and its window/compressor/indexer
state cannot rewind, so every request re-prefilled the whole conversation
(measured: turn 2 re-prefilled 8548 tokens = 677 s on the CPU path, well past
opencode's 30-minute budget once decode is added; an aborted client also left
the engine finishing the orphaned generation, wedging the pipe for the next one).

What this adds (all CPU-side, no CUDA)

• Prefix checkpoints (V4_PREFIX_CKPT, on): a snapshot of the attention
  transaction (the existing ColiV4AttentionSnapshot: window KV + compressed
  slots + compressor/indexer recurrent state) is captured at a shared boundary
  and restored for any later request whose prompt starts with the same bytes.
  Three capture rules: (1) the gateway passes the byte offset where the rendered
  system turn ends (optional 8th SUBMIT header field / prefix_bytes in the
  session options — older engines ignore it), so the first request already
  seeds the checkpoint; (2) fallback: longest common prefix of two successive
  fresh prompts; (3) prompt-end capture after every prefill, because agents
  re-render the assistant reply so "extends everything fed" fails at the reply
  boundary. 4 in-memory slots, LRU, prompt-end evicted first.
• Disk persistence (V4_PREFIX_CKPT_DISK, on): prefix captures go to
  <model>/.coli_ckpt/ (config-fingerprinted, temp+rename, ~140 MB for an 8.3k
  prefix) and are loaded lazily after a serve restart. New
  coli_v4_{attention,compressor,indexer}_snapshot_write/read.
• Segments + cancel + resume: prefill runs in atomic 4096-token segments
  (V4_PREFILL_SEGMENT); a new should_abort poll in the session options runs
  between segments (the gateway wires it to client disconnect), completed
  segments are recorded so an identical retry resumes; the gateway keeps the
  scheduler admission until the engine acks the CANCEL (this is the wedge fix).
• 128-token chunks (V4_PREFILL_CHUNK clamp raised 64→128, +3 % CPU prefill;
  the fp8 batch reference's sums[] grows accordingly).
• Batched indexer selection in prefill (V4_IDX_BATCH, on): advance per
  token, score/select once per chunk past index_topk; per-token numerics
  unchanged. V4_IDX_IDENTITY (off) documents the faster identity
  short-circuit that changes rounding.
• Progress lines per segment, V4_PREFIX_LOG diagnostics, docs section
  "Prefill: segments, chunks, checkpoints", env table.

Numbers (CPU path, 8.3k-token opencode system prompt)

• first turn: full prefill once per model (also survives restarts via disk)
• every later session / turn: seconds (restore + new tokens) instead of a full
  re-prefill
• (with the CUDA tier of the follow-up PR: 8.3k first turn ~4 min, later turns 6–9 s)

Validation

• CPU parity vs dev: 826-token prompt, 48 greedy tokens, --memory-gb 22,
  COLI_V4_ROWS16=0 COLI_V4_AUTOPIN=0 COLI_V4_SAVE_USAGE=0 V4_PREFIX_CKPT=0
  → byte-identical text (md5 9cf553633292 on this branch and on upstream dev,
  2026-08-16). Why those variables — and why greedy text is not a stable
  identity check across kernel/cache configurations — is written up in
  docs/deepseek-v4.md → Validation.
• tiny fixture (make deepseek-v4-tiny-check) token-exact; two-request
  SUBMIT/DATA/DONE flow; cancel-during-prefill exercised through the gateway.

Files: c/deepseek_v4.c (+~1.3k), c/deepseek_v4_internal.h, c/deepseek_v4.h
(should_abort, prefix_bytes), c/openai_server.py (prefix hint, cancel ack,
KeyboardInterrupt), docs.

---

🤖 Generated with Claude Code

[JustVugg]

Reviewed this properly, and I want it — the problem it solves is real and the measurement makes it undeniable: 677 s to re-prefill 8,548 tokens on turn 2, past opencode's entire budget before decode even starts. An engine that re-prefills the whole transcript every turn isn't usable with an agent client, and this is the fix.

The design is right too. Snapshotting the attention transaction at a shared boundary and restoring it for a matching byte prefix is the only approach that works given V4's window/compressor/indexer state can't rewind, and the three capture rules (gateway prefix_bytes hint, longest-common-prefix fallback, prompt-end capture) cover the case that would otherwise break it — agents re-rendering the assistant reply, so "extends everything fed" fails at the reply boundary. The should_abort poll between atomic segments is the right shape for the wedge fix as well: cancellation at a boundary, with completed segments recorded so a retry resumes.

I checked the thing I was most worried about, and you got it right

The V4_PREFILL_CHUNK clamp going 64 → 128. That ceiling was mine, and it was there because five kernels validate batch > 64 and a 256-wide chunk dies at layer 0 within seconds. So I verified the whole set on your branch rather than trusting the summary:

5 guards, all raised:            batch > 128   (none left behind at 64)
the stack array grown with them: __m256 sums[128]   (was sums[64])

Raising a bound and growing the fixed-size buffer it protects, in the same change, is exactly what that needed — a guard raised without the array is a stack smash. Thank you for doing both. And +3% for 128 = 103.9 s vs 64 = 107.3 s, text identical is honest reporting of a small win.

Also noted and appreciated: V4_IDX_IDENTITY shipped off, documented as changing rounding. That's the right instinct in this repo.

One change before I merge: please make the disk half opt-in

V4_PREFIX_CKPT_DISK defaults on, which means every V4 user's model directory starts growing by ~140 MB per prefix without having asked for it. Model directories are somewhere people reasonably treat as read-only inputs — they mount them read-only, they put them on full disks, they sync them, they check free space against the checkpoint size and nothing else.

I'd like:

• V4_PREFIX_CKPT (in-memory) stays on. That's where the 677 s → seconds win comes from, it costs 4 LRU slots, and it disappears when the process does. No reason to hide it behind a flag.
• V4_PREFIX_CKPT_DISK defaults off, opt-in for people who want the win to survive a serve restart. Documented with the size (~140 MB per prefix at 8.3k) so the choice is informed.

If it stays on by default, please at least make an unwritable or full model directory fail soft and silent-ish — a warning, not an error, and certainly not a failed request. But I'd rather have the default flipped: this is the same reasoning I applied to #1058 this week, where LTO-by-default changes something the user didn't ask to change.

Everything else here is ready. clean, 19/19 green, and it stopped carrying #988's diff now that that's landed. Flip the disk default and I'll take it.

Separately: #1055 can rebase now too — both #988 and #1054 (@ZacharyZcR's kernels, with his test and oracle suite) are in dev, so its diff should collapse to its own three commits and finally be reviewable as the engine wiring.

[ZacharyZcR]

Small catch from reading the merged tree: the review asked for V4_PREFIX_CKPT_DISK to default off (same reasoning as rejecting LTO-by-default in #1058), but c/deepseek_v4.c:11493 still ships v4_ckpt_disk_mode = setting ? atoi(setting) : 1 — on by default. The fail-soft behavior you asked for is there (v4_ckpt_disk_write returns silently on an unwritable/full model dir), so this is only the default, not a crash path.

Happy to send the one-line flip (: 1 → : 0) plus a doc line stating the ~140 MB per prefix cost at 8.3k tokens, if that matches what you intended at review time.