docs/features/advanced_kv.md

F5 — Advanced KV Cache Techniques

Parent: ../MASTER_PLAN.md Status: planning Owner: TBD

Problem

The vendored llamafile (F2) handles a single model + a single context window competently, but the KV cache is the dominant memory consumer at long contexts and the dominant time consumer across agentic turns that share a prefix. The research landscape has produced several techniques that attack different parts of this problem:

• Context-aware eviction — drop tokens the model is least likely to attend to again, instead of allocating capacity for every token in the window.
• Per-head GPU/CPU split — selective attention heads stay on GPU, the rest stream from host RAM, dramatically shrinking the GPU memory footprint.
• Asymmetric GPU-CPU pipelining — load-aware scheduling that keeps both compute resources busy.
• Layer-ahead pre-computation — the CPU starts the next layer's attention while the GPU finishes the current one.
• Shared compressed pool — a single asymmetrically- compressed KV pool across concurrent agents.

opencoti adopts these as a sequence of patches against the vendored llamafile/llama.cpp tree, each developed and benched in isolation, then validated as a stack.

Goals

• G1. Land a series of numbered patches under vendors/patches/llamafile/ that implement (or port the ideas of) ReST-KV, HeadInfer, NEO, ScoutAttention, and PolyKV.
• G2. Each technique gets its own per-patch bench in perf/llamafile/<technique>.bench.ts measuring the documented claim against baseline.
• G3. A glue milestone (0050-glue-bench.patch) validates the stack composes and provides a standing bench harness for all techniques + the full stack.
• G4. Patches stay rebasable on upstream pin bumps. Each patch ships with a header naming its bench file, milestone, and (when applicable) the upstreaming PR.

Non-goals (for now)

• LMCache — vLLM/SGLang-bound, llama.cpp portability 1/5. Tracked as M7 (deferred). Re-evaluate once M0–M6 are in.
• Full multi-tenant serving infrastructure. PolyKV ships in single-agent mode at M6 (degenerates to a 1-tenant pool); multi-tenant value lights up only when opencoti adds multi-agent fan-out on Tier 0 (future roadmap).
• A single "advanced features" toggle in user config. Each technique is its own opt-in flag with sensible defaults (most default off pending stability).

Sequencing rationale

ReST-KV first (memory savings reduce load on every later technique). HeadInfer second (per-head GPU residency is a clean orthogonal axis on top of eviction). NEO third (load- aware GPU-CPU scheduling assumes per-head granularity to schedule against). ScoutAttention fourth (layer-ahead pre- compute sits naturally on top of NEO's CPU compute path). Glue fifth. PolyKV immediately after glue (keep momentum from the preceding patch series — deferring loses the implementation context and forces a costly cold-context re-visit later).

llama.cpp-portability scores from research:

Technique	Score	License	Reference
ReST-KV	3/5	unknown	arXiv 2605.09649 (Make Each Token Count) — no public code yet
HeadInfer	2/5	MIT	arXiv 2502.12574 — github.com/wdlctc/headinfer
NEO	2/5	Apache-2.0	arXiv 2411.01142 — github.com/NEO-MLSys25/NEO
ScoutAttention	2/5	unknown	arXiv 2603.27138 — code pending (DAC 2026)
PolyKV	1/5	MIT	arXiv 2604.24971 — github.com/ishan1410/PolyKV
LMCache (deferred)	1/5	Apache-2.0	arXiv 2510.09665 — github.com/lmcache/lmcache

Patch number reservations

0001..0005-*.patch                   F4 M3 lazy-slot (these are F4's, not F5's)
0006-kv-reuse-prefix.patch           F5 M0 (was F2 M4; SHIPPED 2026-05-24)
0007-build-header-deps.patch         build infra (SHIPPED 2026-05-27; not a feature — see below)
0010-rest-kv-eviction.patch          F5 M1 (SHIPPED 2026-05-27)
0011-rest-kv-config.patch            F5 M1 (optional split; NOT built — single-patch v1)
0020-headinfer-per-head.patch        F5 M2 (SHIPPED 2026-05-28)
0030-neo-pipeline.patch              F5 M3 (SHIPPED 2026-05-28)
0031-per-stream-split.patch          F4 M3 Phase 4 (SHIPPED 2026-05-29; task #109)
0032-m2-state-io.patch               F5 M2 hardening / bug-222 fix (SHIPPED 2026-05-29)
0033-cpy-tie-blck-align.patch        F5 bug-225 fix (SHIPPED 2026-05-29)
0034-m2-pinned-host.patch            F5 M7-A foundation — pinned host KV (SHIPPED 2026-05-29)
0035-concat-q-block-aware.patch      F5 bug-226 fix (SHIPPED 2026-05-30)
0036-pcie-probe-consume.patch        F5-opt W1 — PCIe probe consume #293 (SHIPPED 2026-05-30)
0040-iqk-flash-attn.patch            F5-opt W2 — wholesale iqk CPU-FA import #290 (SHIPPED 2026-05-30)
0041-q8kv.patch                      RETIRED 2026-05-30 — Q8_KV dropped (#291); redundant 8-bit (q8_0 covers it), TurboQuant M6 is the KV-quant driver
0042-fused-moe.patch                 F5-opt W3 — op-level fused MoE up+gate+GLU #292 (SHIPPED 2026-05-30; separate-up/gate MoE only — N/A to Gemma-4 fused gate_up)
(M5 glue+bench shipped as TS — perf/llamafile/, NOT a patch; old 0050 slot retired)
0070-rolling-kv.patch                F5 M7 Rolling KV #296 (was LMCache; redesigned — see rolling_kv.md) (SHIPPED)
0071-state-io-cosmocc.patch          F5 M7 cosmocc state-IO open fix #349 / bug-269 (SHIPPED; applies after 0070)
0072-poly-kv-pool.patch              F5 M6 PolyKV S1 SharedKVPool #370 (SHIPPED 2026-06-06; applies after 0071)
0073-turboquant-kv.patch             F5 M6 PolyKV S2 TurboQuant tiers+InnerQ+Level-A+Level-B #371 (SHIPPED 2026-06-06; GGML_OP_TURBO_WHT; applies after 0071)
0090-scout-layer-ahead.patch         F5 M4 ScoutAttention (DEFERRED Aug 2026, #264 — moved off the 0040 slot)

Gaps (0002–0009, 0012–0019, ...) are intentional — they leave room for per-technique refinements without renumbering downstream patches.

Milestones

M0 — KV-cache reuse across agentic turns (absorbed from F2 M4; shipped 2026-05-24)

Reuse the resident KV across agentic turns instead of re-prefilling the shared prefix (system prompt + tool defs + history) every turn.

Key finding (shaped the design). The vendored llama.cpp server already reuses KV for a matching token-prefix: cache_prompt defaults to true, and get_available_slot() (server-context.cpp) picks the idle slot whose cached prefix best matches the incoming prompt (above --slot-prompt-similarity, then LRU). So single-session turn-to-turn reuse at the default --parallel 1 works out of the box — the original "add a server-side prefix-hash cache" plan predates that capability. What upstream lacks is session→slot affinity: under --parallel > 1, a concurrent opencode session can be routed to another session's slot and evict a prefix that would otherwise be reused.

What shipped.

• vendors/patches/llamafile/0006-kv-reuse-prefix.patch (touches only tools/server/{server-task.h,server-task.cpp,server-context.cpp}, disjoint from the F4 0001–0005 lazy-slot patches): an optional session_id request field + a session_id → slot.id affinity map. get_available_slot() prefers the session's own idle slot (its resident KV is that session's prefix) before the existing LCP→LRU path; launch_slot_with_task() records the binding, dropping any stale session on that slot. Regression shield: empty session_id (every non-opencoti client) skips the affinity phase, so upstream slot selection is byte-for-byte unchanged.
• @opencoti/llamafile: a session-aware fetch wrapper (withKvReuseFetch) tags each chat-completions body with the session's id
  • cache_prompt: true; TierStreamRequest.sessionID already carries it, so no runtime.ts / surgical-hook change. Typed, env-backed server flags (--parallel/-sps/--cache-reuse/-c), omitted when unset.
• "(session_id, prefix) reuse" falls out: affinity routes to the slot, the server's existing token-LCP supplies the prefix match — no separate prefix hash needed.

Bench (perf/llamafile/turn-2-latency.bench.ts, live on solidPC, Qwen2.5-Coder-0.5B, CPU): scenario A turn-2 prefill collapses 2061 → 20 tokens (202ms → 12ms; cache_n 2076) vs a no-cache control at 2096 tokens; scenario B (--parallel 2, two sessions sharing a system prefix) each reuse their own KV under affinity (cache_n 1990 each). 5/5 checks pass.

Deferred (optional 0007, evidence-driven): an explicit prefix_hash

• cross-restart disk persistence (--slot-save-path keyed by session_id+hash). Within a server lifetime the resident-token LCP already supplies the prefix match, so the hash adds nothing until restart-durable or distributed reuse is wanted. No plugin, no HTTP surface, no TS client, no opencode surgical hook — the surgical-hook grep count stays 18 (the vendored C++ marker is // opencoti F5 M0, a vendored-source tag, not a counted hook).

Build gotcha (buglog bug-167): server-task.h is included by server-context.cpp, server-task.cpp and server-queue.cpp. Editing that header and rebuilding with bun run build:llamafile:make leaves server-queue.cpp.o compiled against the old task_params layout → ABI skew → bad_array_new_length on every generation request. After a vendored header change, force a consistent recompile (rm -f vendors/sources/llamafile/o//llama.cpp/tools/server/*.o then :make, or the full reproducible build:llamafile).

M1 — ReST-KV / retention-aware eviction (shipped 2026-05-27)

When a slot exceeds n_ctx, upstream context-shift keeps the first n_keep tokens and blindly discards the positional middle n_discard chunk. M1 makes the choice retention-aware: discard the lowest-value contiguous window instead, so long-context generation keeps the evidence that matters.

Reference decision (resolved the open question). The named paper (arXiv 2605.09649) is TRIM-KV — learned retention gates trained by distillation, needing a per-model pipeline + published gates; not a drop-in runtime patch. SnapKV / H2O score tokens from the materialized attention matrix, which Flash Attention never produces (the vendored pin defaults flash_attn_type = AUTO) — incompatible with the FA path opencoti wants on. So M1 substitutes a score-free, FA-compatible criterion: KeyDiff (arXiv 2504.15364) — evict the keys least distinctive by cosine similarity (distinctiveness = mean over heads of 1 - cos(key_head, mean_head); higher = keep). domvox/triattention-ggml is a working ggml reference but scores every decode interval (K copied GPU→CPU each --tri-interval, "the main bottleneck"); M1 avoids that by scoring only at the infrequent context-shift trigger, amortizing the K read.

Depth decision: server-layer window (contiguous). Eviction stays contiguous so the server's CPU token mirror (slot.prompt.tokens) stays 1:1 with KV positions and all prefix-reuse / sampler logic is reused untouched. Only which width-n_discard window is discarded changes — the minimum-score one over [n_keep, n_tokens − rest_kv_recent), not the positional middle. Full per-token scattered KeyDiff (a core evict_to_budget + a server rework to tolerate holes) is deferred — revisit only if M1's bench shows window granularity costs meaningful quality.

What shipped.

• vendors/patches/llamafile/0010-rest-kv-eviction.patch (build-time; no opencode upstream file touched):
  • src/llama-kv-cache.{h,cpp}: concrete seq_key_scores(seq_id, layer_hint) — enumerates the seq's cells in position order, reads each post-RoPE K row to host, reshapes into n_head_kv heads, returns the per-position KeyDiff score over a representative layer (default mid).
  • ABI-safe interface choice: src/llama-memory.h gains a non-pure virtual seq_key_scores with an empty default body — no backend is forced to implement it and no vtable churn/RTTI is needed (the pure-virtual llama_memory_i methods are left untouched).
  • src/llama-kv-cache-iswa.{h,cpp}: delegate to the base cache.
  • include/llama.h + src/llama-context.cpp: C shim llama_memory_seq_key_scores(...) that returns 0 for non-KV memory (or a key type with no float converter) → the caller falls back to positional.
  • tools/server/server-context.cpp: inside the existing shift block, when rest_kv_eviction is on, build prefix sums over the scores and pick the min-sum window start w0; run the existing seq_rm/seq_add/mirror rewrite parameterized by w0 instead of n_keep. Off (default) = byte-for-byte upstream (the off-path never calls the scorer).
  • common/common.h + common/arg.cpp: --rest-kv-eviction / --rest-kv-recent N / --rest-kv-layer N (server example).
• @opencoti/llamafile (config.ts/launch.ts): typed, env-backed (OPENCOTI_LLAMAFILE_REST_KV_*) fields → buildServerArgs emits the flags when set, omits when unset. Global flag, not per-request → no fetch wrapper.

Bench (perf/llamafile/rest-kv-eviction.bench.ts, needle-in-the-middle, live on solidPC, Qwen2.5-Coder-0.5B, CPU): hard checks 3/3 — a context shift fired, the rest-kv eviction path executed only with eviction on, and the off-path never logged it (proving the off-path is upstream-identical). These deterministic, model-independent checks gate the exit code. The quality-delta headline is soft and model-dependent: a 0.5B model either states the needle code before the shift fires (generous headroom → both runs recover) or degrades so far after the shift that neither restates a random code even when the needle survives in KV — demonstrating the delta needs a stronger model. 0011 stays reserved (single-patch v1, mirroring M0's 0006).

Surgical-hook count stays 18 — the C++ marker is // opencoti F5 M1 rest-kv-eviction, a vendored-source tag, not a counted opencode hook.

Build-infra prerequisite (0007, buglog bug-170). Iterating M1's vendored C++ surfaced that the build had no working header-dependency tracking (mkdeps' o//depend can't resolve llama.cpp's -iquote includes → zero src edges → header edits never recompiled → forced from-scratch rebuilds) and that ccache cached nothing (a .d-name mismatch made it abort every compile with "internal error"). Patch 0007-build-header-deps.patch fixes both (-MMD -MF $@.d in the compile rules + -include the per-object .o.d), and build-pipeline runs the full reset-build with CCACHE_RECACHE=1 so a warm cache can't orphan a cosmocc fat object's .aarch64 twin. See docs/features/llamafile_build.md. This is what made M1 (and every later F5 milestone) tractable to iterate.

Deferred / not done: scattered per-token KeyDiff; quantized-K cosine scoring (falls back to positional when K is quantized); the optional 0011 per-request override; the full reset-based reproducible build verification (would git clean -fdx the in-tree vendored state — run when a reset window is authorized; 0010/0007 are reverse-verified against pristine).

M2 — HeadInfer (head-wise GPU/CPU split) (shipped 2026-05-28)

When a model is GPU-offloaded, the entire KV cache for offloaded layers lives in VRAM and caps how much context fits on the card. M2 keeps a configurable fraction of attention heads' KV resident on GPU and offloads the rest to host memory, reassembling the full head set per attention step via ggml_concat on the head axis. Resident VRAM drops proportionally to the offloaded fraction; correctness is preserved.

Reference decision (resolved at impl time). The named paper (arXiv 2502.12574, wdlctc/headinfer) is HF-Transformers; not portable. A parallel investigation of /shared/dev/lightseek — which has deep, working llama.cpp head-splitting surgery — confirmed it is distributed tensor-parallelism (heads split across MPI/RPC ranks, each on its own GPU) on a different llama.cpp lineage (vanilla upstream, not the Mozilla-Ocho llamafile fork). lightseek informed the design (GQA-ratio guard, "logical dims full / physical alloc partial" framing) but its backend-level machinery is reference only, not code to port — porting it wholesale would blow past opencoti's minimal-additive-patch rule.

Depth decision: head-RESIDENCY split, not head-COMPUTE split (foundation-first). M2 delivers offload + scheduler-driven stream-back for correctness; the per-head attention COMPUTE split (run CPU-resident heads' attention on the CPU backend to avoid the per-step stream-back) is the headline perf win and is deferred to optional 0021 — revisit only if the bench shows the stream-back cost dominates. The "heavy" depth option, consciously deferred (mirrors M1's window-vs-scattered).

What shipped.

• vendors/patches/llamafile/0020-headinfer-per-head.patch (build-time; no opencode upstream file touched; 14 files / 28 hunks / 520 lines):
  • src/llama-kv-cache.{h,cpp}: per-layer kv_layer gains optional k_cpu/v_cpu sub-tensors + a gpu_heads count (0/absent = no split). Constructor allocates the GPU half (gpu-head rows) in the device buffer-type and the CPU half (cpu-head rows) in a CPU buffer-type via the existing ctx_for_buft(ggml_backend_cpu_buffer_type()) — a CPU sub-tensor on an otherwise-GPU layer just adds a CPU ctx/buffer pair, no allocator rework. GQA-safe split: clamp gpu_heads to a whole number of KV-head groups (lightseek lesson). Phase-1/2/3 hooks (ensure_cleared, shrink_if_idle) loop the extra sub-tensor.
  • get_k/get_v: when split, build the gpu/cpu views and ggml_concat(ctx, kg_v, kc_v, /*dim=*/1) — the head axis is dim 1, so the concat produces the same [head_dim, n_head_kv, n_kv, ns] shape the non-split path returns. Backend scheduler auto-inserts the cross-backend copy. Off path (no split) is unchanged.
  • cpy_k/cpy_v: split branch slices k_cur by head → two ggml_set_rows
    • a dependency-tying ggml_add so the two scatter nodes propagate in one graph node. Off path unchanged.
  • ggml/src/ggml-cuda/ggml-cuda.cu: one-line fix to CUDA supports_op for GGML_OP_CONCAT — was over-broad (any non-I32/I16) while the kernel ggml_cuda_op_concat is F32-only (three asserts at concat.cu:158-160). Tightened to op->src[0]->type == GGML_TYPE_F32 so F16 concat correctly routes to the CPU backend (which has a real F16 path at ggml-cpu/ops.cpp:1980 concat_f16). This is a genuine upstream bug — supports_op was lying about what the kernel implements.
  • Plumbing: common/common.h + common/arg.cpp add --headinfer-gpu-heads-frac F (server example), threaded through cparams/llama-context.cpp/llama-cparams.h/llama-model.cpp into the cache constructor (mirroring the --no-kv-unified wiring template). Both iswa and hybrid memory ctors forward 1.0f (off) — they don't thread headinfer themselves; the split is unified-mode only by design.
• @opencoti/llamafile (config.ts/launch.ts): typed, env-backed (OPENCOTI_LLAMAFILE_HEADINFER_GPU_HEADS_FRAC) headinferGpuHeadsFrac field → buildServerArgs emits --headinfer-gpu-heads-frac only when set (omitted = server default = off). Three new launch-args tests (omit/emit/precede-extraArgs); typecheck clean, 78/0 tests.

Bench (perf/llamafile/headinfer-residency.bench.ts, live on solidPC, RTX 3090, Qwen2.5-Coder-0.5B-IQ4_XS, -ngl 99 -fa on --parallel 1, ctx 4096): 6/6 PASS —

	baseline frac=1.0	split frac=0.5
CUDA0 KV	48 MiB	24 MiB (-50%)
CPU KV	0 (off-path identity)	24 MiB (+ host half)
greedy decode (temp=0)	reference	byte-identical to baseline

Identical-text is the strongest correctness check possible (implies identical token ids; stronger than a cosine ≥ 0.999 bar). Reproducible against Qwen2.5-1.5B-Instruct at the same frac (linearly scaled, 112 → 56 / 56). The deferred-perf cost is the GPU→CPU→GPU concat bounce per step; option (a) two-ggml_cpy-into-preallocated-dest can replace concat for steady-state perf in a follow-up (post-M2 or 0021).

Off-path identity: frac=1.0 (default, every non-opencoti client) — k_cpu is never allocated, get_k/get_v/cpy_k/cpy_v take the non-split branch; the binary's KV path is byte-for-byte upstream-equivalent.

Surgical-hook count stays 18 — the in-tree marker is // opencoti F5 M2 headinfer, a vendored-source tag, not a counted opencode hook.

Build-infra prerequisite that surfaced (buglog bug-178). Iterating M2's vendored C++ surfaced that an earlier CCACHE_RECACHE=1 from-scratch rebuild (M1-E / #244 lineage) had silently produced a binary with broken C++ exception unwinding — any throw aborts with ud2 / SIGILL. This is a distinct variant of bug-176: that defect's known symptom was a link failure (concomitant .aarch64 file missing); this variant links successfully but the consolidated .eh_frame in the APE is corrupt, so the unwinder traps — masked from M1-E because --version (the only thing M1-E verified) never exercises unwinding. Disambiguated from non-causes: NOT app/M2 source (gated-off headinfer_gpu_heads_frac=1.0 AND pure-upstream std::stoi both trap); NOT toolchain or flags (a 5-line throw/catch compiled with the exact build flags via cosmoc++ catches correctly). Fix: rebuild non-destructively with OPENCOTI_NO_CCACHE=1 bun run build:llamafile:make (the make subcommand calls runMake() which never git-resets the submodule; OPENCOTI_NO_CCACHE=1 drops the ccache CC=/CXX= wrapper so cosmocc itself writes both fat-object arches in one invocation, no stale x86-only twin). Lesson: every llamafile build verification MUST run a real CPU inference or a throw-probe (-n notanum), never just --version.

Deferred / not done:

• Per-stream device buffers + host-pageable (#109 Phase 4): the M2 plan's gated extension. Gate fired — landing per-stream now would expand the cpy_k/cpy_v multi-stream scatter surface (the plan's flagged hot-path regression risk) with no validated need; M2-A/B/C delivered the head-residency headline win on unified mode cleanly without it. Carries to optional 0021.
• Per-head attention COMPUTE split (run CPU-resident heads' attention on the CPU backend to eliminate the per-step stream-back). The headline perf win; consciously deferred per the foundation-first depth decision above. Revisit only if a future bench shows the stream-back cost dominates.
• Quantized-K head-split edge cases; multi-GPU head residency — later tuning.
• Steady-state concat optimization (two ggml_cpy into a pre-allocated F16 destination on GPU): replaces the GPU→CPU→GPU concat bounce with a single CPU→GPU copy + a GPU-side memcpy. Larger surgery, post-M2 follow-up.

M3 — NEO (asymmetric GPU/CPU attention pipelining) (shipped 2026-05-29)

Reference: arXiv 2411.01142, Apache-2.0 at github.com/NEO-MLSys25/NEO. Reference is swiftLLM (PyTorch). Like M2's lightseek precedent, not directly portable — NEO ports the structural split + concurrent dispatch concept into ggml's decode loop rather than wrapping the upstream code. The published design covers three pieces (per-head COMPUTE split + async overlap + load-aware scheduling); M3 bundled all three into one 0030-neo-pipeline.patch per locked decision (absorbing the 0021 "per-head compute split" reservation from M2).

Depth decision: full structural-split + orchestrator + load-aware controller was the planned scope; the orchestrator and structural split shipped; the EWMA load-aware controller (M3-C) was deferred after the M3-D gate fired (see "What shipped" below).

Patch: vendors/patches/llamafile/0030-neo-pipeline.patch. Sits on top of 0020 (M2 head-residency split). The historical "per-head compute split" reservation at 0021 was absorbed by 0030 (M3 bundled it in). F4 M3 Phase 4 (per-stream KV tensor split, SHIPPED 2026-05-29 — see docs/decisions/0001-lazy-slot-context.md §Phase 4 — shipped notes and the patches README row for 0031-per-stream-split) lands at 0031, after 0030, because Phase 4-D lifts M2's dedicated get_k_gpu/get_v_gpu accessors that 0030 introduces for NEO's two-FA dispatch. The originally-reserved 0021 slot stays vacant.

What shipped:

• M3-A structural two-FA split. New llm_graph_context::build_attn_mha_neo in llama.cpp/src/llama-graph.cpp is invoked from build_attn's flash-attn branch when cparams.neo_pipeline_mode != 0 && mctx_cur->headinfer_split_active(il). It slices Q on dim 1 at the GQA-clamped gpu_heads * gqa_ratio boundary (a clean ggml_view_4d with no extra permute), slices mask and sinks correspondingly, builds TWO ggml_flash_attn_ext ops on disjoint head ranges, and concats the outputs on the head axis. The K/V halves come from new llama_memory_context accessors get_k_gpu/_cpu/get_v_gpu/_cpu

  • headinfer_split_active(il) (added in llama-memory.h, implemented in llama-kv-cache.{h,cpp}, delegated through llama_kv_cache_context wrappers). The scheduler auto-routes each FA op to the backend hosting its K/V views (CUDA for cur_g, CPU for cur_c) — no scheduler surgery needed. This eliminates M2's per-step O(n_kv × head_dim_cpu × 2 B) CPU→GPU stream-back and replaces it with an O(n_tokens × head_dim_cpu) bounce on the FA output. The original get_k/get_v paths are unchanged so --neo-pipeline off runs the M2 concat path byte-identically.

• M3-B NEO orchestrator (wired-but-no-observable-win). New vendored source files vendors/sources/llamafile/llama.cpp/ggml/include/ggml-neo-pipeline.h

  • ggml/src/ggml-neo-pipeline.cpp provide a registry of (cur_g, cur_c) pairs + a cudaEvent slot per cur_g. The graph builder calls ggml_neo_pipeline_register_pair(cur_g, cur_c) after constructing each pair. The CUDA backend's ggml_backend_cuda_graph_compute consults the registry: if the split's first node is a registered cur_g, it sets cuda_ctx->curr_stream_no = 1 so all kernels dispatch on the alternate CUDA stream, records a cudaEvent on stream 1 at end-of-graph, then restores curr_stream_no = 0. When a downstream split reads a registered cur_g (typically the concat join), the hook inserts a cudaStreamWaitEvent(cuda_ctx->stream(), event, 0) so the consuming stream waits for the stream-1 work. CUDA graph capture is disabled for NEO-engaged splits; off-path graphs keep the fast cuda_graph path.

• M3-D adapter. --neo-pipeline off|on|auto in common/arg.cpp (env: LLAMA_ARG_NEO_PIPELINE). cparams.neo_pipeline_mode field plumbed through include/llama.h, common/common.{h,cpp}, llama-cparams.h, llama-context.cpp. @opencoti/llamafile typed field neoPipelineMode: "off"|"on"|"auto"|undefined + env OPENCOTI_LLAMAFILE_NEO_PIPELINE mirror the M2 frac template; 3 new launch-args tests pass (13 total).

• Build system. New ggml-neo-pipeline.cpp listed in llama.cpp/BUILD.mk GGML_SRCS_CPP (for the main binary) and in llamafile/build-functions.sh compile_ggml_core source list (for the CUDA DSO). Both rebuild cleanly; 14 ggml_neo_pipeline_* symbols exported from ggml-cuda.so.

Bench numbers — perf/llamafile/neo-pipeline.bench.ts (solidPC RTX 3090

• Qwen2.5-Coder-0.5B IQ4_XS, ctx=4096, predict=128, frac=0.5):

Run	tps	Δ vs M2
B baseline (no split)	326.69	+75%
M m2-only (split, M2 path)	186.96	—
N neo-on (split, M3 path)	187.04	+0.04%
O off-identity (--neo-pipeline off)	151.65	—

Hard checks:

• C1 PASS — N completion === B completion (correctness intact).
• C2 FAIL — N tps - M tps < 5% (observed 0.04%).
• C3 PASS — N GPU KV (24 MiB) === M GPU KV (24 MiB) — residency preserved.
• C4 PASS — O completion === M completion (--neo-pipeline off byte-identical).

Soft check S1 (NEO ≥ 80% of baseline tps): WARN at 57.3% — the CPU FA on half-heads is the wall-time bottleneck; no overlap can recover the baseline because the baseline runs zero CPU work.

Why C2 failed: the M3-B stream-swap mechanism is correct (M3-A cur_g is in fact dispatched on CUDA stream 1, downstream consumers correctly wait on the stream-1 event, cur_c runs concurrently on the CPU backend, the GPU FA finishes long before the CPU FA), but for Qwen-Coder-0.5B at ctx=4096 the per-step GPU FA on half-heads finishes in single-digit μs while the CPU FA on half-heads takes ~10× longer. Total wall time is bounded by max(GPU FA, CPU FA) + concat which is dominated by the CPU FA. Even at long context (ctx=16384, predict=256 with a 4480-char prompt) the speedup stayed marginal (+1.4%, still below C2's 5% threshold). The overlap WOULD pay off on workloads where GPU FA cost approaches CPU FA cost — multi-GPU splits, much larger models with smaller frac, or quantized-K paths where the GPU side does more work per head. M3-B is correct infrastructure for those regimes and for M4 (ScoutAttention) to reuse.

Per the M3-D gate decision: C2 failed but the orchestrator is correct (no crashes, no deadlocks, C1+C3+C4 pass at the short-ctx bench) and the off-path is byte-identical to the M2-shipped binary. The orchestrator ships wired-but-no-observable-win (not dormant-stub as the original plan text described) — the CUDA hook is live and registered pairs route to stream 1 — but the mechanism produces no measurable throughput improvement on this hardware/model. The M3-A structural win (eliminating the per-step CPU→GPU K/V stream-back) is the actual M3 deliverable.

M3-C (EWMA load-aware controller) deferred. The plan called for a per-layer neo_layer_state measuring GPU/CPU FA times and adjusting gpu_heads_active toward load balance. Without an observable M3-B overlap win this is cosmetic — there's nothing for the controller to optimize toward. Revisit when M3-B's mechanism finds a regime where it actually wins (multi-GPU, much larger models, quantized-K), or when M4 ScoutAttention reuses the same orchestrator and the load-balance trade-off becomes meaningful.

Off-path identity proven. M2's perf/llamafile/headinfer-residency.bench.ts re-runs 6/6 PASS against the M3-built binary — the M2 concat path is byte-equivalent to the 0020-shipped binary. 0030 does not regress M2. Surgical-hook count stays 18 (all M3 changes are in vendored llama.cpp source captured into 0030, not in upstream opencode packages/).

Build / capture rules carried forward: OPENCOTI_NO_CCACHE=1 bun run build:llamafile:make for the main binary; bun run build:llamafile:cuda for the CUDA DSO (~24 min build on solidPC). Patch captured via snapshot-diff under .opencoti/snap-m3-pre/ (NOT git diff HEAD; bug-121). Apply-tested clean against the snapshot. Vendored submodule's dirty state intentionally not staged on the substance commit.

Deferred / non-goals (carry forward):

• M3-B win on this hardware — see Why C2 failed; revisit when workload reshapes.
• M3-C EWMA load-aware controller — see above.
• Multi-GPU NEO (orchestrating multiple CUDA devices alongside CPU) — single-CUDA-device + CPU is the current scope.
• Vulkan / ROCm backends — CUDA-only; the orchestrator's CUDA hook is unreachable on other GPU backends, so M3-A's structural split runs without M3-B's stream-swap there.
• Per-stream device buffers + host-pageable (#109) — still carries from M2 / F4 Phase 4.
• Persistence of EWMA state across server restarts — moot since M3-C deferred.
• Reallocating gpu_heads mid-decode — the runtime controller is deferred; physical gpu_heads is still ctor-set.

M4 — ScoutAttention (layer-ahead CPU pre-compute) (deferred 2026-05-29)

• Reference: arXiv 2603.27138 — paper published 2026-03-28, accepted DAC '26 (July 26-29, Long Beach). Authors: Qiuyang Zhang, Kai Zhou, Ding Tang, Kai Lu, Cheng Li, Zhenyu Yang, Peng Xu, Jiguang Wan (HUST + Huawei + Zhejiang Lab).

• Reference-code status (verified 2026-05-29): ABSENT. Authoritative search via gh search repos/gh search code for ScoutAttention, scout-attention layer-ahead, layer-ahead pre-computation, asynchronous periodic recall sparse attention returns zero implementation hits. HF Hub model + space search empty. The full paper HTML has no GitHub/GitLab/artifact link and no code-availability statement. Authors are HUST + Huawei + Zhejiang Lab — academic publication, code not advertised.

• Algorithmic substance (from paper Algorithm 1, recorded so a future picker-up has the spec cold):

  1. Predict next-layer query Q_pred^(i+1) ← W_Q^(i+1) · X^i (residual- similarity trick, cosine > 0.93 across Gemma 3, Llama 3.1, Mistral, GLM 4 tested in the paper)
  2. Block-wise top-k: B_pred^(i+1) ← TopK(Q_pred^(i+1) · K_digest^(i+1)^T)
  3. Identify CPU-resident blocks: B_cpu^(i+1) ← B_pred^(i+1) \ B_gpu^(i+1)
  4. Async spawn CPUAttn(B_cpu^(i+1))
  5. GPU runs full layer i (attn + FFN + QKV projections) overlapped
  6. A^i ← Merge(A_gpu^i, A_cpu^i) via FlashAttention's online-softmax merge (the previously-spawned CPU work joins here)

  Plus asynchronous periodic recall (β = 12% threshold → avg ~8.7 layers between recalls, CPU compute ratio ~8.2%). Hyperparameters tested: block size 32 default, sparse budget 1024-2048, ctx up to 64k, batches 16-64, Qwen3-8B/14B/32B.

• Why portability is harder than M2/M3. Paper ships on SGLang + FlashInfer + IPEX. M2 (lightseek) and M3 (NEO/swiftLLM) were paper-ports onto existing ggml primitives. M4 requires new primitives: a block-wise top-k CUDA kernel (FlashInfer-equivalent), a block-sparse variant of ggml_flash_attn_ext, a K_digest side tensor + its computation kernel, and an IPEX-equivalent CPU worker accepting a selected-blocks input (llama.cpp's flash_attn_ext CPU kernel does full-K/V only). Multi-week kernel work with C1 byte-identity risk on every downstream check.

• Decision (user 2026-05-29): defer until reference code surfaces. Re-evaluation August 2026 (post-DAC). Progression updates to M5 Glue ("compose-validate M0+M1+M2+M3", M4 removed from the compose set until it lands) → M6 PolyKV (has reference code) → revisit M4.

• If/when M4 lights up later, options remain on the table: (a) reimplement from spec, (b) substitute a simpler layer-ahead K_pred-prefetch-only approximation that skips the block-sparse FA refactor, or (c) compose with M3's NEO orchestrator (run NEO's GPU half dense, run the CPU half block-sparse) — the registry surface added in 0030-neo-pipeline.patch is the natural integration point.

• Config flag (when M4 ships): --scout-attention on|off (default off).

M5 — Glue + bench harness (shipped 2026-05-29)

Goal. Compose-validate the F5 stack — M0 (session-keyed KV reuse) + M1 (retention-aware eviction) + M2 (head-residency split) + M3 (NEO concurrent FA). Each prior milestone landed its own bench in isolation; M5 is the first that lights all four up simultaneously. M4 is removed from the compose set per its 2026-05-29 deferral.

What shipped.

• perf/llamafile/_bench-lib.ts — first shared helper module under perf/llamafile/ (293 lines). Exports resolveBin, resolveModel ("qwen" | "gemma" | "auto" | <abs-path>), pickPort, probeGpu (enforces the 2026-05-28 Cerebrum free-VRAM rule), spawnServer (returns {baseURL, stop, log}; stop() is the bug-191 SIGTERM→2 s race→SIGKILL reap with await proc.exited in both branches), waitForHealth, metric, parseKvBuffers, callCompletion, tps. Underscore prefix keeps the file out of *.bench.ts glob runners. Existing four benches stay on their inlined helpers — no regression risk to landed work; future benches (M6 PolyKV onwards) inherit the lib.

• perf/llamafile/advanced-kv-stack.bench.ts — the compose bench (510 lines). Six configurations × one multi-turn needle workload × 9 hard checks + 2 soft checks. CLI: --model qwen|gemma|<abs-path> defaulting to auto→qwen for fast CI; --ctx, --predict, --filler, --frac for one-off tuning. Self-skips on no bin / no model / no GPU+DSO / insufficient free VRAM. Results JSON to .opencoti/m5-stack-bench.json.

Configurations.

ID	Flags	Tests
B	-ngl 99 -fa on --parallel 1 -c 1024	Baseline
A0	B + --kv-unified --parallel 2 -c 2048 --cache-reuse 256	M0 alone
A1	B + --rest-kv-eviction --rest-kv-recent 256 --rest-kv-layer -1	M1 alone
A2	B + --headinfer-gpu-heads-frac 0.5	M2 alone
A3	B + --headinfer-gpu-heads-frac 0.5 --neo-pipeline on	M3 alone
S	B + ALL of A0+A1+A2+A3 simultaneously (parallel=2, ctx=2048, kv-unified, all four flag groups)	The compose validation

Live bench (solidPC RTX 3090, Qwen2.5-Coder-0.5B IQ4_XS, ctx=1024 per slot, n_predict=256, filler=35, frac=0.5):

B  baseline    : t1_tps=328  cuda_mib=12  (per cell 0.0117 KiB)
A0 M0 alone    : t1_tps=350  cuda_mib=24  (per cell 0.0117 KiB; -kv-unified at parallel=2)
A1 M1 alone    : t1_tps=337  cuda_mib=12  (per cell 0.0117 KiB)
A2 M2 alone    : t1_tps=159  cuda_mib=6 + cpu_mib=6  (per cell 0.0059 KiB — 50% saving)
A3 M3 alone    : t1_tps=160  cuda_mib=6 + cpu_mib=6  (NEO on; identical residency to A2)
S  STACK       : t1_tps=159  cuda_mib=12 + cpu_mib=12  (per cell 0.0059 KiB — M2 saving preserved)

Hard checks: 9/9 PASS  (C3, C7 cleanly SKIPPED — workload didn't force
                       context shift on this model + ctx + filler combo;
                       the bench mirrors M1's skip-not-fail pattern.)
Soft checks: 2/2 PASS

Key findings.

1. M1 + M2 compounding cost is negligible. Phase-1 explore flagged a soft compose risk — M1's seq_key_scores dequantizes K rows at context-shift; M2's get_k concatenates k_cpu on every step. The bench measured S t1_tps 158.55 vs A2 alone 158.90 — within noise (soft check W1 PASS). No follow-up needed.

2. M2's per-cell GPU saving is exactly preserved in the stack. Soft check W2: S per-cell CUDA = A2 per-cell CUDA = 0.0059 KiB/cell. The M2 split is identically engaged in S as in A2 alone.

3. --kv-unified must be explicit when --parallel > 1. First compose-conflict surfaced: llamafile's llama_context_params defaults kv_unified = false, so --parallel N > 1 → n_stream = N, and M2's split gate (patch 0020 line 176, n_stream == 1) silently fails. M0 + M2 wouldn't have composed otherwise. The bench encodes the fix: A0 and S pass --kv-unified explicitly. This is exactly the kind of finding M5's compose-validation exists to surface. A future M5 follow-up could make --kv-unified the default when --headinfer-gpu-heads-frac < 1.0, or document a hard error if both --parallel > 1 and frac are set without --kv-unified. Tracked as a soft TODO in the bench's config block.

4. M3 stack equality. Hard check C5: A3 turn-1 completion === A2 turn-1 completion (byte-identical). The 0030 patch's off-or-on structural equality on M2 is preserved when other flags are on.

5. No pathological combined slowdown. C9 floor at 40% of baseline (A2 alone runs at 48% of B — the cost of M2's CPU stream-back). Stack tps lands at A2 alone's level (158 vs 159 — no additional M0/M1/M3 cost stacked on M2's floor).

Deferred / non-goals.

• Helper-module extraction in M0/M1/M2/M3 benches. Per the M5 scope-locking decision, leave the four existing benches on their inlined copies to protect landed work. Follow-up task can refactor them when convenient; the lib API was designed compatible.
• Force shift on Qwen workload. Tuning (filler, predict) to cross n_ctx mid-decode under Qwen2.5-Coder IQ4_XS hit a BPE merge-breakpoint at filler ≈ 36-37 that jumps prompt tokenization from ~750 → ~1050. Filler=35 leaves prompt at ~740 + decode 256 = 996 < 1024 — under the shift threshold. C3 / C7 skip cleanly rather than false-fail; override with --filler N --ctx N --predict N from the CLI for one-off shift-forced runs. The M1 bench already exercises shift in isolation, so M5's skip is fine for compose-validation.
• CI runner / fixed GPU image (the F2 M6 dependency from the original M5 sketch) — out of scope; the bench self-skips when no GPU is present, so it's CI-safe today.
• Make --kv-unified the auto-default when M2 is engaged. Could land in a tiny 0030-tier follow-up patch; deferred until either M6 PolyKV or a user request makes it worth touching the cparam wiring.
• Gemma profile run. Bench supports --model gemma (probes for 14 GB free VRAM); Qwen-default run is the CI path. Realistic-perf Gemma numbers can land in a future one-off study.

Hook count unchanged (git grep "opencoti-hook:" packages/ docs/ returns 19 — same as before M5, where the 19th is a docs-text mention from the M3-H ship, not a real new surgical hook). No new marker in opencode source, no llama.cpp source change, no patch. 0050-glue-bench.patch reservation slot in vendors/patches/llamafile/README.md stays empty — M5 is pure application-side perf harness.

M6 — PolyKV (shared compressed KV pool)

Child plan: poly_kv.md — M6 SHIPPED 2026-06-06. Staged fast-first: S0 asymmetric KV-compression via stock -ctk q8_0 -ctv q4_0 through M7's dequant-on-lift (zero new kernel, banked correct at 4k) → S1 shared read-only prefix pool (multi-tenant core, O(1) in agents, 6.9× fan-out memory win) → S2 TurboQuant TURBO{2,3,4,8}_0 family + InnerQ + Level-A dequant-on-lift + Level-B fused FA-VEC (GGML_OP_TURBO_WHT; 4/4 tiers logit-equiv) → S3 glue bench + compose gate + patches 0072/0073 (captured + byte-identical-proven). Key insight: M7's dequant-on-lift (S3d) lets opencoti decompress the pool in-attention, a win the PyTorch reference (decompress-then-attend) can't get. Remaining: S4 MTP draft head (#373).

• Reference: arXiv 2604.24971, MIT at github.com/ishan1410/PolyKV. Single asymmetrically-compressed KV pool across multiple concurrent agents; ~97.7% memory reduction at 3–4+ concurrent agents.
• Sequenced immediately after M5, not deferred. Reasoning: deferring loses the implementation context from the preceding patch series — the same code paths (slot KV management, allocator hooks, attention kernel modifications) get re-visited with cold context if PolyKV lands in a separate work unit later. Keeping it in this work unit reuses the mental model. The fact that opencoti's multi-agent fan-out on Tier 0 isn't shipped yet is not a blocker — PolyKV is a memory-pool architecture that can land in single-agent mode (degenerates to a 1-tenant pool) and light up its full value once fan-out arrives.
• Smoke at M6: single-tenant correctness. Multi-tenant bench follows whenever multi-agent fan-out lands.
• Compose-with-glue: re-run M5's advanced-kv-stack.bench.ts with PolyKV on the stack; document any conflict resolution.

M7 — LMCache (DEFERRED)

• Reference: arXiv 2510.09665, Apache-2.0 at github.com/lmcache/lmcache. Cache offloading + prefix-decode disaggregation; modular connector for vLLM/SGLang.
• Deferred because: vLLM/SGLang-bound architecture, llama.cpp portability 1/5. Re-evaluate when M0–M6 are landed and there's a clear marginal-win case.
• Superseded 2026-05-29: the M7 slot is now Rolling KV (streaming double-buffered KV pipeline, M2-as-runtime-tactic), designed in rolling_kv.md, patch 0070. LMCache remains an external reference only.

F5-opt — concurrency=1 optimization round (task #287)

RESOLVED 2026-06-04 — superseded by M7. A fresh concurrency=1 re-baseline of the shipped M7 binary (default --kv-residency-mode auto) vs vanilla-0.10.1 on the prefill-bound regime that defined this task lands at parity: prefill 4k/32k ratio 1.00/1.01, decode 1.03/0.94 — all in [0.94, 1.03]. The old worst case (niah@32k prefill, 3.6× = 0.28) is now 1.01. The cause is structural: M7's GPU_RESIDENT-by-default retired the always-on M2 CPU-half this round was created to optimize (server log: GPU_RESIDENT=5/25, POSITION_WINDOW=0, CPU_FA_TAIL=0 at every tested ctx). The lone residual — decode@32k 0.94 (~6%, per-layer tactic-table bookkeeping on the resident path) — is a future micro-opt, not a blocker. Full matrix + artifact note: .opencoti/m7-c1-rebaseline/RESULT.md. The four workstreams below shipped en route and remain the building blocks (W1 probe, W2 CPU-FA fallback, W3 fused-MoE, W4 the Rolling KV residency redesign itself).

The 2026-05-29 RULER F5 comparison held quality at 100% (4k/32k) but ran 2.6–4.7× slower than vanilla at concurrency=1, dominated by the M2 CPU-half. Four workstreams attack compute / transfer / residency before the deferred RULER 256k re-run (see the program plan). Shipped so far:

• W1 (0036, #293) — PCIe/ReBAR probe consumption. Boot-time pcie_profile reader + --pcie-autodetect/--pcie-bw-gbps; feeds W4's tile sizing. SHIPPED 2026-05-30. Measured proof (#367, 2026-06-04, RTX 3090, perf/llamafile/rebar-probe.cu). Pinned vs pageable host↔device bandwidth (GB/s), pinned/pageable ratio:

  transfer	pinned H2D/D2H	pageable H2D/D2H	pinned/pageable H2D · D2H	ReBAR verdict
  1 MiB	6.397 / 6.361	5.545 / 4.437	1.154× · 1.434×	partial/degraded
  16 MiB	6.656 / 6.590	6.400 / 6.334	1.040× · 1.040×	active
  64 MiB	6.669 / 6.593	6.537 / 6.468	1.020× · 1.019×	active

  Link = x8 @ 8.0 GT/s (PCIe 3.0 x8; iGPU absorbs 8 of 16 lanes), ReBAR active (max x16 @ 16.0 GT/s). The pinned win is size-dependent: large at small, latency-bound transfers (+15% H2D / +43% D2H at 1 MiB), negligible once the x8 link saturates at 64 MiB (~6.67 GB/s, the cached effective_bw_gbps). This validates the #289/M7-A pinned-host KV residency: M7 streams the spill tail in small tiles — exactly the regime where pinned staging earns +15–43%, not the 2% a bandwidth-bound read would suggest. effective_bw_gbps (pinned plateau, 6.67) is the M7 tile-sizing input (tile_bytes_max = compute_ms × eff_bw × 0.8). Raw JSONs: .opencoti/rebar-probe-solidpc-dev0-{1,16,64}MiB-20260604-*.json.

• W2 (0040, #290) — wholesale ik_llama.cpp CPU-FA engine. Vendored behind GGML_IQK_FLASH_ATTENTION + --iqk-flash-attn on|off (default off) + one dispatch hook (opencoti-hook: f5-opt-cpufa). SHIPPED 2026-05-30. Verification: flag-OFF byte-identical (neo-pipeline C4 PASS); flag-ON RULER vt+niah @ 4k/32k (Gemma-4 A4B + M2 + full stack) 100% all 8 cells, ON-vs-OFF wall-time flat. The flat result is the key finding: on prefill-bound RULER cells the CPU-FA kernel speed is not the lever — the CPU split existing at all while VRAM is free is. That motivated the M7 GPU_RESIDENT-by-default / maximize-VRAM redesign (rolling_kv.md Decision 4): start GPU-resident, engage CPU/stream relief only under VRAM pressure, release it as load drains. W2 stays valuable as the CPU-spill fallback kernel. #291 Q8_KV DROPPED 2026-05-30: recon found it needs a full mainline-ggml-type registration (it is the pseudo-type (ggml_type)151 inside the engine, not a real type), not the "nearly-free" flip assumed — and it is a redundant 8-bit option (q8_0 covers it; TurboQuant TBQ3_0/4_0 in M6 is the aggressive-compression KV-quant driver). 0041 retired; the engine's dormant Q8_KV FA path stays gated off at zero cost.

• W3 (0042, #292) — op-level fused MoE up+gate+GLU. SHIPPED 2026-05-30. New ggml op GGML_OP_MOE_FUSED_UP_GATE collapses a MoE FFN's two separate per-expert projections (up_exps @ cur, gate_exps @ cur) + the GLU into ONE decode-time op, dispatched to our existing fused mmvq kernel (ggml_cuda_mul_mat_vec_q + fusion.gate/glu_op) — eliminating the gate_up HBM round-trip + the standalone GLU launch per MoE layer per token. Zero ik_llama code vendored — only the op-level-fusion idea (ik_llama PR #229/#520, MIT). Off by default (--fused-moe-up-gate on). Scope discovery: the hook only fires for separate-up/gate layouts (gate_exps && !gate_up_exps, same type, SILU|GELU, decode n_tokens==1) = Qwen2-MoE / Qwen3-MoE / OLMoE / Mixtral-style. The original target Gemma-4 A4B already fuses gate+up into one ffn_gate_up_exps matmul (gemma4-iswa.cpp:158 passes gate/up = nullptr), so W3 is N/A to Gemma-4 — it gains nothing (only a cheap elementwise GLU remains, <1% ceiling, confirmed by flat OFF/ON tok/s on Gemma-4). CPU backend aborts the op (CUDA-only). Validated on OLMoE-1B-7B Q4_K_M (nsys can't trace llamafile's dlopen'd CUDA → engagement via GGML_SCHED_DEBUG=2 --verbose graph dump: 34 op-nodes on CUDA0 ON / 0 OFF): greedy decode byte-identical OFF vs ON, +2.4% decode tok/s (265.9→272.2). ABI GGML_OP_COUNT 96→97 (binary + ggml-cuda.so rebuilt & paired). Adapter --fused-moe-up-gate after --iqk-flash-attn (87 tests).

• W4 (0070, #296) — M7 Rolling KV — IN FLIGHT (2026-05-31). The residency/transfer redesign the W2 data points to. Shipped on dev in cosine-gated rungs: Rung 0 (M7-A auto-residency — GPU_RESIDENT when KV fits the VRAM budget, spill only under pressure) + Rung 1 (M7-B runtime layer_tactic[] table + rolling_kv_plan, populated/logged, inert on the compute path) + Rung 2 R2-a (the GGML_OP_STREAMING_FLASH_ATTN op + the get_layer_tactic(il)==GPU_STREAM dispatch shell — CUDA forward reuses ggml_cuda_flash_attn_ext, byte-identical to GPU_RESIDENT/CPU_SPILL). Canonical 0070 captured (47 hunks / 19 files, 5 f5-rolling-kv markers; bug-250 patch-chain blocker fixed first, #309). The streaming kernel proper is still unwritten — R2-a is only the op-dispatch shell (n_slots == 0, KV still device-resident, no tile loop). Next: pinned-host residency + slot pool + tile loop + inter-tile online-softmax via flash_attn_combine_results (M7-C/M7-D), then double-buffer overlap (R2-b) + scheduler (R2-c), then M7-F/G/H. See rolling_kv.md "Rung 2 implementation blueprint." bug-226 not yet superseded.

Surgical-hook footprint

• Zero hooks in the opencode source tree from F5.
• Each milestone adds 1 (or 2) new patch files under vendors/patches/llamafile/. The patch protocol header is enforced by vendors/patches/llamafile/README.md.

Open questions

• ReST-KV reference choice. Resolved at M1 start: pick the cleanest retention-scoring paper-or-code combination available at that time. SnapKV and GraphKV are both live candidates if Make-Each-Token-Count remains code-less.
• ScoutAttention publication date. Resolved 2026-05-29. Paper published 2026-03-28; reference code confirmed absent at M4-start. Decision: defer M4 to August 2026 post-DAC. See M4 section above.
• PolyKV ↔ HeadInfer conflict. Research suggests these conflict on shared-pool vs per-head residency assumptions. Resolution path is M6's compose-with-glue step; if the conflict is irreconcilable, document it and let users pick one or the other via flags.

Risks

• Each patch starts against the pinned upstream SHA; later patches modify the same source files. Lexical apply order is enforced by the NNNN- prefix protocol. Gaps in the numbering leave room for per-technique tweaks without renumbering downstream patches.
• M4 (ScoutAttention) has no public reference code. Riskiest milestone. Same escalation path as above.
• Llama.cpp portability scores ≤ 3/5 for all six techniques. Each technique is a port-from-PyTorch-or-vLLM, not a direct drop-in. Estimate scope generously when picking the next milestone after M0.
• Multi-tenant value of PolyKV doesn't land until Tier 0 multi-agent fan-out ships. M6's M5 compose-step assertion is the contract; the bench is only meaningful once concurrency