# F5 — Advanced KV Cache Techniques

> Parent: [../MASTER_PLAN.md](../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](https://github.com/wdlctc/headinfer) |
| NEO | 2/5 | Apache-2.0 | arXiv 2411.01142 — [github.com/NEO-MLSys25/NEO](https://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](https://github.com/ishan1410/PolyKV) |
| LMCache *(deferred)* | 1/5 | Apache-2.0 | arXiv 2510.09665 — [github.com/lmcache/lmcache](https://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](https://github.com/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](https://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`](../../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](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](https://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](https://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`](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`](../../.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`](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
  > 1.