Skip to content

Conversation cache: keep the chat between requests, read only what is new - #8

Merged
Niko1221 merged 1 commit into
Niko1221:mainfrom
Mirtraxxx:conversation-cache
Sep 26, 2026
Merged

Niko1221 merged 1 commit into
Niko1221:mainfrom
Mirtraxxx:conversation-cache

Conversation

@Mirtraxxx

Copy link
Copy Markdown
Contributor

What it does

Right now every request starts from an empty session and reads its whole prompt again. In a long chat or an agent
loop, every turn waits for the full prefill. That's the main complaint in the r/LocalLLaMA thread ("throws away the
KV cache every prompt").

With this change the engine keeps the conversation and reads only the new part:

  • Live session. If the new prompt starts with exactly the tokens the session already holds, it just continues.
  • Checkpoints. Up to 6 are kept in RAM, about 118 MB each. One is taken at the last <|im_start|> of every
    prompt, where the history ends and the new assistant header begins, and one every 16K freshly read tokens. A
    rebuilt chat history (the client re-sends the messages, the template re-renders them) still matches the one
    taken before its last turn, so only the last reply and the new message are read.
  • Only the running state is copied. That's the 36 GDN recurrences and conv histories, the 12 QSA indexer tails
    and the PLE history: the same set Verifier::commit rolls back. The KV cache and the pooled indexer keys are
    positional. A query only reads cells below it, and dead for its own block, so stale cells past a rewind point
    are never read and get overwritten.
  • Safety. A checkpoint is used only if the prompt starts with exactly its tokens and pictures. Images are
    compared by a hash of their embeddings and grid, so a different picture of the same size never matches.
    Checkpoints whose cells the next request will overwrite are dropped.

Measured

RTX 3090 24 GB, Ryzen 7 5700X, 64 GB DDR4-3000, Windows 11, Swift 1.5 IQ2_XS, --expert-cache auto, MTP --spec 4.

Agent-style session through serve/server.py: OpenAI streaming, one tool, thinking on, reasoning sent back each turn.

Turn Prompt Reused Read Time to first token
1 19,355 0 19,355 56.2 s
2 (tool result added) 20,882 19,407 1,475 5.6 s
3 21,189 20,877 312 2.9 s
4 21,355 21,184 171 2.6 s

Is it exact?

  • STRATA_CKPT_REREAD=1 makes the engine read a checkpoint's tokens again from position 0, in the same chunks the
    request that saved it used, instead of restoring it.
  • STRATA_STATE_HASH=1 prints a fingerprint of every part of the session after each request: GDN states, PLE
    history, indexer tails, pooled keys, the KV cells [0, L) of all 12 QSA layers, the MTP K/V, and ple_prev.
  • A 4-turn chat (19K-token first prompt, two rebuilt follow-ups, one exact continuation) with --adapt-swaps 0, so
    the VRAM expert set is fixed. Restore and re-read gave identical fingerprints for every component after every
    turn
    , and identical tokens and draft counts.
  • Two runs of the restore path were also identical to each other.

Protocol (backward compatible)

  • New stdout lines: RESUME <n> before reading, PP <pos> <n> <ms> <fresh tok/s> per chunk, and REUSED <n> once
    the prompt is read.
  • DONE gains <drafts accepted> <drafts offered> <reused> after the existing fields.
  • serve/server.py ignores lines it doesn't know. One change: PP lines now count as a heartbeat. They arrive
    every few seconds, which kept resetting the 10 s wait, so a long prompt sent no keep-alives at all: 0 in a
    41-second prefill before this change, 6 after.

Options

  • --prompt-cache N: checkpoints kept (default 6; 0 restores the old behaviour).
  • --prompt-cache-every N: periodic checkpoint spacing (default 16384; 0 = only at the turn boundary).
  • --turn-token ID: the token that opens a turn (default 248045, <|im_start|>).

Limits

  • One conversation is cached at a time. Two chats that take turns overwrite each other's state and re-read.
  • Small reads still pay the prompt path's fixed cost (lending and refilling the borrowed cache slots): about 1–2.5 s
    even for a few dozen tokens.
  • Only tested on Windows / one RTX 3090. Images through the cache are designed for but not tested end to end.

Made with Claude Code (Claude Opus 5.5); tested on the hardware above.

🤖 Generated with Claude Code

… what is new

Every request used to start from session_zero and read its whole prompt again,
so a long chat or an agent loop waited minutes per turn. The engine now keeps
the live session and up to 6 checkpoints of the running state in RAM and reads
only the part of a prompt after the longest one it starts with.

- A checkpoint is the state that cannot be rewound - the 36 GDN recurrences and
  conv histories, the 12 QSA indexer tails, the PLE history (~118 MB). The KV
  cache and the pooled indexer keys are positional (a query reads only cells
  below it, and `dead` for its own block), so they stay in place.
- Taken at the last <|im_start|> of each prompt (the prompt is read in two parts
  there: the history, then the new turn's header) and every 16K fresh tokens.
- Used only when the prompt starts with exactly its tokens and pictures (images
  are compared by a hash of their embeddings and grid); checkpoints the next
  request would overwrite are dropped.
- New lines on stdout: RESUME <n> before reading, PP <pos> <n> <ms> <tok/s> per
  chunk, REUSED <n> after; DONE gains <drafts accepted> <drafts offered>
  <reused>. The server ignores what it does not know; PP lines now count as its
  heartbeat, so keep-alives keep flowing during a long prompt.
- --prompt-cache N (default 6, 0 = off), --prompt-cache-every N, --turn-token ID.
- Checks: STRATA_CKPT_REREAD=1 reads a checkpoint's tokens again instead of
  restoring it; STRATA_STATE_HASH=1 prints a fingerprint of every state part
  after each request.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@Niko1221
Niko1221 merged commit 23f6b15 into Niko1221:main Sep 26, 2026
Niko1221 added a commit that referenced this pull request Sep 26, 2026
…gh the decode windows (#8, #10)

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

Copy link
Copy Markdown
Owner

Honestly this is huge, thank you. I was halfway through my own cache and dropped it for yours, since yours does more (live continuation, pictures, the reread/hash checks). On my 5070 follow-ups went from ~13 s to 0.3 s, and restore vs reread gave the same tokens and the same state hashes on every turn. Out in v0.1.3.

MistVVK referenced this pull request in MistVVK/XeStrata Oct 3, 2026
MistVVK referenced this pull request in MistVVK/XeStrata Oct 10, 2026
…prefill and decode attention keep the cached rows FP16 in shared memory (34 KB, was 66 KB, over Turing's 64 KB); found and diagnosed by @dummerjindabin (#8). AMD signal fences (#2, @boxwrench)

The F32 prefill attention asked for 67,584 bytes of shared memory and ignored the failed opt-in, so on Turing every
launch failed and the engine died on any prompt over ~1k tokens; the decode split attention made the same request
and quietly fell back to the one-pass kernel. Both now stage the FP16 latents as they are in the cache and read
them as F32: the same values and arithmetic, 34,816 bytes (no opt-in needed). Tesla V100, Maya-S, 3000-token
prompt + 200 greedy tokens: the same tokens before/after on the default path and on the forced F32 prompt kernel;
the F32 kernel's attention check against the tensor-core kernel prints identically before/after.
MistVVK referenced this pull request in MistVVK/XeStrata Oct 10, 2026
…2k figure was the first prompt after start-up, caches cold; #8), and prompt-only requests show just their prefill time in --bench / --report (not a 96,000 tok/s decode)
MistVVK referenced this pull request in MistVVK/XeStrata Oct 10, 2026
… the accumulator registers

Second half of #8 - the FP16 staging fix for the same issue shipped in v1.0.7, so this is the
remaining piece.

mla_attn_tc_kernel asked for 91136 bytes of dynamic shared memory a block (sO 32 KB + sQ 16 KB +
sL 32 KB + sPart 8 KB + sP 1 KB), over Turing's 65536 ceiling: the opt-in failed, tc_ok went -1
and sm_75 always ran the F32 kernel, so the faster prefill never reached RTX 20-series.

The context O now rides the wmma accumulator fragments: a warp owns four 16-wide column tiles
(TC_OT = KV / 16 / 8 - eight warps cover the 512-wide context exactly once) for the whole cell
walk, and the per-chunk online-softmax rescale multiplies them in place. In the m16n16k16 f32
accumulator layout lane l holds rows l>>2 in x[0],x[1],x[4],x[5] and rows (l>>2)+8 in
x[2],x[3],x[6],x[7], which is where the per-row factors go. The tiles are stored into ctx at the
end, and a __syncthreads() before the row normalization makes sure every warp's tiles are there
first. The block's shared memory drops to 58368 bytes, which the same opt-in grants (58368 <=
65536), so the kernel runs on Turing and on Volta unchanged.

Measured on 2x Titan RTX (sm_75, CUDA 13.0, driver 610.57.04), Maya-S IQ2_XXS, 96K context, one
18476-token prompt, greedy. The engine logs which prompt kernel it picked ("prompt attention on
the F32 cores" / "on the tensor cores"), so the A/B is unambiguous:

  f32 forced : prefill 62 s, prompt 3.4 ms/tok (~294 tok/s), decode 12.6 tok/s
  this kernel: prefill 58 s, prompt 3.1 ms/tok (~323 tok/s), decode 11.3 tok/s

+7..10%, inside this machine's run-to-run spread. STRATA_GLM_PREFILL_ATTN_CHECK=1 over 110
chunk checks: rel L2 5.3e-5 .. 1.6e-4 (median 1.1e-4), worst head 2.5e-4 .. 1.3e-3 (median
6.4e-4). The prompt's three buried codes come back 3/3 and the runs have no invalid argument /
illegal memory access / planned 0 of routed rows lines. No sm_70 card here, so the V100 A/B is
still the gate.
gopinath87607 pushed a commit to gopinath87607/Strata that referenced this pull request Oct 20, 2026
From the Maya checkout (`project-maya`), local commit `cc58309` (2026-10-14): `STRATA_GLM_WARM_PROMPT` there, and
that name is accepted here too, so a config written for the Maya checkout runs unchanged.

A session's first prompt pays what no later one does: the first chunk carries every kernel's first launch, the
prompt path's own first-chunk work and the file tier's cold reads. Maya's issue Niko1221#8 reported a first prompt at
95-133 ms/token against 23-44 warm; that is a number from that checkout on its rig, and whether the knob wins on
any other rig is what the first-prompt A/B with it on and off decides - so the default stays 0 and nothing runs
without it.

The warm prompt goes through the whole part chain before READY with the request-bound callbacks moved aside and
put back around it (on_chunk/on_stage_chunk read the request's own tokens, print its PP lines and save its
checkpoints: a warm prompt has none of those), and it takes the same loan a request takes - each participant's
own rows in the slots the prompt path is laid out in go non-resident, the residency table goes out, and after
the run the slots are refilled from the arena into the same slots through their own cache. A slot left holding
a prompt buffer while still marked resident would feed the next decode an expert that is not there; a slot whose
refill fails stays non-resident, which is the truth for it.

Every participant's state is zeroed before and after the run, as the request path zeroes at resume == 0 - the
session's memory is not zeroed at start, so a warm prompt read without that would read whatever is there - and
the knob never fails the start: a prompt path that cannot run prints why (or the run's own error) and the first
prompt pays it. The elastic K/V is grown to the warm prompt's cells first, as the non-serve prompt path does.
The watchdog guards the run as it guards a request: busy is set around it, and the prompt path beats a layer at
a time.

Not measured here: what it buys and what it costs on a given rig is the A/B the knob exists for - the first
request's time to first token (or its PP line), with STRATA_WARM_PROMPT=<n> and unset.

Co-Authored-By: Claude Code <noreply@anthropic.com>
svgop added a commit to svgop/Strata that referenced this pull request Oct 20, 2026
…v6 M1 +5.2%/conc +5.5%, bit-exact, rollback STRATA_GR_NORM_FAST=0); staged-down 4-warp widen REJECTED (M1 -6.99%, reverted); UPM 16->8 proved bit-exact but under bar solo (+3.83%) and FLAT in stack - flagged default-off (STRATA_GR_UP8); LFUSE device-plan arm adjudicated DEAD by code read (F-206 sweep + FUSE_HITS displacement, T16 Niko1221#8 resolved); proof-first gr_parity byte-equality at arm state before every boot + per-boot hc selftest in every battery leg; t15c-battery.sh/t15c-score.py harnesses + 4 battery logs (30 legs)
RandoTeam referenced this pull request in RandoTeam/4gbVramStrata Nov 1, 2026
[Docs] the numbers of the engine this release ships
qwertie added a commit to qwertie/Strataq that referenced this pull request Nov 4, 2026
… off

DSpark's 752 MB of dense weights stayed loaded with drafting off,
memory the page pools could use. strata-ds41 now frees them when
drafting is switched off (or never on) and reloads them when it comes
back, without a restart; the shared paged engine gives the freed bytes
to the pool and takes them back (537 MB after chunk rounding). A
switch while a conversation generates waits until that answer ends.
Answers continue unchanged across on -> off -> on, pauses and rewinds.

memory_report.hpp and the Memory panel (strata_server/web/memory-panel.js)
now show the drafter's experts (7.2 GB) apart from the main model's, so
users see what drafting costs.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
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.

2 participants