Repository navigation
Conversation
Closed
Whole-session snapshots (attention KV in every format + idx_pooled/idx_block_pos/ idx_dead/idx_tail + GDN/PLE running state + MTP drafter host KV) keyed by the exact token prefix, stored under --kv-nvme DIR: - automatic synchronous cascade on DONE (idempotent; one growing file per conversation via strict-prefix supersede), fsync'd - automatic promote on prefix match in the serve resume path - no client calls (a failed promote falls back to the clean session_zero path, never a mixed state) - LRU byte cap (--kv-nvme-max GB); corrupt/foreign store files are refused, never converted, and can no longer crash the server - atomic restore: the whole snapshot is validated before anything is applied - STRATA_STATE_HASH refactored into state_hash_line() and shared with a new post-restore oracle (STRATA_NVME_HASH); verify.cpp: add <cmath> for Linux Oracles: tools/nvme_p0_test.sh (bit-exact restore + negative control), tools/nvme_steps123_test.sh (cascade/supersede, promote-after-restart: 13 fresh tokens read in ~350 ms after a restart vs a full re-prefill; LRU cap evicts). Known engine property (documented): decode is run-to-run nondeterministic on v0.1.13 (pre-existing, reproduced with the old binary), so continuation token equality is informational; restore-exactness is the hard gate. Design: docs/nvme-kv-cache-design.md; review: docs/nvme-kv-cache-critique.md (glm-5.3, design) and docs/nvme-kv-impl-review.md (glm-5.3, implementation, P1s addressed).
A chat client re-sends history without the model's hidden reasoning tokens, so a full consumed-state snapshot could never full-prefix-match the next turn (the first live HTTP needle test showed every request falling back to a full re-prefill). The longest ConvCheckpoint supplies the running state as of the end of the prompt; KV cells / pooled rows / dead key truncate to that position. The next request's prompt extends exactly that prefix, so the match is exact. Live-server verification (three ~110k-token sessions, needle ~900 tokens in): all three needle turns promoted from NVMe (reused 110k+, 39-42 fresh tokens read, vs 110 s full prefill cold), TTFT 110 s -> 1.3-1.9 s (snapshot read dominates), needle retrieved. Config: --prefill auto.
…e review annexes The design doc now matches the shipped implementation (turn-boundary snapshots, payload digest footer, refuse-never-convert, graceful promote fallback, LRU cap) and the test record: P0 bit-exactness, steps 1-3, the live 110k needle rotation (110 s -> 1.3-1.9 s TTFT), and the short correctness suite (chaining, branching, shared-prefix isolation, rotation, idempotence, corruption fallback). Known follow-ups: streamed promote, per-conversation supersession, sparsity.
FNV-1a over the payload (everything after the header), written as an 8-byte
footer and verified by nvme_restore before anything is applied: a flipped or
truncated byte is refused ('integrity check failed') instead of silently
decoded. open() skips files too small to carry a footer.
Found by the corruption-fallback test: the first version of this check was
itself broken twice (the turn-boundary tail blobs were written unhashed, and
the restore's hash scope included the header the writer never hashed) - every
snapshot failed its own digest, deterministically. Both fixed; the store's
files now verify externally and in-engine. Old footer-less files are refused
and must be re-dumped.
Check the walk end and the drafter-array count before the digest, with the numbers in the message: a snapshot written by an older engine layout used to misreport as 'corrupt snapshot' when the real cause was a sizing change (idx_pooled_rows / PLE / drafter ring). Also validated on the rebased engine (v0.1.15): P0 ALL PASS (bit-exact restore, resume, negative control) and the HTTP correctness suite ALL PASS (chaining, branching, shared-prefix isolation, rotation, idempotence, corruption fallback).
… (step 2, unfinished) This is a safety checkpoint, not a finished step. The step-2 child was killed by a server restart before it could commit, and its work existed only as uncommitted changes in the worktree. Committing it here so the next child builds on a durable base instead of re-deriving it. What it does: - imports the shared core from jrich/feat/conversation-cache-shared-core: conversation_cache.hpp, conversation_snapshot.hpp, conversation_memory.hpp, conversation_state.cpp, conversation_snapshot.cpp, conversation_checked.hpp - ConvCheckpoint / ConvStateSizes / ImgKey become aliases of the shared types; checkpoint_save, checkpoint_restore and conv_state_sizes become one-line wrappers - NvmeRunning raw pointers replaced by const ConversationCheckpoint*, so the snapshot key and its running state cannot disagree about which point in the conversation is stored - turn-boundary checkpoint blobs are validated against conversation_state_sizes instead of trusted because raw pointers were handed over - static_assert pins sizeof/alignof of ConversationImageKey to the old pair, proving the v2 envelope's 16-byte image record did not move - dead / block_pos still read from the live device arrays; collisions C4 and C5 deliberately untouched - their RAM cache policy is imported dormant Verified: build/strata links with 17 strata::core::conversation_* symbols.
include/strata/platform/kv_nvme.hpp:30 tells the reader to see
docs/nvme-kv-cache-convergence.md ("Where step 2 leaves NvmeHeader"). That section
did not exist: the closest thing was C9, which states the rule but never inventories
what our format actually does. The reference was dangling.
Written against the code, not against the design: the four ways NvmeHeader breaks the
shared core's disk-adapter boundary (raw struct memcpy at kv_nvme.cpp:176/288/412, an
implicit ABI whose padding the payload digest deliberately does not cover at :362-363,
no versioned envelope - the restore re-derives every segment length from the live
engine and reports drift instead of negotiating at :344-356, and a geometry tag that is
a derived projection of eight computed fields rather than the shared core's 18-field
geometry_key, which their own type says is not a disk schema either), and what step 3
has to turn it into.
… conversation with a picture was unreachable
What was wrong: KvNvmeStore::dump re-keyed the snapshot at the turn boundary (key = at->ids,
T tokens) but still passed the LIVE image list down to nvme_dump_at. That list is
imgs_below(req_imgs, live.size()) - every picture in the whole conversation, including the
ones in the user message being answered now, whose `start` is at or past T. So the file's
image segment described a longer prefix than its key did.
The consequence was total, not cosmetic. The resume match (generate.cpp:2939) compares
imgs_below(req_imgs, e.L) - the request's pictures strictly below the entry length - against
what the entry holds. A stored list containing a picture at start >= e.L can never equal
it, so an entry for any conversation that had an image was skipped on every request, and the
idempotent exact-match skip in dump() never fired for one either. The tier silently did
nothing for image sessions while reporting a successful dump.
Pre-existing (it predates 114300f: the old code passed d_imgs, built from the same live
list, alongside &at->ids), but 114300f is what makes it fixable in one line: the boundary
now arrives as ONE ConversationCheckpoint, and that checkpoint already carries the filtered
list - checkpoint_at builds c.imgs as imgs_below(req_imgs, c.ids.size()). Taking at->imgs is
the completion of that retype, not a new design: C1's rule ("a spill re-keys at the turn
boundary; it does not copy the parked image") was already decided, and the image list is part
of the key, not a format decision. No field, offset or sizing formula moves.
nvme_dump_at now also refuses an image record outside [0, L) outright, so a caller cannot
write a snapshot whose pictures contradict its own prefix.
…in the tier compared against nothing What was wrong: NvmeEntry carries `imgs`, and two places read it - the resume match (generate.cpp:2939, `imgs_below(req_imgs, EL) == e.imgs`) and the idempotent exact-match skip in KvNvmeStore::dump. Nothing ever wrote it. KvNvmeStore::open read the header and the ids from each snapshot and stopped there; KvNvmeStore::dump built the new entry from path/L/cvec/ bytes/mtime and left the vector empty. Every entry in the tree, scanned or freshly dumped, claimed to be a snapshot of a conversation with no pictures. So for any session that had an image, imgs_below(req_imgs, EL) was non-empty against an empty e.imgs: the entry was skipped on every request, and dump() re-wrote the same snapshot on every turn instead of refreshing its recency. Text-only sessions worked, which is why it survived every test run so far - the needle bench and the P0/step tests are text. It also made `live_imgs = best->imgs` (generate.cpp:2956) clear the live picture list after a promote. Pre-existing since 0dd4e3e, and it is the other half of the previous commit: filtering the stored list to the boundary only helps if the entry actually carries it. open() now reads the h.n_imgs 16-byte records that follow the ids (the size check above already guarantees the file has room for them), dump() stores the list it wrote, and restore() compares the restored list against the entry's - the same "the file changed under us" guard it already applies to L and cvec.
…y at every call site
What was wrong: 114300f routed the checkpoint types and the byte counts through the shared
core, and that core reports every refusal as a string naming what it refused
("conversation snapshot: invalid model geometry", "... invalid checkpoint running-state
payload", "... indexer byte count overflow"). All four wrappers built the string and dropped
it. Two of them dropped the RESULT too: kv_nvme.cpp's sizes_of ignored
conversation_state_sizes' bool, so a geometry the core refused left every byte count at zero
and the envelope was still laid out, written and walked with zero-length gdn/ple/tail/dead/
block_pos segments - the failure surfaced, if at all, as a layout mismatch on the next restore
rather than as the field that was rejected.
sizes_of now returns bool and both envelope paths (nvme_dump_at, nv_restore) refuse before
sizing anything. The generate.cpp wrappers keep their bool signature - every caller already
handles false - and print the core's reason, which is what the generic "saving a conversation
checkpoint failed" was hiding.
…, not its LAYOUT What was wrong with it: the assert 114300f added compares sizeof and alignof of ConversationImageKey against std::pair<int64_t, uint64_t>. That catches a field being added or retyped, but a field REORDERED (hash first, start second) keeps both numbers at 16 and 8 and silently swaps the two halves of every image record in every v2 file - and the resume match compares those records field by field. The comment claimed the bytes were pinned; the assert did not pin them. Adding is_standard_layout closes the gap: a standard-layout class lays its non-static data members out in declaration order, so with the existing comparisons the record is 16 bytes, 8-aligned, and (start, hash) in that order. No offset is hardcoded here - stating the file's field offsets is step 3's envelope, not this step's.
What was wrong: the envelope comment called block_pos "the per-token scratch [that] rides along harmlessly". It is not per-token and it is not harmless-by-accident: it is ONE int32 per QSA layer of device-internal scratch - the pooling kernel writes the completed block's first-cell position into it and the rotation reads it back on the device, which is the only reason the field exists at all (qsa.hpp:199-204). It is safe to carry the live value across a turn-boundary snapshot because it is rewritten before anything can read a stale one, and that is the reason worth having in the comment. Checked while I was there, since the envelope now sizes this segment from the shared core instead of a literal: conversation_state_sizes sets block_pos to sizeof(int32_t) = 4, the same 4 the dump and the restore hardcoded before 114300f, so the segment did not move. No code change; whose copy of dead/block_pos the envelope owns is still C5.
… step 4 so it cannot be lost
Two things this branch must not lose:
- WHY the two image fixes are in step 2 and not in C1. C1's rule was already decided ("a spill
re-keys at the turn boundary; it does not copy the parked image"); the image list is part of the
key, not a format choice, and a snapshot whose pictures describe a longer prefix than its ids do
is unreachable rather than merely wrong. What C1 still owes is the DECISION about spilling a
parked (RAM-evicted) session, which is untouched here.
- THE FIXTURE. Both bugs survived every run this tier has had because every test we have is
text-only. The fixture step 4 must add is named exactly: a boundary that is not idx_block-
aligned, an image whose start is at or past that boundary, and a second request that re-sends
the prompt without the model's reasoning tokens - asserting a promote, not a successful dump.
The three existing nvme scripts all need the GPU this branch may not take, so until that fixture
exists the image path of the tier is untested, and this branch says so out loud.
Also lists what step 3 inherits (C1's spill policy, C4/C5/C6 as the consumers of the NvmeHeader
section, C7-C11 untouched) and states that the RAM policy is dormant ON PURPOSE: only
conversation_state_sizes and conversation_checkpoint_{save,restore} are live on our side.
…e device arrays (C5) At a turn boundary T the live idx_block_pos names a block COMPLETED BY TOKENS PAST T (qsa.cu:216 rewrites it with each completed block's first-cell position), so reading it off the device wrote a running-state value that does not describe the prefix the file is keyed by. dead is the cell-0 key (qsa.cu:187, pos == 0 only) and is constant for the sequence, so the two copies agree there - but the pair now comes from the same source as gdn/ple/tails: the boundary checkpoint when one is given, the live device arrays otherwise, where the live arrays ARE the state at L. The checkpoint's dead/block_pos blobs are now size-checked against conversation_state_sizes before a byte of them is written, like its other blobs. The segments did not move: same sizes, same order. Only their source changed, and a reader cannot tell the two apart from the bytes, so NvmeHeader::version goes 2 -> 3 and version 2 is REFUSED rather than reinterpreted - by nvme_restore (naming the version found) and by the startup scan, which now counts a stale-format store separately from a malformed one and prints the version it found.
Two formulas sized the same segment: ours min(L/idx_block + 2, idx_pooled_rows), theirs upto/idx_block + 1 (conversation_snapshot.cpp:45). Theirs wins, and the row it keeps is the load-bearing one: qsa.cu:213 and native_qsa_indexer.cu:93 keep pooled[n_bid] == dead, and qsa.cu:260 and native_qsa_score.cu:74 read row n_bid straight out of the pool - so a snapshot that stopped at the completed rows would leave the block in progress scored with a stale key. Our +2 wrote one row past that (what the live array happens to hold at dump time) which no reader can reach: qsa.cu:253 gates b > n_bid, qsa_select.cu:33 gates b > n_bid and reads dead for b == n_bid (so its highest pooled read is n_bid - 1), native_qsa_score.cu:74 gates row <= full. The writer does touch row n_bid+1 when a block completes (qsa.cu:213 seeds it with dead), so a stale value there is overwritten before any read. Their refuse-instead-of-clamp rule comes with the formula: a live pooled array too small for the snapshot is now a refusal naming both counts, where the old min() wrote a short segment the restore reported as layout drift. The formula is stated ONCE (kv_nvme.cpp snapshot_pooled_rows) and both the dump and the restore call it, so the writer and the reader cannot drift apart again. Part of format version 3.
…ly the resumes it owns (C6) nvme_restore now calls refill_drafter_ring after its apply pass: the same [b1 - n_slots, b1) blocks the serve loop's mtp.kv_restore computed, clamped to min(L, max_cells) because that is the window the adapter restored. The --nvme-restore spike's mtp.kv_restore(rL) is deleted. Not their core: reaching kv_ring_restore there means routing the drafter KV through a ConversationKv staging vector, which contradicts the bounded-staging answer we give to their open question. Not the serve loop: the refill is part of the residency contract for bytes the tier just wrote, and a call-site rule is what a new spill or restore path forgets. The adapter already owns the main layers' kv_stream_reset, so this extends behaviour it has rather than adding a second owner of that contract. Collapse condition written into docs/nvme-kv-cache-convergence.md: it folds into conversation_kv_restore the moment the adapter adopts that wholesale (C10 / staging). The serve loop's mtp.kv_restore(resume) stays for the resumes that never touched the adapter - a RAM checkpoint (from_live false) or the live session. Deleting it was proposed and rejected: their conversation_checkpoint_restore restores gdn/ple/tails/dead/block_pos only, no KV and no ring, so the RAM tier would keep ring cells from a longer or diverged turn. A promote is indistinguishable from a live resume at from_live (both set it true), so the promote now sets from_nvme and the rule is gated on it. Cells past the promoted L are left in the ring deliberately, and that was checked rather than assumed: the ring table is static (block -> block % n_slots, kv_stream.hpp:76-79), so the slots a longer previous turn could have clobbered are exactly the blocks below b1 - the ones the refill repopulates - and the drafter's attention reads only cells below the one it is writing (n_kv = pos + 1, layer.cpp:813), so every cell past L is written before anything can read it. A failed promote never reaches the refill: the apply pass is the last step of an atomic restore, and the fallback (resume = 0, full re-read) is the path today's `resume > 0` gate already took.
…on_geometry_key verbatim The tier keyed a file on a DERIVED PROJECTION (n_qsa / n_gdn from n_layers + qsa_interval, idx_dim from idx_key_dim, page_size / idx_block from qsa_real_shapes); the shared core keyed the same conversation on geometry_key(g)'s 18 raw ModelGeometry fields. Two keys answer two questions, so the two tiers could refuse different mismatches on the same conversation. There is now one: NvmeHeader::geometry is strata::core::conversation_geometry_key(g), the same 18 fields in the same order, compared as a single array by the dump, the restore and the startup scan. page_size, idx_block and max_cells stay as SEPARATE header fields. They are not model identity and their key does not contain them, but a disk reader must validate them rather than re-derive them from a live engine - re-deriving is exactly what made a sizing change surface as "layout mismatch ... idx_pooled_rows / PLE / drafter ring changed?" instead of as a refused tag. To keep it one key rather than two, conversation_geometry_key left conversation_state.cpp's anonymous namespace and is declared in conversation_snapshot.hpp (the only edit to that imported file; its 18 fields and their order are unchanged, and it gained the conversation_ prefix its neighbours already have). The core's three call sites and SavedConversation::geometry are untouched in meaning. The header is 208 bytes and that is now pinned: static_asserts on sizeof/alignof and on the offsets of L, geometry, page_size and mtp_host. The payload digest deliberately covers only the payload, so the header - the one thing that can move every segment after it - has to be checked at build time. tools/nvme_p0_test.sh and tools/nvme_steps123_test.sh read the ids at a fixed offset and take it from one named variable now instead of hardcoding 104 twice. Format version 3 (same bump as C5/C4; no v3 file has been written yet, so this is one format, not two).
…bid path does NOT mask it The question step 2 left open was whether the stale row at L/idx_block in a turn-boundary snapshot matters, because qsa_select.cu:35 reads `dead` rather than the pool for the block in progress. It matters: that masking exists in one reader only (and there its highest pooled read is n_bid - 1). qsa_index_kernel guards `b > n_bid` and then reads pooled[b * idx_dim] (qsa.cu:253, :260), and native_qsa_score.cu:74 gates `row <= full`. At resume n_bid = n_kv / idx_block = L / idx_block (qsa.hpp:156, layer.cpp:813) - exactly the row the snapshot holds stale, because the pooling kernel rewrote it with the completed block's key (qsa.cu:211) when tokens past the boundary finished that block. So the invariant pooled[n_bid] == dead, which the writers maintain at every block completion (qsa.cu:213, native_qsa_indexer.cu:93), has to be RESTORED rather than assumed. nvme_restore now re-publishes `dead` into row L / idx_block of every QSA layer after its apply pass - the same rule conversation_checkpoint_restore applies (conversation_state.cpp:186), so both tiers leave the indexer in the same state. For a full-L dump the row already equals dead and the copy is a no-op. Row index safety follows from C4: snapshot_pooled_rows guarantees L/idx_block + 1 <= idx_pooled_rows. The step 4 fixture (non-block-aligned boundary, distinct idx_dead) is still owed; this commit is the rule, not its proof, and docs/nvme-kv-cache-convergence.md says so.
The rule was already the rule since 0c7e6f7, but it lived in a doc section rather than in the two places a reader of the tier looks: the API comment on nvme_dump_at (done with C5) now says it is the ONLY path to disk, that it writes the boundary's ids, pictures and running state, all describing position L, and that a parked conversation's consumed ids - which include the model's hidden reasoning tokens - are what a spill may NOT copy; docs/nvme-kv-cache-design.md gains it as correctness rule 4, next to the turn-boundary keying it generalises. The envelope states the same property in its fields rather than in prose: L, the geometry key, the ids, the image records and every running-state segment are the boundary's (C5), and the pooled-row count is a function of L alone (C4). Refuse-never-convert now covers a format version too, so rule 5 says that explicitly. What C1 still owes is recorded as unchanged: whether a RAM-tier eviction also spills through nvme_dump_at. That is a policy question about the RAM cache, which is still dormant here on purpose.
Two header comments still asserted things that are no longer true: conversation_state.cpp said the disk adapter "still uses its own L / idx_block + 2", and conversation_snapshot.cpp said the adapter lays out "their own ... geometry tag" and that step 3 "moves" it. Both formulas are the core's now (snapshot_pooled_rows; NvmeHeader carries conversation_geometry_key verbatim), so the comments say that, and say what is still the adapter's - the envelope itself and, because of it, the drafter-ring refill (C6's collapse condition). Comments only; no behaviour and no format change.
…ns it (C8) state_hash_line() is the tree's only fingerprint and both call sites already used it; its FIELD SET was not the one issue Niko1221#57's main() prints. Adopted theirs: idx_dead hashed per QSA layer and printed, the pooled span widened to the spare row, the stale clamp on the session's max_cells, no PLE bytes hashed when a session has no PLE history, and the fp16 KV branch reads the fp16 pools instead of the int8 array names (which are null in that mode). The widened span is the row C3 re-publishes. A fingerprint that could not see it could not tell a restore that re-publishes the spare row from one that leaves it stale - the oracle would be blind to the fix it exists to prove - and it is the row the file already carries (C4). The count is now one named function, strata::kernels::qsa_pooled_rows, read by the envelope's pooled segment, by the fingerprint, and (unchanged) by the shared core's own sizing.
…stant that pins it tools/nvme_p0_test.sh and tools/nvme_steps123_test.sh read a snapshot's ids at sizeof(NvmeHeader), and both were told the number by hand (104, then HDR=208 when step 3 widened the header). They now source tools/nvme_header_layout.sh, which reads kNvmeHeaderBytes / kNvmeFormatVersion from include/strata/platform/kv_nvme.hpp - the same constants the static_asserts pin - and refuses to run if either cannot be read. nvme_snapshot_ids also checks the VERSION FIELD OF THE FILE against the header's version, so a snapshot written by a different build is caught instead of parsed at the wrong offset (the oracles run /local/strata/build/strata, which is not necessarily the tree they live in). Fixed while doing it: nvme_p0_test.sh's snapshot-ids step passed two arguments to a python reader that indexed argv[3] - an IndexError on every run. It was edited for the v3 header and never run.
conversation_cache_test (35 checks), conversation_memory_test (23), conversation_validation_test (780) and its wrapped variant conversation_transfer_test (1020) are their files, verbatim, ported to this tree. The validation fixture is where their conversation_checked overflow tests live: 'running-state arithmetic overflow rejected', 'zero layer interval rejected before division', 'KV byte overflow rejected'. They compile against our imported core unchanged - the only difference between our copy and theirs is comments plus conversation_geometry_key leaving the anonymous namespace. conversation_memory.cpp is imported for the memory fixture ONLY, not into strata_engine: the RAM-admission guard stays declared-and-unwired in the engine, which is what step 2's 'declared boundary only' was protecting and what C10 will change. New switch STRATA_BUILD_CONVERSATION_TESTS, not STRATA_BUILD_TESTS: the published source leaves tests/ out, so turning the suite on in this worktree fails on the missing subdirectory. The CUDA-linked pair runs with CUDA_VISIBLE_DEVICES=-1 - they link strata_engine and never initialise it, and neither loads a model.
…e it can assert (step 4, parts 2-3) src/platform/kv_nvme_host_test.cpp drives nvme_dump_at / nvme_restore / KvNvmeStore with a synthetic session whose device arrays are host buffers: the target links CUDA but never initialises it (GNU link wrapping on cudaMemcpy / cudaMemcpyAsync / cudaDeviceSynchronize / cudaGetLastError, the shared core's own technique), and the fixture sets CUDA_VISIBLE_DEVICES=-1 itself so it cannot reach the busy GPU. 113 checks, 9 of them asserting a specific refusal. What it pins: * a turn boundary that is NOT idx_block-aligned (10 tokens, idx_block 4): the key is the boundary's 10 ids, not the consumed 26, and the entry's pictures are the boundary's, not the live conversation's; * a picture at or past the boundary is REFUSED, naming its start; the boundary's filtered list is accepted; * the next request re-sends the prompt without the reasoning tokens and kv_nvme_match PROMOTES the boundary snapshot - and the same session dumped WITHOUT a boundary is unreachable by that request (the control that makes the promote mean something); * idx_dead differs from every pooled row, the file carries the STALE spare row a boundary dump produces, and after the restore pooled[L/idx_block] == dead while the completed rows are untouched and row L/idx_block+1 was never written (C3, C4); * block_pos and the tails are the boundary checkpoint's, not the live arrays' (C5); the drafter's ring slots hold the window the adapter restored, ceil(min(L, max_cells)/page), not ceil(L/page) (C6); * version 2, a geometry-key field, a flipped payload byte, a truncated file and a too-small pooled array each refuse with their own message, and the ids sit exactly at kNvmeHeaderBytes. kv_nvme_match is the serve loop's promote rule moved into the tier's own header (templated on the token type like the shared core's conversation_prefix), so the fixture asserts the code the server runs rather than a copy of it. Behaviour is unchanged; a malformed entry whose ids are shorter than its L is now skipped instead of read past the end. Four mutations were run to check the fixture is not decorative: dropping the spare-row re-publish, dropping the image filter, widening qsa_pooled_rows to +2, and keying the dump on the consumed ids. Each one fails it.
The convergence record gets its step-4 section (the four commits, the fixture step 3 owed, how a fixture drives the tier with CUDA linked and never initialised, the checks that ran, the ones skipped and why, and the list of what stays unverified). C3 and C1's fixture bullets now point at the proof instead of promising it. docs/nvme-kv-cache-design.md gains the spare-pooled-row rule in section 4, the host fixture in its test record, and an explicit note that its P0 / Steps 1-3 results are the v2-format record: neither shell oracle has been run since step 3 widened the header.
…ture proves it (C7, part 1) `nvme_restore` and `KvNvmeStore::restore` now return the shared core's `ConversationRestore` (`restored` / `invalid` / `transfer_failed`) instead of a `bool`: one vocabulary for the RAM tier and the disk tier, which is what "one failure contract" has to be spelled in. The classification is not a guess - it follows from where in the function the failure happened, and the fixture now measures that. Tier side: - every pre-apply refusal (magic, format version, geometry, header sizes, segment walk, layout end, payload digest, a live array too small, a null target buffer) returns `invalid`, and makes NO CUDA call at all; - the apply pass now BEGINS with a `cudaDeviceSynchronize()`, as `conversation_snapshot_restore` does (conversation_state.cpp:258): a device that is already unusable fails before a byte is written; - a `cudaMemcpy` in the apply pass or in the spare-row re-publish, and either sync, return `transfer_failed`, naming WHICH segment and how many bytes. The old message was the anonymous "H2D"; an operator could not tell a corrupt indexer segment from a corrupt GDN state; - every one of those paths CONSUMES the CUDA error it handled. `cudaGetLastError()` returns the last error and clears it; left unread it is sticky, and the next `cudaGetLastError()` on the serve loop's drop-and-restart path is `kv_stream_reset`'s `check()` (kv_stream.cu:199-202), which exits the process - so today a failed promote aborts the re-prefill with a message naming the reset. `pinned.cu:170-183` is this tree having already paid for exactly that trap with a `cudaHostRegister` error that made an unrelated kernel launch report "out of memory". Consuming it does not make a transfer failure recoverable - a sticky context error comes straight back - it makes the report true. - `KvNvmeStore::restore`'s own case (the file applied cleanly but no longer matches its index) is `invalid`: the restore's final sync succeeded, which is the proof a recovery needs, and the session holds a complete digest-verified snapshot rather than a half-applied one. The fixture (113 -> 247 checks) models the CUDA last-error state instead of suppressing it, and can now inject a failure at a chosen copy or sync: - 9 pre-apply refusals: `invalid`, zero cudaMemcpy, zero CUDA calls, every session buffer still at a poison sentinel, no prefix handed to the caller, no error left pending; - 6 transfer failures: `transfer_failed`, no error left pending, and the session PROBABLY half-written - fail at the 2nd device copy and the GDN state is the snapshot's while the PLE history is not; fail at the last one and every segment landed while `pooled[L/idx_block] != dead`, i.e. a session that looks restored and is one invariant short. That is the same shape the shared core's own fixture asserts for its restore (conversation_validation_test.cpp:167-169); - a store reports the same classes, and the stale-index case is asserted to be fully applied (12 copies, 2 syncs, the new snapshot's last token in the PLE window); - a failed dump leaves no entry, no partial file, and the session byte-identical - the dump side needs no contract because it only reads; - a store of version-2 files: `open()` succeeds, 0 entries, 0 bytes against the cap, the operator is told the version found and what to do, the files stay on disk, and the request that WOULD have promoted from those same files at version 3 gets nothing - a re-prefill. Mutation-tested: dropping `consume_cuda_error()` -> "left a CUDA error pending (invalid argument) for the next caller to misread"; dropping the pre-apply sync -> the pre-apply case reports the POST-apply message; classifying a transfer failure as `invalid` -> "the first copy failing is a transfer failure" fails.
…ills (C7, part 2) The serve loop now branches on the class the tier reports, so one contract covers both tiers: - `transfer_failed` -> `ERR …` + `return 1`, the same consequence the RAM tier already gives a failed checkpoint restore (generate.cpp:3010-3012) and the same one issue Niko1221#57's core calls fatal. No drop (a transfer failure says nothing about the file), no re-prefill: the apply pass is a loop, so the session is half-applied by construction, and nothing has shown the CUDA context still answers. `serve/server.py:686-695` turns that exit into an actual recovery - the next request restarts the engine, and a new process is a new CUDA context, a new pinned arena and new graphs. - `invalid` -> unchanged behaviour (drop the entry, `resume = 0`, full re-read), now justified by evidence rather than by hope: the fixture asserts each of those refusals makes zero CUDA calls and writes nothing, so the clean path has nothing to undo. - The startup `--nvme-restore` message names which class it was ("transfer" vs "refused"): same consequence, different fix - one is the file, the other is the machine. The stale-store scan message now states the operator's two options instead of one: re-dump the version-2 snapshots with the binary that wrote them, or remove them and let the store rebuild. It also says the consequence out loud - they are skipped, nothing promotes from them, every request re-prefills. A store full of unreadable files is a RECOVERABLE condition reached before any transfer, and it is not swept into the fatal rule.
… reset would have to prove (C7, part 3)
docs/nvme-kv-cache-design.md gains §5 "The failure contract (C7)" (Test record / Known limitations / Prior art
renumbered to §6 / §7 / §8):
- §5.0 states the finding plainly, not softened: the clean-reset fallback this document has described since §3
has never run. A failed promote died inside the fallback's own first step - session_zero -> qsa_state_zero ->
kv_stream_reset (layer.cpp:676) -> check("reset") (kv_stream.cu:199-202) -> std::exit(1) - printing the
RESTORE's error under the label "reset". Neither our claimed recovery nor their fatal rule: an unintended exit
at an unrelated point with a misattributed diagnosis, and it survived every recorded test because every one of
them corrupts a FILE (a refusal, which never reaches the transfer pass) and none injects a failing COPY.
- §5.1 `invalid` - recoverable, with the evidence that makes it safe (zero CUDA calls, nothing written) and the
exact operator line.
- §5.2 `transfer_failed` - fatal, why the apply pass makes it so, the exact operator lines, why the snapshot is
NOT dropped, and why stopping is recovery rather than defeat: serve/server.py:686-695 restarts the engine, and a
new process is a new CUDA context, a new pinned arena and new graphs.
- §5.3 a stale-format store - recoverable, not swept into the fatal rule, with both operator actions named:
re-dump the v2 snapshots with the binary that wrote them, or remove them and let the store rebuild.
- §5.4 what a future clean reset must prove before it may exist: on device, that a real failed H2D cudaMemcpy
leaves a non-sticky context error; that the captured graphs survive replay after one; that the pinned arena and
page table are back in the residency contract. Names the device-usability probe as the mechanism and marks it
unvalidated on this machine - deliberately not implemented.
- §3's Promote bullet no longer claims the fallback works, §6's test record describes the 249 checks and what
they do not cover, and §7 says out loud that a transfer failure ends the request and the process.
docs/nvme-kv-cache-convergence.md: the C7 section now carries the settled contract as a three-row table with the
operator-visible messages, and a "Where step 5 stands" section records the commits, what the fixture proves, the
three mutations, what ran, and what the GPU still owes - including that the P0 corruption negative control now has
a CLASS to check: a corrupt file must report `nvme promote refused`, not `FAILED (transfer)`.
Code comments updated to point at §5 (they named a section number that did not exist).
…hot is Its corruption control already greps for "integrity check failed", but the two failure classes now have different consequences - one re-prefills, the other stops the engine - so the oracle has to tell them apart: a corrupt FILE must report `nvme_restore failed (refused)` and must never report `(transfer)`. Edited, not run: the script needs the model and the GPU, which hold a live ~24 GB engine. `bash -n` clean.
…ay named apart A fourth tab that exists only when the tier does: the button starts hidden and /metrics' cache.enabled unhides it without a reload, so a server whose engine has no --kv-nvme renders today's three tabs. The cards, gauge and totals render from the 1 s /metrics block (summary(), no filesystem work) and /cache is fetched on its own 2 s timer that starts when the tab is shown and stops when it leaves - an unwatched Cache page costs no directory walk, which is the whole point of the summary()/detail() split. "Store used / 100 GB cap" carries a tooltip saying it is the engine's cap accounting (a chunk shared by three manifests counts once per manifest, on purpose); "On disk - N files · X GB" is the walk's footprint. Equalising them is the mistake this page exists to avoid, so they are two numbers with two names. The events table is design §7's three failure classes as badge tones, with the honest shortenings the payload forces: the KV line carries no <reason> for a refusal, and §7 itself forbids promising a restart the page cannot observe, so the transfer row stops at "engine stopped". Review fixes: the events header said "ms" while the new dur() prints "1.84 s" (now "Duration"); the prefixes table's "Last used" column had no field behind it and §8 forbids inventing one from an mtime, so the column is gone rather than rendered as a dash; and "Show all" capped at 12 rows like the summary it replaced - it now shows every row /cache kept. Tests 68 -> 70 (the markup, the data-tab="cache" button, the #i-cache glyph in page and sprite, the app.js render path, and /web/cache, /cache/app.js, /cache/../web/app.js all still 404). node --check passes; a mock-engine server starts. Owed: the manual both-themes / narrow-breakpoint pass the plan asks for - the CSS reuses only existing tokens and collapses at the Monitor's own 1000 px and 640 px breakpoints.
…eb plan step 7) The Monitor's request table gains the highest value-per-line change in the plan: `Reused` cannot tell a 4,107-token RAM resume from a 4,107-token promote, and the KV line can. RAM / NVMe / disk / cold, with a tooltip carrying the tokens, the duration and the bytes staged in RAM; a dash - not a badge - when the request reported no KV line, because unknown is not cold (§3 note 1). The Cache tab's events table keeps naming the tier (v3 / delta); this column answers a different question: did this turn read disk. About gains the card a user reads the caveats from, built only from the /metrics cache block - About never fetches /cache and never walks the store. Two plan-listed rows are deliberately absent rather than invented: the stale-file count is a walk field (it is on the Cache tab), and a weight fingerprint is not a field of anything the serve side has. telemetry samples the engine process's own resident bytes, which is how the ~2 GB a promote stages becomes a bump on the Monitor instead of a footnote. `pid` is resolved every sample, and StrataEngine.pid is a property over self.proc, not a number captured at startup: restart() replaces the process and a transfer failure is exactly the path that does it (§5.2). No psutil, no engine, or a process that is gone -> the key is absent, never 0; the sparkline draws nothing rather than a flat zero. Tests 70 -> 79 (EngineRss x5, the pid-across-restart case, the Cache column's data and markup, the About card, and an assertion that the Reused column is byte-identical). Verified against a live process: rss_used 33,452,032 with a history of readings; absent for pid=None and for a dead pid.
… (web plan step 8) setup.py offers the NVMe cache (store dir <root>/kvstore, cap 100 GB = the engine's own default, delta tier on) instead of leaving it to a hand-edited args list, and it refuses to offer it where it would be a lie: under a layer split (a snapshot would be written and never read), where it would force back on the KV streaming setup just refused (WSL's pinned-RAM limit, or a RAM copy that does not fit), and where the cap is more than the volume holds - there the cap is clamped and the clamp is printed. Answering "no" leaves args exactly as it was; under --yes the recommended answer is taken without a prompt and without holding the install up. C11: "NOT settled" becomes "settled as a build", naming what exists (TierActivity, the KV line, kvcache.py behind /cache and /metrics, the Cache tab) and what it deliberately is not (no store mutation - /cache is a GET, the only POST paths are /settings and the two chat routes, and the engine's stdin takes QUIT/STOP/GEN/ GENI only; no conversation identity across a restart). §6 records the page's metric definitions so a number on the page is a number an oracle already measures - above all that entries_bytes/delta_bytes are cap accounting (a chunk shared by three manifests counts four times, on purpose) while the walk reports the disk footprint, and that endurance is §10.2's ~6.5x, not the handoff's 33x. Recorded honestly as unrun: the KV-line assertions in nvme_steps123_test.sh and nvme_failure_contract_test .sh are written but never executed - both need a GPU and this machine's GPU holds the production engine. And why short_tests.py can never gate on the line: the server gives the engine's stderr to the log and keeps stdout on a pipe (server.py:159-160), so a log tail never sees a KV line. setup.py's install path was never executed; /local/strata was read only.
Found by RUNNING them on silicon (they had not been executed since the delta tier became the default cascade - kv_delta = 1 at generate.cpp:298, which predates this branch; verified at 7c779de). With --kv-nvme alone a turn appends chunks + states + a manifest and writes NO v3 snapshot, so: - nvme_steps123_test.sh globbed $STORE/kv-*.bin, found nothing, and its python helper died on the empty path before writing full.txt; Step 2 then reported RESUME 0 for a store it believed was populated - a silent pass on a store that never had a snapshot in it. - nvme_failure_contract_test.sh copies a snapshot to $STORE/kv-1-1.bin and corrupts it: with no snapshot there is nothing to refuse and no transfer to hook. - nvme_delta_p0_test.sh's runv3 has said "the V3 cascade (delta off)" in its comment the whole time and never passed the flag, so the §5.2 comparison's other half ($OUT/store-v3/kv-*.bin) was empty, V3_SNAP was "", and BV died on `--nvme-restore ""` - "needs a value", rc=2. Its own rund passes --kv-delta 1 explicitly, which is the intent written down twice and the omission once. Each now names the tier it asserts on with --kv-delta 0; the delta path stays gated by rund's explicit --kv-delta 1 and by short_tests' --kv-delta 1 run. Also: steps123's step3 KV PASS message printed `head -1` of every evict= match, so it read "evict=0 evict_bytes=0" while asserting evict>=1 - the assertion was right and the record was wrong; it now prints the line that satisfied it.
…(found on a live server)
A tier-on server reported `enabled: true` with every store fact null - promotable.bytes, entries,
delta_entries, cap, the RAM tier - until some request's line happened to arrive. The cause is ordering,
and no unit test could see it: main() builds the KvCache AFTER the engine is constructed, the engine prints
its `KV start=1` state just after READY, and the pump thread had often consumed it by then with
`self.cache` still None. Every test attached the cache by hand, before the line, so every test passed.
attach_cache() now does what _kv_line could not: it takes the cache and replays `store_kv`, which the pump
fills whether or not a cache exists yet. Either path delivers the state exactly once - replay when the
line is already read, direct when it is still to come - so there is no double count and no window where the
page knows the tier is on and nothing else about it.
Proven on the live server, not only in the test: after the fix, /metrics with ZERO requests carries
`promotable {prefixes: 24, entries: 5, delta_entries: 19, bytes: 6,170,642,908, cap: 200 GB}`, while
`ram_tier` stays honestly null - checkpoints and live tokens ride the request line, not the store line.
Tests 79 -> 80 (KvPump: the startup line read with no cache, then attach_cache, then the facts present;
attach_cache(None) is not an error).
The web plan's eight steps are shipped, so §6 stops describing the live `KV` gates as written-and-unrun and records what they said: steps123 3 KV lines per 3 DONE with one start=1 each and the cap turn's `evict=1 evict_bytes=178042216`; the failure contract's promote-path line carrying `transfer=1 total_transfer=1` on stdout line 4 with the ERR line at 5, on all three CUDA hooks; nvme_p0 `RESUME 4104` with the post-restore hash equal to the dumper's DONE hash; delta_p0 14 checks with the two tiers keying the same 2355 tokens and identical restored-prefix hashes. Recorded as findings, not as footnotes: - three v3 oracles had been asserting a tier they never asked for - the delta tier is the default cascade, so no snapshot existed to supersede, to corrupt, or to restore; the worst one reported `RESUME 0` off a store it believed was populated. A gate that cannot find its fixture must fail loudly. - the page against a real engine, which no host test can substitute for: a restart-then-promote seen through /metrics (`src=delta resume=1453 promote_ms=581 staging_bytes=140293564`, 1751 ms -> 652 ms), and §5.1's recoverable class arriving as `refuse` + `sweep` rows because a chunk short_tests corrupted is SHARED into another conversation's manifest - which is the cap accounting counting it once per manifest, explained. - the two books on that live store are 1.89x apart (6,170,642,908 vs 3,267,409,228 over 3,650 files). - needle_bench 3/3 at 32K with the tier on, and honest about why: all three are `src=none` - three depths are three different prompts - and the speed-up those runs saw was the RAM prompt cache. A page that showed only wall-clock would have credited the disk with a saving it did not make. Also fixed the plan's own test matrix, which named short_tests.py as the KV gate: it tails the engine's log, which carries stderr, and the line is on stdout.
…ing plus the recompute's own test 6648be7 made the sweep recompute the tier's total from the disk, so delta_bytes is no longer a monotone over-count: it drifts above the footprint as shared references accumulate (the safety direction - over-evict, never under-evict) and snaps back at every sweep. The three places that still said "the cap cannot lie about the disk" as if the over-count were permanent now say what the number actually is (serve/web/app.js:137 tooltip, serve/kvcache.py _promotable_locked docstring, design doc §6), including that the two books agree exactly only right after a sweep - a coincidence the page must not present as an invariant. New host fixture sweepbooks: after a sweep, assert total_bytes() equals a recursive disk walk of the delta subtree - the recompute's promise, checked against the disk itself. kv_delta_host_test green (110471 checks, fuzz seeds random per run).
…d 644/1176 figures retired) Outage-1 redeploy verification, 2026-09-30, merged binary + serve tree: - four GPU oracles ALL PASS (steps123, failure_contract, p0, delta_p0) with NVME_ENGINE=build/strata; short_tests ALL PASS; needle_bench 3/3 at 32k. - zero-request /metrics: enabled=true, delta_entries=63, delta_bytes 21198619632, ram_tier nulls (the start=1 replay, bf69e20, proven live). - one conversation lands a KV line (cascade event, 118.9 MB). - delta_p0 fork shape under grouping: 10 chunk FILES for 28 references. - the two books re-measured live: delta_bytes 25,008,877,848 == the walk's delta subtree, byte-exact - since 6648be7 the sweep runs after every turn's dump (kv_delta_enforce_cap sweeps unconditionally), so the live sawtooth sits at its snap-back point; the pre-merge 1.89x was the old binary's drift. Kept as history, not as a current claim. - promote re-verified live: src=delta resume=3018 promote_ms=560 promote_bytes=164173592 staging_bytes=164172720 (restore path untouched). - web-design: the fact table now states block AND chunk span (64 blocks, 256 tokens, per-manifest blocks_per_chunk, 0 = legacy); the cap-accounting section carries the sawtooth wording and the merged-tree line numbers.
…RE_TIMING=1) One stderr line per delta_restore under the env flag: manifest scan / chunk read+digest / assemble / apply, plus the process's VmHWM beside the RSS held at entry (so the promote's OWN transient is the difference), with T, chunk count and read-back volume for context. Unset - the default - nothing is printed; the flag is read live like STRATA_DELTA_FAIL_AT. The RSS parse uses fgets/sscanf per line: a bare fscanf over /proc/self/status desyncs on the non-numeric lines and silently reads nothing. Before, measured on-device on the merged binary (143,495-token conversation, 560 chunks, 2.15 GiB read, server stopped, gen-driven): manifest 0.5 ms, read+digest 3115.3 ms, assemble 2378.8 ms, apply 1887.8 ms, rss peak 50606 MB (entry 46202 MB) - ~7.4 s wall, ~4.4 GB own transient. Recorded into the restore handoff §1 as the A/B baseline. Host fixture green with the flag on (110410 checks).
…kers
The restore keeps its shape (validate -> read chunks -> assemble ->
nvme_restore_image) but the middle stages run on a small worker pool
(std::thread + an atomic index, no new dependency): each worker takes the
next chunk, does its open/read + FNV digest + slice-math size check, and
places that chunk's slices into the image with memcpys. The walk that used to
copy the chunk bytes now RECORDS each slice's destination - same walk, same
running offset, so the destinations are byte-for-byte the ones the sequential
assembly produced - and the workers' writes are disjoint by construction.
- the workers touch NO CUDA (R2): host reads, digests and memcpys only, and
nvme_restore_image still runs on the caller's thread after the join;
- a worker keeps ONE payload buffer and reuses it across chunks (the chunk
staging is gone: no vector-of-vectors of payloads);
- refusals stay deterministic: each chunk's error is collected, the pool is
joined, and the FIRST failure in CHUNK ORDER (lowest index) is reported as
invalid - same strings delta_read_chunk produces, the workers call it
itself;
- image_bytes (staging_bytes) now reports only once every refusal is past, so
a refused restore still stages NOTHING - the fixture asserts this and caught
the one-commit-late placement;
- the State record is read on the caller's thread first (it is the assembly's
non-chunk half), validation unchanged everywhere.
Gates: ctest 7/7 (kv_delta_host_test 108930 checks, incl. every Phase-3
refusal case and the fault matrix), nvme_delta_p0_test.sh ALL PASS on device
with NVME_ENGINE=build/strata. Evidence, 143,495-token promote on device:
before: manifest 0.5, read+digest 3115.3, assemble 2378.8, apply 1887.8 ms,
rss peak 50606 MB (entry 46202) = 4.4 GB own transient
after: manifest 0.5, read+digest 3021.3, assemble 1713.2, apply 1811.1 ms,
rss peak 48597 MB (entry 46265) = 2.3 GB own transient
read+assemble 5494 -> 4734 ms; the transient is halved because the chunk
staging is gone. The read itself is disk-bound (~0.7 GB/s over 560 small
files) - which is also what the v3 tier pays for its sequential 2.2 GB, and
the target B reaches.
…r-line)
R1's one permitted extraction: nvme_restore_image's tail, from 'the apply
pass begins with a sync' through the final synchronize, becomes a second entry
point that takes the applies list and the layout/prefix facts as parameters.
The body did not move in the file - only the function boundary did, so the
diff is the boundary plus three parameter substitutions (h.n_imgs -> n_imgs
twice, h.cvec != 0 -> cvec_flag). Everything else is the same lines:
- the STRATA_TEST_FAIL_CUDA hook stayed inside the extracted body, so the
fault matrix keeps its names, its counts and its positions;
- nvme_restore_image calls it with the applies its own walk recorded.
Two types moved to the header so a caller can name them: Sizes (out of the
file's anonymous namespace - linkage only, no field or use changed) and the
apply-segment record, now NvmeRestoreApply (the file's local Apply is an alias
of it), so the delta tier's streaming restore (B2) can build its own list in
the v3 walk's order and names.
Proof the extraction moved nothing: the fixtures are green UNCHANGED -
ctest 7/7, kv_nvme_host_test 286 checks (18 refusals, 6 transfer failures
asserted - the fault matrix runs against the extracted body), kv_delta_host_test
green including the nvme_restore(file) == delta_restore(store) cross-check.
No assembled image any more. The restore validates, reads the State record, then each worker reads one chunk and PLACES it: the KV pages straight into st.host.* at their page offset, the drafter's pages into its host arrays, and the pooled rows into a small per-layer staging buffer - the only bytes that cannot land in one copy, because they live on the device and each chunk holds only its own span of rows. The staging (~20 MB at 142k) plus the State record plus one payload buffer per worker IS the transient; the 2.2 GB image, the second hash pass and the extra full copy are gone. - THE APPLY PASS is B1's nvme_restore_apply with OUR applies list, in the v3 walk's order and under its names: gdn, ple, then pooled/tail/dead/block_pos per layer (R4). The fixture pins the streamed restore at exactly 2 + n_layers*5 device copies - the v3 fixture's number - and checks the fault hook still skips them BY NAME (pooled rows, spare-row re-publish). - VALIDATION IS A STRICT SUPERSET of what the image walk refused: the manifest's header/count facts, the LATTICE (one chunk per sealed span - a short ref list would have left blessed holes in the image), the pooled-rows capacity refusal (kv_nvme.cpp's snapshot_pooled_rows rule, which streaming would otherwise have written past), every chunk's existence/size/key/range/digest in the workers, and the State's sizes in step 2 - all before the first CUDA call. - THE REFUSAL UNDO: streaming wrote verified chunks into the host arrays before it could know about a later chunk's refusal, so the refusal path re-runs the same placement as memset over exactly those ranges (and the State's) - a fresh session comes back all-zero, which is what the contract means and what the fixture asserts. - Two layout traps from the handoff kept: separate per-array strides for the chunk side (span/page pages) and the State side (tail pages), and pooled rows per SPAN, not per block (the fixture caught this one: 64x off under grouping). - staging_bytes now reports what the path really allocates - the State record + the pooled staging + one buffer per worker - not a whole snapshot. Evidence, 143,495-token promote on device (560 chunks, 2.15 GiB read): before 3115.3 read+digest / 2378.8 assemble / 1887.8 apply ms, 4.4 GB transient after 2637.7 read+digest+place / 0.0 stage / 24.2 apply ms, ~340 MB transient = 2.7 s wall, ~53k tok/s - at or better than the v3 tier's ~50k. BIT-EXACT at this scale: the v3 tier's own snapshot of the same boundary restored with --nvme-restore prints the same STATE_HASH line (all nine fields). Gates: ctest 7/7 (kv_delta_host_test green incl. the cross-check, every refusal case, the fault matrix), nvme_delta_p0_test.sh ALL PASS with P1 bit-exactness, serve.test_server + test_mcp 104 green, node --check. The page's staging claims follow: 'Promote staging', the About card's 'Promote cost' row, and the C10 record are now mode-dependent and measured.
…target we did NOT hit
B3 evidence on the deployed binary (build/strata == engine/strata, md5
c83b790e2ad8faca0ce1e22ecb40f497), server stopped per §9 for the oracles:
- nvme_steps123, nvme_failure_contract, nvme_p0, nvme_delta_p0: ALL PASS,
delta_p0 including P1 bit-exactness.
- 143,495-token promote, on device, from the KV line itself:
KV src=delta resume=143495 promote_ms=2668 promote_bytes=2309000692
staging_bytes=356179248 (2.7 s, ~54k tok/s)
- zero-request /metrics after redeploy: enabled=true, delta_entries=85,
ram_tier nulls; needle_bench 3/3 at 32k on the deployed binary.
- short_tests: ALL PASS against a FRESH store with this binary (20/20).
ON THE RSS TARGET, stated plainly because it is not what the handoff guessed:
staging_bytes is 356,179,248 at 143k tokens, not ~100 MB - 115 MiB State
record + ~209 MiB pooled-row staging + 16 MiB of worker buffers. The process's
own transient peak is ~340 MB (VmHWM 46,597 vs 46,257 at entry) against the
assembled path's ~4.4 GB, so the 13x reduction and the bounded, chunk-count-flat
staging are real, but the number is O(conversation) in the pooled rows: R4's
one pooled apply per layer needs each layer's whole row span in host memory at
apply time. Cutting that would mean per-chunk pooled applies (560 x n_layers H2D
copies), which changes the applies list's shape and count - the owner's call,
recorded as the next item rather than done quietly.
Also recorded: short_tests.py's S3 section flips a byte in every chunk and state
of the store it is pointed at and does not repair them, so a second run against
a store a previous run poisoned reports S1/S2 prefill failures that are the
TOOL's negative control working as documented, not a regression - proven by the
same binary scoring ALL PASS on a fresh store.
`docs/nvme-kv-web-merge-handoff.md`: what merging nvme-kv-web into nvme-kv-cache actually involves, for whoever does it from the other side. Two conflicted files (one hunk in kv_delta.cpp, five in the delta host test) resolved from a probe of the merge tree - no working change - so the hunk line numbers and both sides' text are quoted rather than guessed. The .cpp answer is "keep both": upstream's total_ = kept recompute and my act.swept reporting, dropping my total_ -= bytes because that subtraction IS the double-count 6648be7 fixed. The host-test answer is "upstream's boundaries and counts, my reporting assertions re-derived on top" - and the assertions that survive are the ones that compare against the disk walk instead of a literal. Two sections exist because the merge is not only textual. 6648be7 makes delta_bytes a SAWTOOTH (per-reference while records are appended, snapped back to the disk by every sweep), so the page's tooltip, kvcache.py's docstring and the tier doc's metric definitions all need rewording - the two books stay two books, but they now agree exactly when a sweep has just run, which the page must not present as an invariant. And c70c119's 64-block chunk grouping invalidates every file-count figure this branch measured (3,607 chunks, 644 for 1,176 refs, the 1.89x divergence), so those are labelled pre-merge and listed for re-measurement; the cap will also evict more aggressively once its books stop drifting low, which is 6648be7 working and must not be attributed to TierActivity. Also, in the handoff's last section and fixed in the web design's §5 mock-up: the Store block said "Chunk BLOCK lcm(page_size, idx_block) = 1,024 tokens". It is 4 - the live engine shows it three ways over (7,590 chunks for T=30363, 363 for T=1453, and both oracles' sealed() values) - and the number was never measured because the page never rendered it. A mock-up figure becomes a documentation figure becomes a belief; if the page grows a BLOCK row, it must come from the engine, not from a literal.
… scaling with the conversation The restore staged every layer's WHOLE pooled-row span in host memory so the apply pass could carry one 'pooled' apply per layer, exactly the v3 walk's shape. That made the staging O(conversation): 209 MiB of indexer rows at 143k tokens, and the total 356 MB. No ordering of the applies avoids it - the apply pass is a BATCH, so every source it copies from is resident in host memory when it starts, and 'one pooled apply per layer' means the whole row span has to be. So the rows no longer go through the applies at all. After the fan-out (and only after it: every chunk has been validated, so nothing can refuse from here on), the caller's thread synchronizes - the same proof the frozen pass makes before its first write - and then, chunk by chunk, reads that chunk's pooled slice per layer and copies it straight into idx_pooled at the chunk's row offset, with one buffer reused for the next chunk. The bytes and their destinations are exactly what the staged shape produced; the applies list keeps the v3 walk's order and names for what it still carries (gdn, ple, then tail/dead/block_pos per layer), and the pass's spare-row re-publish still runs after the rows are in place. Measured on a 143,495-token promote on device (560 chunks, 2.15 GiB read): before staging_bytes 356,179,248 promote_ms 2668 transient ~340 MB after staging_bytes 135,802,160 promote_ms 2621 transient ~130 MB and the STATE_HASH is bit-identical to the v3 tier's restore of the same boundary, as it was before this change. What is left is the floor: the timing line now prints the State record's own breakdown, and it is 112 MiB of gdn+ple - the v3 walk's first two device segments, which must be host-resident when that batch runs - plus ~15 MiB of reader buffers. The host-bound slices (KV and drafter tail pages) are 1.9 MiB and are consumed before the apply. So the ~100 MB the handoff guessed is below what this geometry's gdn segment alone costs (117,669,888 bytes at the 262,144-token context); below that line, only the reader buffers are ours to trim. Gates: ctest 7/7 (kv_delta_host_test 109707 checks, the cross-check, every refusal case and the fault matrix - now counting the pooled copies and still skipping by name), nvme_delta_p0_test.sh ALL PASS with P1 bit-exactness.
docs/nvme-restore-perf-evidence.md keeps the verbatim output of every measurement behind the docs - the four phase lines (7.4 s / 4,404 MB -> 2.6 s / 130 MB), the KV line with promote_ms=2621 and staging_bytes=135,802,160, the byte-identical STATE_HASH pair (v3 tier's restore vs the streamed delta restore of the same 143,495-token boundary, all nine fields), the State record's own composition, and the gate table - so none of these numbers has to be taken on a doc's word. The handoff, the tier design doc and the page's 'Promote cost' row now carry the measured numbers and the floor: 135,802,160 bytes of staging is 112.2 MiB of gdn+ple (the frozen apply pass's own first two device segments) plus reader buffers, so the restore no longer scales with the conversation - the 219 MB of pooled rows that made it O(prompt) cross the bus per chunk now.
…e (main() closed early) The integration rebase onto 0.1.28 (bbaaabb) left an extra closing brace at the end of the "--kv-nvme needs the KV host copy" block, so main() ended there and the ~206 errors that followed ("expected unqualified-id before 'if'") were all one defect. The file is brace-balanced again, as it is in cfbbb06, in a7578f9 (this series' root) and in the a790805 trial. Both guards the integration is meant to keep are untouched and still present here: the k8v4 --kv-resident refusal above, and the --kv-nvme streamed-floor force that follows.
The image-promote oracle died on "the store wrote no snapshot" with a run that had answered the image request perfectly (rc=0, 39 tokens). The store was NOT empty - it held a 118 MB delta state and a manifest - but this oracle globs $STORE/*.bin, the v3 cascade's layout, while the run used --kv-nvme alone, and the delta tier is the DEFAULT cascade since Phase 6 (kv_delta = 1, generate.cpp:307, commit 0924c1c). So it was reading the delta store's file format as if it were an NvmeHeader: the state file's magic is 0x54415453, not 0x5E564D45, and its offset 8 is a byte count, not L. A fixed glob would not have helped; only the tier flag does. This is d2cb8f6's bug again, one oracle later: that commit named the tier on the steps123, failure contract and delta p0 oracles and left this one behind, so it alone regressed. With --kv-delta 0 the oracle passes on the integrated binary at exactly the numbers the 09-29 record carries: L=19 of 22 ids, n_imgs=1, the run-2 promote resumed 19 of a 73-token conversation, and the different-image negative control still does not match. The image contract is not weakened by naming the tier: NvmeEntry::imgs, the two historical defects (the never-filled list, the dump of the LIVE list) and the ImgKey the negative control probes all live on the v3 path this oracle keeps asserting, and the promote, the boundary key and the n_imgs == 1 check are unchanged. The delta cascade's own write path stays gated by tools/nvme_delta_p0_test.sh, which passes on the same binary.
…o the KV format
The v3 header carried the shared core's 18-field geometry key and the KV format, and that was
enough to keep a store from promoting a file the engine could not read - except that none of it
was the question a same-geometry file actually raises. Geometry is a property of the MODEL, so two
weight sets of one architecture share it exactly: the same key, the same segment layout, the same
prompt. A whole-session snapshot of a DIFFERENT model's conversation therefore had everything it
needed to match, verified its integrity digest (which proves only that the file was not corrupted on
its way to disk), and applied cleanly into a session that had never seen those weights. Nothing in
the file could have caught it: the KV bytes are just bytes, and no arithmetic over them tells one
network's attention state from another's. Only the writer knows, so the header has to say.
v4 adds the two fields that make it say so, and nothing else about the envelope's layout changes:
weights_fp the weight-set fingerprint of §5.8 - FNV-1a over each shard's resolved path, its
size and its first/last 64 KiB, read once per process when the store opens. It is
the SAME computation the delta tier keys its manifests with, and it now lives in one
place: nvme_weights_fp (kv_nvme.cpp) holds §5.8's formula and kv_delta_weights_fp is
a forwarder to it. Two writers of KV bytes into one store directory must refuse on
one definition of "these weights", and a second copy of a fingerprint is a second
chance to drift.
kv_format filled from qsa_kv_key instead of qsa_kv_format. The split is the k8v4 part: a
KEY has to name every format the engine can hold, and qsa_kv_format answers a
different question - "which layout may the block movers walk" - and refuses a hybrid
state outright, because no (pointer, length) pair of runs can move a K-INT8/V-Q4_0
pool. Keying storage with the refusing accessor would have meant either crashing on
a hybrid session or filing k8v4 as fp16: 816 B/cell read as 1,056. kKvK8V4 is now
in the enum and the key is total over all four formats; the movers still never see
it, which is what keeps that refusal honest.
Both mismatches are `invalid`, decided before the first CUDA call, so nothing is written and the
caller re-prefills. Neither is converted, and there is no repair path - the bytes are wrong in a
way no arithmetic can fix. The scan checks weights_fp first, so a foreign snapshot never becomes an
entry, never reaches the match and never holds a byte against the cap, and it gets its own line:
a store full of files that pass every other check and are still not promoted is a case none of the
other scan messages explains. The store's key decides what becomes an entry; the live engine's key
decides what may be applied - so pointing an engine at a store written under a different --kv
promotes nothing instead of reading a file at the wrong stride.
A v3 file is refused BY VERSION, never re-read as "fingerprint 0" - which is precisely the value a
foreign store would otherwise have been accepted on. The cost is one re-prefill per conversation
after the upgrade; the files stay on disk and §5.3's operator actions are unchanged. The header
grows 208 -> 216, moving every field after kv_format by 8.
The two static_asserts under the struct were restated field by field rather than loosened into a
range - the pins exist to catch a field gaining padding, and a permissive assert stops catching
exactly that. They fired on the first build here, which is the discipline working.
The host fixture gains a weight set of its own (a real file, because that is what an engine
fingerprints) and the refusals v4 is about: same geometry + different weights is not promoted and
leaves nothing written; k8v4 against int8 is refused in both directions, including the case the
split exists for, where a store opened as k8v4 keeps the file but an int8 engine still refuses to
apply it. 329 checks, 21 refusals (was 286/18). The delta fixture's byte-identity oracles write
their v3 controls at the same fingerprint the delta records are read with, since the header is
part of the bytes they compare - its count is unchanged at ~110k, though it is seeded and varies by
a few hundred between runs.
One oracle still spelled the header size out by hand: nvme_failure_contract_test.sh corrupts 64
bytes at 208+4L+64, and 208 is now wrong. It reads $HDR like the others do. Worth recording the
measurement, because the obvious conclusion is the wrong one: at v4 that literal lands 56 bytes
into the GDN segment rather than before the payload, so the oracle would have kept corrupting GDN
bytes and kept passing. It was right by luck - 8 bytes of drift per header change, on an assertion
whose whole job is to be exactly where it says it is. tools/nvme_header_layout.sh picked the new
numbers up with no edit (HDR=216, NVME_VERSION=4), which is the version of "the constant has one
spelling" that works.
The delta tier needed none of this. Its manifest already carried weights_fp and kv_format and
already refused a foreign one at match time; it was written after the lesson and took it. What
changed there is only that it no longer owns the fingerprint's implementation, so a store's two
tiers cannot drift on what they are looking at.
…s the 0.1.29 rebase The rebase spliced upstream's inline .map((m) => tail onto our named metricCard const, and the terminator lost its backtick, so the literal ran on into dur() and the whole page failed to parse. Both features are intact: upstream's prefill card + the m.title tooltip, and our CACHE_METRICS block and dur formatter.
maedoc
force-pushed
the
nvme-kv-cache
branch
from
September 30, 2026 17:03
c5108f8 to
21c967d
Compare
Contributor
Author
|
Rebased onto 0.1.29 (d6708a4) and fully re-gated: host suites green (kv_nvme 329 checks, kv_delta 110520, serve 96 passed) and all five device oracles rc=0. |
Owner
|
Thanks, the delta tier idea is good, but this PR carries its own copy of the conversation core, which now conflicts with the one shipped in 0.1.30 (#189), plus ~3k lines of planning docs, web UI and setup changes - too much to gate. Please bring the weight fingerprint, the k8v4 format key and the delta writes onto #190 (the disk tier on the shipped core) or as a follow-up PR on it. Closing this one. |
Contributor
Author
|
ok 👍 |
ScrambledEggsOnToast
pushed a commit
to ScrambledEggsOnToast/Strata
that referenced
this pull request
Oct 19, 2026
Niko1221#259: execution_plan.hpp resolves the option-level contract once - requested and effective speculative widths, slot groups, proposal shape, lookup eligibility, the foreground budget, concurrent-reader exclusions and prompt suspension - and generate.cpp consumes the resolved widths and budget decision. EPLAN reports it before READY; --report-plan prints it before any model allocation. The frontend records it for /metrics and /v1/status. Niko1221#260: ResidentSlots exposes a read-only identity and owns(). Legacy resident admissions and terminals carry request=/incarnation=; BSTOP fences by that identity and ignores a stale owner; the frontend records the admitted identity, sends it on cancellation, and never applies a mismatched terminal. Niko1221#264: opt-in STRATA_DECODE_SERVICE selects ready requests by a weighted 3:1 service cycle (one ticket = one decode window) and keeps an interactive window free of background rows. Measured and left off by default. Niko1221#265: a refused yielded-prefix continuation is preflight-only and ends one request with an ERR terminal; the engine continues serving its peers. Docs: DETAILS.md gains the resolved-plan and decode-service sections. Tests: execution_plan_test, decode_service_test, resident identity checks, frontend lifecycle/stale-terminal cases and the offline roundtrip fixture.
ScrambledEggsOnToast
pushed a commit
to ScrambledEggsOnToast/Strata
that referenced
this pull request
Oct 19, 2026
Niko1221#259: lookup is conditional - it executes in a solo serial serving window after the first window, not in resident decode or the BGEN single admission window. Both earlier flat variants were too broad. Niko1221#264: preserve the cursor's position within a request's weighted ticket block across ready-set churn, so membership changes cannot starve a continuously ready background request. Add the churn regression. Niko1221#265: conversation_snapshot_save's internal synchronize failure is classified as a device fault by its error reason, so the capture is not silently dropped as durability-only.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Builds on the work reviewed in #52, with the two asks it left open addressed: snapshots bound to the weights that wrote them, and
--kv k8v4in the format key.Summary
The v3 header carried the model geometry, and geometry is not identity - two weight sets of one architecture share it exactly, so a snapshot of another model's conversation had everything it needed to match, verify its digest and apply cleanly. v4 makes the header say whose bytes these are.
--kv k8v4in the format key - the key is nowqsa_kv_key, total over all four formats, so a k8v4 store is keyed apart from every other layout instead of being filed as fp16--kv-nvmeis passedTwo things worth flagging: this converges onto #189's shared core as soon as that lands - this branch currently carries the imported copy - and the disk tier in #190 is acknowledged; the coordination is the open item.
The one re-prefill per conversation after the upgrade is the deliberate cost of the version bump: a v3 file is refused by version, never converted.