# F5 M7 — Rolling KV: position-windowed KV residency **Status:** SHIPPED (2026-06-04, `dev`). > The design below (locked 2026-05-29) describes the head-axis streaming > double-buffer pipeline. **That approach was superseded mid-flight** — on > Gemma-4's 2-KV-head global layers the head axis can only express a binary > 0/50/100 spill, and the streaming op was inert under attention softcap. The > **shipped** design decomposes residency by **KEY POSITION**: a device-resident > window `[0, wc)` + a pinned-host tail `[wc, n_kv)` on plain flash-attention, > merged by an online-softmax combine (keystone: LSE-emitting `launch_fattn`). > The tactic auto-selects at cache construction (eligibility predicate + > measured PCIe-bandwidth/compute crossover; host-agnostic, nothing hardcoded) > and is driven by **`--kv-residency-mode {auto,head,window}`** (default `auto`). > The dormant CPU_FA_TAIL tactic stays in-tree for sub-x1 PCIe links but is > gated out of auto-select. Acceptance is **logit-distribution equivalence vs > vanilla** (`.opencoti/rolling-kv-equiv-gate.sh`), not greedy byte-equality. > Ships as `0070-rolling-kv` (29 files); the orthogonal cosmocc state-IO fix > (bug-269) ships as `0071-state-io-cosmocc`. See `STATE_SUMMARY.md` ("W4 — M7 > Rolling KV … SHIPPED") and the `0070`/`0071` rows in > `vendors/patches/llamafile/README.md` for the as-shipped record. **Patch slot:** `0070-rolling-kv.patch` (was LMCache, repurposed — LMCache stays deferred under a future slot if it returns). **Prerequisites:** #289 pinned host buffer **(SHIPPED, 0034)**, #290 CPU FA TG threading, #293 rebar / PCIe probe **(SHIPPED, 0036 — runtime consumption + PCIe-link autodetect; see M7-B)**. **Replaces:** the static GPU/CPU partition in F5 M2 (`--headinfer-gpu-heads-frac`) — M2 becomes a runtime tactic the scheduler picks per-layer, not a user-facing flag. **Sister docs:** [`advanced_kv.md`](advanced_kv.md) (the M0-M6 series this builds on). ## Why this exists opencoti's M2 puts a static fraction of attention heads on GPU and the rest on CPU. At Gemma 4 A4B + Q4_K_M + 256k + q8_0 KV that pins ~19 GB on a 24 GB card and leaves 4-5 GB GPU idle for the entire decode while half the KV computes on CPU through a slow CPU-FA path. The card has spare capacity it can't reach because the policy is static, the CPU side never returns to GPU compute, and there is no streaming staging between host and device. The result observed today (RULER vt@262144, F5 stack): 0.0 score, semantic collapse — task tracked separately as bug-226, but the deeper architectural gap is what M7 fixes. The right design for sequential, memory-bandwidth-bound attention is classic double-buffered streaming: hold all KV in pinned host memory, keep a small pool of GPU ping-pong slots, run compute on one slot while DMA loads the next. Both PCIe and SMs are always active. VRAM is always hot. ## Decisions locked (2026-05-29) 1. **Tile axis = composite (PCIe-auto-adapted).** Per-layer would have been the cleanest mapping, but on solidPC (PCIe 3.0 x8, ~6.5 GB/s effective) a single layer's KV at 256k q8_0 is ~530 MB, transferring in ~80 ms while per-layer compute is ~5-15 ms. Per-layer pipeline is PCIe-bound 6-15×. The auto-adapt loop at server boot picks the tile axis: per-layer when compute ≥ transfer; position-window sub-tile otherwise. solidPC will land on position-window; a future PCIe 4.0 x16 host falls into per-layer automatically with no flag changes. 2. **M2 retires as a flag, lives on as a runtime tactic.** `--headinfer-gpu-heads-frac` is removed. The pipeline scheduler picks per-layer between two tactics: - **GPU-stream tactic** (new) — fetch this layer's tile via the ping-pong pool, compute FA on GPU - **CPU-spill tactic** (M2-style) — keep this layer's KV pinned on host, run CPU FA inline User sets `--vram-target` (e.g. `--vram-target 22000` for 22 GB); the scheduler aims to saturate that without exceeding it. 3. **PCIe 3.0 x8 is permanent on this host class.** solidPC has a Ryzen 5600G; the iGPU absorbs 8 lanes and the board caps the discrete slot at gen3. This is common (5600G/5700G class). Auto-adapt is permanent infrastructure, not edge-case handling. 4. **Default tactic is GPU_RESIDENT; relief is load-adaptive and reversible.** *(locked 2026-05-30 — supersedes the boot-immutable tactic model below; motivated by the W2 RULER ON-vs-OFF run, where iqk speeding the CPU-half FA moved wall-time ~0% because the CPU split should not have been engaged at all with 23 GiB VRAM free.)* The engine **maximizes VRAM by default** and only sheds load to slower tactics when it must: - The server starts a fixed pool of N parallel slots (`--parallel N`, e.g. 4/8/…) but allocates KV **dynamically**, never at a fixed CPU/GPU split. - With 1–2 active requests whose KV fits, **every layer is GPU_RESIDENT** — KV lives on the card, zero DMA, zero CPU FA: vanilla's fast path, using all available VRAM natively. - As concurrency rises and the working set approaches the VRAM budget, the engine engages graduated **relief efforts, in order**, only for the layers/streams that no longer fit: (a) tile streaming (ping-pong double-buffer, `GPU_STREAM`), then (b) sub-tile streaming (`GPU_STREAM_SUBTILE`), then (c) CPU/GPU KV split (`CPU_SPILL`, M2-style, consuming the #290 iqk engine). - As requests drain, the engine **transitions back** toward GPU_RESIDENT, reclaiming VRAM for the survivors. Residency is therefore a runtime function of *(free VRAM, context size, active-request count)*, re-evaluated as load changes — **not** a boot-time immutable table. `--vram-target` is a **cap** (leave headroom for siblings), defaulting to "all free VRAM − reserve" (maximize). The fixed `--headinfer-gpu-heads-frac` split is never the default; CPU spill is the last resort, reached only under genuine VRAM pressure. ## Architecture ### Storage layer All per-layer KV lives in **pinned host memory** (`ggml_backend_dev_host_buffer_type(cuda_dev)` → `cudaHostAlloc(cudaHostAllocPortable)`). Pinning unlocks full PCIe bandwidth on async DMA and is the foundation #289 builds. With M7 in place, M2's CPU-half pinning generalizes to *all* KV: the GPU view of each layer is a transient slot, never a permanent home. A small "hot-tile" reserve at the GPU side can hold the most-recently- streamed tiles across tokens (M1-retention-style affinity), avoiding redundant DMA when consecutive tokens both attend to the same layer range. Default reserve = 0 (start simple), tunable. ### Slot pool Boot-time computed: ``` free_vram = nvidia-smi free at server start, minus user --vram-target slack reserve = model weights + compute scratch + cuBLAS workspace (~1-2 GiB) slot_budget = free_vram - reserve - hot_tile_reserve tile_bytes_max = (compute_ms_per_layer × effective_bw_GBps × safety_0.8) n_slots = clamp(slot_budget / tile_bytes_max, 2, 4) ``` Two slots is the minimum for a ping-pong; three lets the scheduler prefetch one ahead of compute; four covers the case where compute and copy are unbalanced enough that staggered prefetch helps. Above four gives diminishing returns since each extra slot just delays host-buffer reuse. ### CUDA stream graph Two dedicated streams per slot pair: - `copy_stream` — issues `cudaMemcpyAsync(slot_i, host_kv_layer_L, ...)` - `compute_stream` — runs FA over `slot_i` Event sync: - `copy_done[i]` signals after copy into slot i — `compute_stream` waits before consuming slot i - `compute_done[i]` signals after FA over slot i — `copy_stream` waits before overwriting slot i Forward pass per token (sketch): ``` for layer L in 0..n_layers: slot = next_slot() if not slot.holds(L): # hot-tile miss copy_stream.memcpy_async(slot, host_kv[L]) # ~tile_ms record(copy_done[slot]) compute_stream.wait(copy_done[slot]) compute_stream.fa(Q[L], slot, V_slot) # ~tile_ms record(compute_done[slot]) # next iteration: prefetch L+1 while compute on L+0 finishes ``` When `tile_bytes_max < per_layer_kv_bytes`, the inner loop runs N position-window sub-tiles per layer with online-softmax accumulation across them. FA's running max + denominator carry between sub-tiles naturally — flash-attention's exact math, just with the K/V chunk source coming from a rotating slot instead of a contiguous tensor. ### M2 as runtime tactic The scheduler holds a `layer_tactic[L]` table picked at server boot: ``` for layer L in 0..n_layers: if total_resident_kv + kv_bytes(L) ≤ slot_budget: # it fits → keep it hot layer_tactic[L] = GPU_RESIDENT # zero DMA, zero CPU — vanilla's fast path else if tile_bytes_max ≥ kv_bytes_per_layer(L) AND has_slot_headroom(): layer_tactic[L] = GPU_STREAM # ping-pong double-buffer else if cpu_fa_throughput(L) > pcie_throughput(L): layer_tactic[L] = CPU_SPILL # M2-style — KV pinned, CPU computes (#290) else: layer_tactic[L] = GPU_STREAM_SUBTILE ``` The table is **re-evaluated as load changes** (active-request count, working-set vs VRAM budget) — *not* immutable for the server lifetime (superseding the original boot-immutable model, per Decision 4, 2026-05-30). With 1–2 requests that fit, every layer is `GPU_RESIDENT` and the engine runs at vanilla speed on all available VRAM; as concurrency grows it demotes the layers/streams that no longer fit `GPU_RESIDENT → GPU_STREAM → GPU_STREAM_SUBTILE → CPU_SPILL` in that order; as requests drain it promotes them back toward `GPU_RESIDENT`. Slot-level adaptivity (which tile in which slot, prefetch ordering) rides underneath the tactic-level transitions. `GPU_RESIDENT` is a first-class tactic (it also absorbs the iswa-window layers noted in Open Questions). CPU-spill consumes the work from #290 (CPU FA TG threading port) and is the **last-resort** tactic, reached only under genuine VRAM pressure. ## Auto-adapt loop (server boot, ~100 ms total) 1. **PCIe topology** — read `/sys/bus/pci/devices//{current,max}_link_{width,speed}`, compute `effective_bw_GBps = width × gtps × 0.8 / 8`. Cache the value for the server's lifetime. 2. **Compute microbench** — synthesize one layer's worth of K/V tensors at f16, run a 1-step `ggml_flash_attn_ext` with `seq_len = target_ctx`, measure wall time on `compute_stream`. Yields `ms_per_layer_at_target_ctx`. 3. **Tile sizing** — apply the formula above. 4. **Tactic table** — populate `layer_tactic[L]` from per-layer KV sizes (Gemma 4 iswa has uniform `n_kv_heads`; trivial. Mixture architectures with variable per-layer KV are the only case where per-L matters). 5. **Log** — print the chosen tile axis, size, n_slots, and tactic histogram (`N_layers_gpu_stream, N_layers_gpu_subtile, N_layers_cpu_spill`) so operators can sanity-check. The probe binary (`perf/llamafile/rebar-probe.sh`) handles step 1's discovery once at install / first-boot; the server reads its cached JSON. The compute microbench (step 2) runs inside the server itself, no external tool. This keeps the boot-time cost bounded and avoids shelling out. ## Phase breakdown ### M7-A — Pinned-host KV residency Consume #289's pinned-host-buffer-type swap + staging pool. Generalize M2's CPU-half buffer to all KV: every layer's K/V starts in pinned host memory regardless of tactic. ### M7-B — Auto-adapt loop Implements PCIe topology read + compute microbench + tile sizing formula + tactic table population at server boot. Logs decisions. > **Step 1 (PCIe topology read) is already shipped** by W1 / patch 0036 > (#293). At server boot `pcie_profile_init()` (`common/pcie-profile.{h,cpp}`) > resolves `effective_bw_gbps` via the 4-tier cascade (manual override → > measured probe JSON → `nvidia-smi` → conservative PCIe3 x8 default) and > logs it. M7-B consumes the resolved `pcie_profile_get().effective_bw_gbps` > directly — no new topology code needed; M7-B adds only the compute > microbench (step 2), tile sizing (step 3), and tactic table (step 4). ### M7-C — Tile streamer + slot pool Allocates `n_slots` GPU buffers of `tile_bytes_max`. Owns the slot rotation policy (round-robin with hot-tile affinity if reserve > 0). ### M7-D — Compute/copy CUDA stream graph Two CUDA streams per slot pair, event-based sync, async DMA enqueue, graph captured into ggml's existing compute backend so FA consumers see no change. ### M7-E — M2 retirement / runtime-tactic scheduler Removes `--headinfer-gpu-heads-frac` from `common/arg.cpp`. Adds `--vram-target` (number, MiB). Plumbs the tactic table through `build_attn_mha` so per-layer kernel dispatch picks GPU stream vs CPU spill from the table, not from a flag. ### M7-F — TS adapter + config schema `packages/opencoti-llamafile/src/config.ts`: drop `headinferGpuHeadsFrac`, add `vramTargetMiB`. `buildServerArgs` emits the new flag. Update adapter tests. ### M7-G — Verification 1. **Cosine ≥ 0.999 unified-mode equivalence** vs pre-M7 binary at small ctx where tile sizing doesn't activate (M5 stack bench unchanged). 2. **256k+q8_0+Gemma 4 quality ≥ vanilla−2pp** — RULER vt@262144, `-ctk q8_0 -ctv q8_0`, ours-streaming vs vanilla. NOTE (2026-06-01, corrected): this is a *vanilla-q8_0 parity* gate, NOT a "bug-226 fix" gate. bug-226's real defect (concat_q block-scrambling) was already fixed by 0035; q8_0 is confirmed coherent on this Gemma at small ctx (vanilla, resident, AND S3d streaming n_head_cpu=4 all = 100% on RULER vt@4k, templated). So this cell tests that the streaming pipeline holds q8_0 parity *at 256k* — the only unconfirmed cell — and is gated behind a 256k f16 canary per the standing "no 256k re-run until M7 lands + canary" instruction. A reference vanilla-q8_0@256k baseline must be captured here (it was never run — the kvtype-split confound). Do NOT treat ours≈vanilla as a tautology: vanilla q8_0 scores ~100% at coherence-checkable ctx. 3. **VRAM target enforced** — server holds within `--vram-target` ± 200 MiB across the run. 4. **Throughput ≥ vanilla at 256k** — at concurrency=1, M7 should at least match vanilla (PCIe-bound floor) and ideally beat it through better VRAM utilization (more GPU compute, less CPU spill). 5. **Hook count audit**: removes M2's hook surface (frac flag); adds M7's vram-target hook. Net delta tracked in `docs/protocols/UPSTREAM_SYNC.md` Registry. 6. **Dynamic residency transition** (Decision 4) — drive a load ramp: 1 request (expect all-`GPU_RESIDENT`, throughput ≈ vanilla, VRAM near `--vram-target`), ramp to `--parallel N` saturation (expect graduated relief: stream → subtile → cpu-spill engages, VRAM held within target), then drain back to 1 (expect promotion back to `GPU_RESIDENT`, throughput recovers). Assert the tactic histogram shifts with load in both directions and never exceeds the VRAM cap. This is the core test of the maximize-VRAM-by-default behavior. ### M7-H — Substance commit + docs ship Snapshot-diff `0070-rolling-kv.patch`. Update `advanced_kv.md` M7 section (was LMCache → now Rolling KV). Move `headinferGpuHeadsFrac` to a deprecated config field with a clear upgrade note. Update STATE_SUMMARY, `.wolf/anatomy.md`, `.wolf/memory.md`. Task #296 → completed. ## Rung 2 implementation blueprint (code-grounded, 2026-05-30) Rungs 0 (M7-A auto-residency) and 1 (M7-B tactic table) have shipped on `dev` (commits `7fe88b8`, `8957fc9`; patch `0070-rolling-kv.patch`). The inert per-layer `headinfer_tactic` table + `rolling_kv_plan` + `get_layer_tactic(il)` are in place. Rung 2 wires the **GPU_STREAM** tactic — and the exploration below pins down what that actually requires before any kernel is written. ### The load-bearing finding: GPU_STREAM is inseparable from online-softmax A "single-tile GPU_STREAM" (copy the whole layer's KV to one slot, run a standard `ggml_flash_attn_ext`, no running-max accumulation) **does not exist as a useful tactic**. GPU_STREAM only has value when a layer's KV *exceeds* what fits resident on the card — if it fit one slot, the correct tactic is GPU_RESIDENT (keep it; don't re-DMA every step). So GPU_STREAM *always* means KV > slot ⇒ multiple tiles ⇒ **cross-tile online-softmax** (running max + denominator in F32). There is no low-risk subset of the streaming win; the precision-critical kernel is the whole deliverable. This is why Rung 2 is gated separately and why the plan flags it HIGH-risk. ### The discrete-GPU linchpin ⇒ a custom op, not graph nodes On the RTX 3090 (non-integrated), CUDA `supports_buft` returns FALSE for pinned-host buffers (`ggml-cuda.cu` ~5144), so if GPU_STREAM were expressed as ordinary graph nodes (KV on a pinned-host CPU buffer feeding `ggml_flash_attn_ext`), the ggml scheduler inserts a **sequential** host→device copy (`ggml-backend.cpp` ~1269) — no overlap, no win over CPU_SPILL. Therefore M7-D must be a **custom streaming-FA op** that owns its own `cudaMemcpyAsync` on a dedicated copy stream. It cannot be plain nodes. ### Integration points (mapped) - **Dispatch site:** `src/llama-graph.cpp:2306`. Today: `if (cparams.neo_pipeline_mode != 0 && mctx_cur->headinfer_split_active(il))` → `build_attn_mha_neo` (two-FA GPU+CPU); else → `build_attn_mha` (single FA). Rung 2 reads `mctx_cur->get_layer_tactic(il)` here and adds a third branch for GPU_STREAM. **Requires** exposing `get_layer_tactic` on `llama_kv_cache_context` (delegating to the cache, like `headinfer_split_active` at `llama-kv-cache.h:514`). - **Stream/event mechanism to extend:** the M3 NEO orchestrator (`ggml/include/ggml-neo-pipeline.h` + the CUDA `graph_compute` hook at `ggml-cuda.cu:4067-4191`) is the exact precedent. It already: registers an op (`register_pair`), routes it to a non-default stream (`curr_stream_no=1`, `cuda_ctx->stream(device,1)`), records a `cudaEvent` (`cudaEventRecord`, 4189), and makes downstream consumers `cudaStreamWaitEvent` on it (4099-4101). M7-D needs the *copy* analogue: a `copy_stream`, `cudaMemcpyAsync(slot[i], host_kv_tile)`, a `copy_done[i]` event the `compute_stream` waits on, and a `compute_done[i]` event so the next tile's copy can reuse the slot — i.e. a per-slot ping-pong of the same `cudaEventRecord`/`cudaStreamWaitEvent` primitives already in the file. - **Bandwidth for tile sizing:** the graph context already holds `cparams` (it reads `cparams.neo_pipeline_mode` at 2306). So `cparams.pcie_bw_gbps` (resolved from `pcie_profile_get().effective_bw_gbps` in `common.cpp:common_context_params_to_llama`, threaded `cparams.h` → `llama.h` → `llama-context.cpp` like `vram_target_mib`) is readable at dispatch — **no kv-cache ctor plumbing needed**, the kernel sizes tiles itself. This 4-file ABI thread lands *with* the kernel, not before (cohesion; nothing reads it until the kernel exists). - **New op:** add `GGML_OP_STREAMING_FLASH_ATTN` to `ggml.h` + a CUDA forward impl (the tile loop + online-softmax) + `supports_op` (CUDA-only; CPU aborts, like 0042's fused-MoE). Slot pool (M7-C) allocates `n_slots` GPU buffers of `tile_bytes_max` in the CUDA backend context, lifecycle paired with the events. ### Decomposition (each its own DSO rebuild + cosine gate) 1. **R2-a — op + slot pool + dispatch, NO overlap.** New op, single copy_stream, sequential tile loop with online-softmax, dispatched from `get_layer_tactic(il)==GPU_STREAM`. Assign GPU_STREAM in `build_rolling_kv_plan` when KV doesn't fit but `--vram-target` allows a slot pool (else CPU_SPILL). Gate: cosine ≥0.999 at small ctx (GPU_STREAM never assigned — inactive), and **correctness of the online-softmax at large ctx** (the precision gate — vs CPU_SPILL output on the same prompt). 2. **R2-b — double-buffer.** Add `copy_done`/`compute_done` events so tile L+1's DMA overlaps tile L's compute. Pure latency; output must stay bit-stable vs R2-a. Gate: same cosine + throughput delta. 3. **R2-c — M7-E scheduler + load ramp.** Promote/demote tactics as `--parallel` load ramps (Decision-4 dynamic residency). Gate: M7-G load ramp + RULER vt@262144 (bug-226) + VRAM cap. The cparams `pcie_bw_gbps` thread + `get_layer_tactic` on the context land in R2-a (smallest cohesive unit that compiles + runs). M7-C's slot pool is folded into R2-a (it has nothing to hold before the op exists — see task #303). ### R2-a actual-state + streaming-kernel refinement (2026-05-31) What **actually shipped** as R2-a (commits `7f73270`, `2489317`; canonical `0070`) is only the **op-dispatch shell**: `GGML_OP_STREAMING_FLASH_ATTN` exists, dispatches at `get_layer_tactic(il)==GPU_STREAM`, and its CUDA forward is `ggml_cuda_flash_attn_ext(ctx, dst)` **verbatim** — `n_slots == 0`, KV still device-resident, **no tile loop, no copy stream, no inter-tile online-softmax.** It is byte-identical to GPU_RESIDENT precisely because it is monolithic FA on a resident tensor; the "online-softmax retired by reuse" insight covers only `launch_fattn`'s **intra-tensor** `parallel_blocks` combine, *not* the inter-tile combine across host-streamed slots. So the precision core is **unwritten**. **The load-bearing architectural finding** (grounded: `build_rolling_kv_plan` `llama-kv-cache.cpp:2197`; dispatch `llama-graph.cpp:2372-2389`): - GPU_STREAM is assigned **iff `layers[j].gpu_heads > 0`** — i.e. exactly when the layer already has an **M2 head-split**, so its spilled heads' K/V are **already in pinned host** (`k_cpu`/`v_cpu`, allocated with the pinned `cpu_buft` at `llama-kv-cache.cpp:461-528`). **There is no whole-layer-host residency reorder** — Rolling KV *rides the existing M2 split*. The ctor-ordering issue (`build_rolling_kv_plan` runs at `:598`, after buffer alloc) is therefore moot for residency: the pinned-host half already exists before the plan runs. - Two combine axes, do not conflate them: - **Head axis (M2):** device head-group vs pinned-host head-group are *independent* outputs → combined by a plain **`ggml_concat`** (exactly what `build_attn_mha_neo` already does, `llama-graph.cpp:2382-2389`). **No online-softmax.** - **Key axis (Rolling KV):** the host head-group's long-context KV is streamed in **key-dimension tiles** to a device slot → FA per tile → merged with **running max + denom in F32** (`flash_attn_combine_results`). **This** is the precision-critical online-softmax; it engages only when the host head-group's KV exceeds one slot (large ctx). One tile ⇒ inactive ⇒ byte-identical. So the real M7-D kernel is **`build_attn_mha_neo` with the CPU-half replaced by an op-owned, key-axis-tiled, async-DMA GPU FA**: stream the spilled head-group's host KV tiles into the slot pool (own `cudaMemcpyAsync` on a `copy_stream`), FA each tile on the `compute_stream`, online-softmax-combine across key tiles, then head-axis-concat with the device head-group's FA. This supersedes the R2-a "feed the full concatenated K/V to one streaming op" framing — that path only ever reproduced monolithic FA (the scheduler's sequential host→device copy of the concat is the very thing M7-D must replace). ### Implementation decomposition, refined (S1 → S2 → S3, each a DSO rebuild + gate) - **S1 — structural split, byte-identical. ✅ DONE 2026-05-31.** Routed GPU_STREAM to the **split** accessors (`get_k_gpu`/`get_v_gpu`/`get_k_cpu`/`get_v_cpu` + `headinfer_gpu_heads`, like NEO), not `get_k`/`get_v`. `build_attn_mha_streaming` restructured to the NEO shape: device head-group via resident `ggml_flash_attn_ext` + host head-group via the `GGML_OP_STREAMING_FLASH_ATTN` op (forward still plain FA at S1) + head-axis `ggml_concat`; no NEO pair-registration. **Slot pool + op-owned DMA deferred to S2** — they have no correctness purpose at a single tile and belong with key-tiling, which keeps S2 as pure precision math. Binary-only `:make` (op forward untouched ⇒ no DSO rebuild). Gate **PASS**: GPU_STREAM (`OPENCOTI_M7_STREAMING=1`) ≡ NEO byte-identical greedy decode on Qwen2.5-1.5B at frac 0.5 (`GPU_STREAM=28` layers engaged; A==B==C all byte-identical). Owning the path surfaced + fixed **bug-167**: the `--flash-attn auto` resolver (`llama-context.cpp`) `GGML_ASSERT`ed on the M2/NEO/M7 split FA naming `__fattn___g-` (latent since M3 NEO) → rewritten to skip non-plain FA ops. - **S2 — key-axis tile loop + inter-tile online-softmax. ✅ DONE 2026-05-31 (byte-identical).** The open sub-design was resolved by a hard CUDA finding: the **mma kernel our 3090 selects does NOT use `flash_attn_combine_results`** — it normalises in-kernel and combines via the `stream_k` `dstk_fixup` path (only tile/vec/wmma use the `parallel_blocks`/`dst_tmp_meta` machinery; fattn.cu dispatches head_dim=128 to MMA on any `turing_mma_available` GPU, VEC fast-path gated `cc ≥ ADA_LOVELACE`). So extending `launch_fattn` to reuse its combine was invalid for our path. **Resolution (user-approved): LSE-weighted host-tile combine.** Per key-tile, run the STOCK `ggml_cuda_flash_attn_ext` (untouched — resident path byte-identical, every kernel family covered) for normalised `O_t`, plus a small `streaming_lse_kernel` for per-row `lse_t = m_t + log d_t`; merge in F32 via `O = Σ_t O_t·exp(lse_t − M)/Σ_t exp(lse_t − M)`, `M = max_t lse_t` (exact online-softmax). Single-tile short-circuits to plain FA (= S1). Lives in `fattn.cu` (`ggml_cuda_streaming_flash_attn` + 2 kernels), dispatched from `ggml-cuda.cu`; env `OPENCOTI_M7_STREAM_TILES` (default 1). Gate `.opencoti/s2-gate.sh`: T2(2-tile)/T4(4-tile) **byte-identical** to NEO ref at a 1969-token prompt (n_blk=8, genuine cross-tile), needle recalled, reap clean — stronger than the cosine≥0.999 bar. Speed: zero penalty on any working path (single-tile short-circuits; multi-tile only runs where resident can't hold KV; the LSE pass reuses the already-streamed slot K → no extra DMA). NB: scoring coverage is scale+mask only (max_bias==0, softcap==0, no sinks) → resident fallback otherwise; gate model qwen2.5-1.5b satisfies all three. - **S3 (= blueprint R2-b) — double-buffer overlap.** Staged S3a → S3b → S3c so correctness is gated before throughput, throughput before the memory win. - **S3a — slot-pool machinery (DONE, 2026-05-31, byte-stable).** Per-key-tile round-robin GPU slot pool (`OPENCOTI_M7_STREAM_SLOTS=N`); each tile's K/V is lifted into a packed slot and FA reads the slot, proving the slot path is a numeric no-op vs S2. **Single-stream by design**: copies issue inline as `cudaMemcpyAsync(…, cudaMemcpyDefault, ctx.stream())` on the op's assigned stream — NO dedicated copy_stream, NO events. This sidesteps the M2 `concurrent_events` stream collision (bug-253): at frac<1.0 under eager exec the head-split forks work onto streams 1..n, and a hardcoded second stream raced it, scrambling output; `ctx.stream()` auto-rides whatever stream the op was forked onto. Zero events ⇒ CUDA-graph-capturable ⇒ the bug-251 graph-disable was reverted. Gate `.opencoti/s3a-gate.sh`: T2s1/T2s2/T4s2 all **byte-identical** to T2(S2)==REF(NEO), needle recalled in all 5, reap clean. - **S3b — double-buffer overlap (DONE, 2026-05-31, byte-stable).** Dedicated copy_stream at fixed index `GGML_CUDA_MAX_STREAMS-1` (=7) carries tile L+1's K/V DMA while FA/lse for tile L run on `ctx.stream()`, gated by a pre-created `copy_done[]`/`compute_done[]` event pool (created ONCE with `cudaEventDisableTiming`, never during capture — record/wait ARE legal during capture, so NO graph-disable needed; the audit corrected bug-251 here). Index 7 avoids the QKV `concurrent_events` fan-out (streams 1-3, which JOIN before our op runs) and NEO's stream 1 — hardcoding stream 1 was bug-253. Two latent bugs surfaced and were fixed during S3b: **bug-255** (cross-op WAR hazard — slot pool buffers reused across layers; copy_stream not ordered vs the compute stream across op invocations → layer L+1's copy clobbered a slot layer L's FA still read → prefill garbage; fixed with a per-op-entry `cs_sync` event ordering copy_stream behind the compute stream) and **bug-254** (the shared `streaming_lse_kernel` hard-casts K to `half*`; an S3b-0 stop-gap f16 guard keeps any non-f16 K resident until S3d's dequant-on-lift lands — S3b stays **f16-pure**). Gate `.opencoti/s3b-gate.sh`: B2s2 (overlap) and B4s2 (overlap + slot-reuse + prefetch) **byte-identical** to S3a==REF(NEO), needle recalled; the `--parallel 2` P2 collision-stress is SOFT (warn-only) — its collision axis is already covered by B2s2/B4s2 byte-identity, and its boot- fail is a pre-existing llamafile graph-node-pool ceiling under `--parallel` (bug-256), gated separately at M7-G. Regression shield green: NEO 4/4 hard, advanced-kv-stack 8/9 (sole fail = pre-existing non-S3b C5, task #280). - **S3c — scheduler bypass (the memory win, ✅ DONE 2026-05-31, byte-stable).** On a discrete GPU the ggml scheduler can't run a `cuda_host` buffer as an op input (`ggml_backend_cuda_device_supports_buft` → false), so it auto-inserts a full sequential host→device copy of the pinned CPU-half KV before the streaming op every step (`ggml-backend.cpp`), then rewrites the op's `src` to the device dup — no memory win, and the S3a/S3b lift was device→device. S3c adds an **op-scoped, env-gated** suppression of that copy (`OPENCOTI_M7_STREAM_BYPASS=1`, keyed on `GGML_OP_STREAMING_FLASH_ATTN` + K/V src + host buffer) so the op's `src` stays the pinned-host **permuted** view, and `fattn.cu` detects the host pointer (`cudaPointerGetAttributes`), forces the slot path, and lifts each tile with a **strided `cudaMemcpy2DAsync`** (`spitch = nb[1] = n_head_cpu·row`, since the permuted host view's key rows are NOT contiguous, unlike the scheduler's device dup). Device srcs (bypass=0) keep the exact 1-D lift → byte-identical bisection lever. **The streaming op never drifts to CPU** (CUDA-only op type). Gate (`.opencoti/s3c-gate.sh`): S3a==REF, B2s2==S3a, B4s2==S3a, **S3c==B4s2 byte-identical**, and the **no-copy-node proof** — under bypass=1 the scheduler copy is suppressed for every layer's K/V (instrumentation count >0, all `host=1`), under bypass=0 never (==0). NOTE: the CUDA0 compute-buffer PEAK is insensitive to the suppressed ~CPU-half MiB at 4k (FA softmax scratch dominates + allocator reuse — bug-257), so the steady-state VRAM win is measured at 256k in M7-G, not at this gate. Regression shield: neo-pipeline 4/4, advanced-kv-stack 8/9 (sole fail = owned non-S3c C5, task #280 — the bench enables streaming on zero configs, so the entire S3c path is structurally inert there). Still f16 (the bug-254 guard stays; dequant is S3d). - **S3d — dequant-on-lift. ✅ DONE 2026-06-01, validated incl. n_head_cpu>1.** The slot is now ALWAYS f16: a quantized K/V is dequantized into the slot on the lift via `ggml_get_to_fp16_nc_cuda` (the exact primitive fattn-common.cuh uses; strides in elements = nb/type_size, packed f16 output), keeping `streaming_lse_kernel` type-agnostic and removing the bug-254 f16 stop-gap. Two safety guards: an unsupported quant type (no nc converter) and a quant type with no slot both stay resident. Gate (`.opencoti/s3d-gate.sh`, Qwen2.5-1.5B): f16 anchors byte-identical (S3a==REF, B2s2==S3a, B4s2==S3a, S3c==B4s2 — refactor inert on f16), NOCOPY (896/0), **DEQUANT byte-identical to resident-q8_0** (streaming q8_0 == resident q8_0). Council audit (7-agent) cleared the dequant stride math against the stock resident reference. **n_head_cpu>1 validated** on Gemma-4 A4B (8 KV heads → frac 0.5 = 4 GPU/4 CPU on 25/30 layers, all streaming): RULER vt@4k q8_0 = **100%**, equal to resident-q8_0 ground truth and vanilla — so the latent-stride fix (`Kt.nb[1]=k_row`, lse `k_nb1=k_row`) is load-bearing and correct, not inert (`.opencoti/s3d-gemma-nhead-validate.sh`). **bug-226 reframed:** the real opencoti defect was the concat_q block-scrambling, already fixed by 0035; q8_0 works on this Gemma at coherence-checkable ctx (vanilla, resident, AND S3d streaming all = 100% templated). The earlier "vanilla q8_0 collapse" was an off-template (raw `/completion`) confound, retracted. The only open cell is **256k q8_0 quality → M7-G** (deferred per the post-M7 + 256k-canary instruction); S3d is NOT blocked on it. - **E1 — stream-by-default auto-engage. ✅ DONE 2026-06-01 (E1a+E1b+E1c).** Turns the dev-env-gated S-ladder into the product's auto-selected relief tactic. **E1a (host):** `OPENCOTI_M7_STREAMING` is now tri-state — `"1"`→force GPU_STREAM, `"0"`→force CPU_SPILL, **unset → GPU_STREAM** (stream-by-default for any spilled layer); `frac=1.0` → no spill → all GPU_RESIDENT → vanilla byte-identical (flip inert). **E1b (CUDA DSO + host):** the streaming op self-configures — tile count `n_tiles=clamp(ceil(n_kv/2048),1,n_blk)` auto from n_kv, `slots=2` default, overlap default-on at slots≥2, S3c scheduler-copy bypass default-on (host-buffer-guarded). All knobs keep their env override as bisection levers. **E1c (adapter):** `launch.ts` default-emits `--headinfer-gpu-heads-frac auto` + `--vram-target`; config gains `vramTargetMiB` (#332). Gate (`.opencoti/m7e-b-gate.sh`, Qwen2.5-1.5B): RESIDENT==VANILLA + KAUTO==BASE + **FIDELITY AUTO==RESIDENT byte-identical** (the auto streaming faithfully reproduces the no-split reference — the pathological-prompt "degeneration" is correct, matching vanilla's own filler-echo; the old "AUTO must diverge" gate premise was retracted) + MULTI-TILE-ENGAGED (auto n_tiles up to 4) + DETERMINISM + NOCOPY (1400/0) + needles incl. q8_0. Shields held: advanced-kv 8/9 (C9 = the #287 M2 slowdown M7 exists to fix, not a regression), neo 3/4 (C2 = M3 "no observable win" perf-noise; C4 off-path byte-identical passes). **Gemma-4 A4B n_head_cpu=4 auto path** (`.opencoti/m7e-gemma-auto.sh`): RULER vt@4k q8_0 = **100%**, equal to vanilla + resident — the full-auto product path is coherent on the bug-226-class multi-CPU-head split. bug-259 (streaming-op warmup/-fit abort under stream-by-default) fixed: CUDA `supports_op` for STREAMING_FLASH_ATTN validates head_dim only (the op repacks K/V to contiguous f16 slots, so raw-src stride-sensitivity was wrong); host-side CPU-supports-op=false + pass-4 CUDA pin kept as inert defense-in-depth. **E2 (dynamic load-ramp under `--parallel`) is next** — deferred behind the M7-G 256k canary (needs runtime active-occupancy infra the codebase lacks; const-after-ctor `layer_tactic[]`). ## Open questions (resolution path during impl) - **Hot-tile reserve default** — does keeping 1-2 recently-used tiles GPU-resident across tokens give measurable gain, or is the DMA cost already negligible vs compute under the auto-tune? Decide at M7-C first bench. - **Sub-tile online-softmax precision** — when running N position-window sub-tiles per layer, the running max + denominator accumulate in F32. Verify f32 accumulator precision holds at 256k (probably fine; FA already does this internally). - **iswa interaction** — Gemma 4 uses interleaved sliding-window attention. Some layers attend to a small recent window only. Those layers may not need streaming at all (the working set fits in one slot). Tactic table should detect and route iswa-window layers to a trivial `GPU_RESIDENT` tactic — a third tactic alongside `GPU_STREAM` and `CPU_SPILL`. - **PolyKV M6 composition** — if M6 ships first, M7 streams the *compressed* tiles, decompressing in-slot before FA. Tile size math changes (smaller bytes per slot, same compute time → bigger effective tiles). M7's auto-adapt loop becomes `tile_bytes_max × compression_ratio`. Plan to land M7 first if bug-226 forces it; otherwise M6 → M7 is natural. ## Risks - **Online-softmax across sub-tiles** is the trickiest math. FA already does it within a kernel; we extend it across multiple kernel invocations using the saved max + denom. One precision bug here destroys output quality. Mitigation: keep the M0/M1/M2/M3/M5 regression-shield benches; gate ship on cosine ≥ 0.999 at small ctx where the sub-tile path doesn't activate. - **Async DMA + ggml scheduler** integration — ggml's existing scheduler doesn't expect "this tensor's data arrives mid-graph from another stream." Need to register events as scheduler dependencies or pre-stage at graph-build time. Risk: race conditions invisible in single-tenant tests but visible under `--parallel`. - **Probe + boot ordering** — the auto-adapt loop reads PCIe topology from the cached probe JSON. If the probe hasn't run (fresh install), fall back to a `current_link_speed` read inline; if that fails, use a conservative tile size (1 MiB) and warn. Never fail boot on topology read. - **Re-tuning when VRAM availability changes mid-run** — if a sibling process grabs VRAM, M7's slot allocations may need to shrink. v1 doesn't handle this; the existing `--vram-target` assumption is static. Document as a known limitation. - **M2 flag removal is a user-visible break.** Anyone scripting `--headinfer-gpu-heads-frac` will break. Ship a deprecation cycle: M7 lands accepting both flags (frac → vram-target translation), a release later the old flag warns, a release after that it errors. ## Why land this above PolyKV M6 M7 is the foundation for "KV doesn't have to fit on GPU." PolyKV adds compression to that. Without M7, PolyKV's compressed pool still has to fit on GPU; with M7, even an uncompressed pool can stream arbitrarily large KV. **M7 first lets PolyKV's compression be pure gravy on top of unlocked KV capacity, rather than a hard requirement to fit.** Sequence: bug-226 root cause → M7 design → M7 impl → PolyKV M6 (composed atop M7). Adjust if bug-226 turns out to need M7 urgently; otherwise the optimization round (#287, tasks #289-#293) ships first as M7 prerequisites. ## S0 cliff-mechanism verdict (2026-07-01) — graphs are NOT the cause; redirect #586/#587/#588 Before building the pending rolling-KV perf work (#586 shared staging ring / #587 cross-layer run-ahead / #588 perf gate), an S0 investigation nailed **what actually causes the overflow decode cliff** (resident ~30–46 tps → spill ~0.3–1.6 tps). The working hypothesis had been "a host-sourced tail tile disables CUDA-graph capture for the whole decode step → all `n_kv/32768` streaming tiles run as eager host-issued launches." **That hypothesis is REFUTED.** **Method:** the binary already emits a native per-slot metric `graphs reused = N` (no rebuild needed — this is the "read the target's native diagnostics first" lesson). A/B on the 3090 (Qwen3-8B-Q8_0, ctx 16384, ~15k-token prompt so the spill genuinely pushes occupied cells into the tail): | config | split | **graphs reused** | decode tps | | --- | --- | --- | --- | | resident (vt 20000) | FULLY RESIDENT | **23** | 41.99 | | spill (vt 2000) | 3072 resident / 13312 tail (1872 MiB) | **23** | 1.56 (27× collapse) | CUDA graphs are captured and **replayed 23× in BOTH cases** — the graph is *not* disabled by a host tail. So the ~640 ms/token penalty happens **inside the replayed graph**: it is real streaming-path *work*, not launch/eager overhead. **Mechanism (grounded):** a host tail flips the *entire* attention — including the device-resident window — from the fast dense-FA path onto the slow per-tile streaming path (`two_region` re-tiles window *and* tail, fattn.cu). Each tile pays a strided per-head copy (`cudaMemcpy2DAsync` ×`n_head_kv`×2) + the `cs_sync` barrier (fattn.cu:~1375) which serializes the copy stream against all prior compute — so tiles execute serially even under graph replay. Cost ∝ tile count (matches the earlier 2-tile 632 ms vs 7-tile 3146 ms scaling), and the device-resident window pays the streaming tax it would NOT pay when fully resident. **Redirect (the levers are structural, none is "graph capture"):** - **#586 shared staging ring** — collapse the per-head `n_head_kv×2` strided copies into one bulk transfer (cuts node/copy count). - **#587 cross-layer run-ahead** — drop/loosen the per-op `cs_sync` barrier so the tail DMA overlaps compute instead of serializing (was mis-framed as "restore graph capture"; the real target is the barrier). - **structural** — when only a tail spills, run the device-resident window as ONE dense FA (the same fast kernel the fully-resident case uses) emitting its lse, tile ONLY the tail, and do one online-softmax combine. Removes the streaming tax from the window (the bulk of the KV); benefit ∝ window/tail ratio (largest when occupied barely exceeds VRAM — the real serving regime; deep overflow stays bandwidth-bound). - **#588 perf gate** — re-run the 3090 cliff curve; success = spill decode within ~10–15% of resident at small-tail (occupied ≈ VRAM) configs. See bug-1342 (0089, the boot OOM that unblocked this curve) and bug-2094 (0090, the `--parallel ≥2` multi-session boot fix found while probing tail amortization). ### Multi-session amortization (2026-07-01) — the cliff does NOT amortize under --parallel With bug-2094 (0090) unblocking `--parallel ≥2` boot, the original question — does batching hide the tail penalty? — was answered on bs2 (RTX PRO 6000, 14B-1M-Q8_0, controlled per-slot window: p2 vt=2×p1 vt so each slot's window ≈ equal): | config | host tail | per-slot decode tps | aggregate | | --- | --- | --- | --- | | p1 (1 session, vt 8000) | 31232 cells | 1.08 | **1.08** | | p2 (2 sessions, vt 16000) | 27136 cells/slot | 0.28 + 0.66 | **0.94** | Two concurrent spilled sessions aggregate to **less** than one (0.94 < 1.08 tps) — batching does not recover throughput; the sessions **contend on the same serialized streaming path** (shared copy stream + the per-op `cs_sync` barrier), and per-slot tps is asymmetric (0.28 vs 0.66) from uneven interleave on that one copy stream. **Consequence:** for the multi-session serving target ("more users per card"), spilling collapses per-card throughput regardless of concurrency, so the #586/#587 streaming-path rework (kill the serialization) is **critical**, not a nice-to-have. ## 2026-07-05 — bug-1843 shipped (patch 0098): the spill-decode recompute is dead The S0-cliff redirect (#586 ring / #587 barrier-drop) proved perf-inert (bug-1838); the real lever was the **per-tile `streaming_lse_kernel` recompute** that every `head_dim<=256` decode tile paid because `decode_lse` was gated on `head_dim > 256` (bug-1843, ~80% of spill-decode wall — 1.14 vs 6.05 tps NOOP A/B, Qwen3-8B ctx40960/vt14000, 3090). Patch `0098-rolling-kv-lse-decode` arms the Stage-3a LSE channel for D≤256 at all three streaming-FA sites and recomputes only when `opencoti_fattn_dst_lse_written` says no finalize ran. Shipped gates (3090, DSO df46dc36): spill decode **1.14 → 6.04 tps (5.3×, == NOOP ceiling)** with needle PRESENT; resident 42.5 tps unchanged; `OPENCOTI_LSE_NOOP` now speed-inert (6.02 — recompute structurally off the path). Correctness: teacher-forced logit-equiv (Qwen3-8B D128, ours-resident REF vs POSITION_WINDOW at 256/16384 cells resident): real_frac=0.0, frac_full_agree=1.0, mean_tv 0.0068 → PASS. Follow-up: #588 re-runs the 3090 overflow→tps cliff curve on the fixed binary (go/no-go for further rolling-KV investment). Note bug-2115 while gating: bs2-built DSOs do NOT load on solidPC (glibc ≥2.38 vs Debian 11) — build the 3090 DSO locally or the run silently falls back to CPU. ## 2026-07-06 — #588 perf gate: post-fix cliff curve — spill decode is at the PCIe floor (rolling-KV perf line CLOSED) Curve on the bug-1843-fixed DSO (`df46dc36`, 3090, Qwen3-8B-Q8_0, ctx 16384, ~13.1k-token prompt = 2304 MiB f16 KV, needle PRESENT in every cell): | vt (MiB) | window split | occupied tail | decode tps | ms/token | | --- | --- | --- | --- | --- | | 22000/16000/14000/13000 | FULLY RESIDENT | 0 | 42.3–42.5 | 23.6 | | 12000 | 11008 / 16384 cells | ~2.1k cells ≈ 295 MiB | 15.5 | 64.6 | | 8000/4000/2000 | 256 / 16384 (floor) | ~12.9k cells ≈ 1815 MiB | 3.33 | 300 | **The PCIe model now fits within ~5%:** `t(token) ≈ 23.6 ms + occupied_tail_bytes / 6.5 GB/s` (vt12000 predicted 67 ms vs 64.6 measured; floor predicted 303 ms vs 300). The S0 anchor cell (vt 2000) improved 1.56 → 3.33 tps and the former compute-side ceiling (streaming_lse recompute) is gone — spill decode is now purely bandwidth-bound, i.e. at the structural floor for a design that re-streams the tail every token. **Gate verdict: NO-GO for further re-stream perf work.** The "within 10–15% of resident at small tails" bar is only reachable for tails ≤ ~25 MiB (~170 cells) on a 6.5 GB/s 3090 (~8× more on bs2's 50 GB/s link) — no copy-plumbing lever (#586-class) can beat the link itself. What remains valuable: - **Window mode as graceful-overflow fallback** (shipping default): linear degradation ∝ spilled bytes, correctness clean (needle + logit-equiv real_frac=0), no cliff pathology left. - **The real capacity levers are residency-side**, per #582: quant-KV decode reclaim (#620), auto KV-tier boot policy (#621), and position-axis mixed KV with a *compressed VRAM tail* (#622) — which avoids the per-token re-stream entirely instead of optimizing it. ## 2026-07-06 — bs2 tail-spill curve (the real-bandwidth one): bug-2116 bulk-H2D unlocks the 50 GB/s link The #588 verdict above is on the **3090's 6.5 GB/s link** — where the tail is so bandwidth-starved that no plumbing lever can help. bs2 (RTX PRO 6000, ~50 GB/s link) is the host where window-mode spill is actually *usable*, and it is also the host where the **bug-2116 bulk-H2D fix** (batched staging vs per-tile serial `cudaMemcpyAsync`) matters most. Curve on **Qwen2.5-14B-Instruct-1M Q8_0**, ctx 131072, ~106k-token prompt (KV 9984 MiB f16-resident), needle PRESENT in every cell. Both curves share the same VT→window-split mapping, so they are directly comparable against the real host-tail byte count: | host tail | BEFORE (per-tile H2D) | AFTER (bug-2116 bulk-H2D) | | --- | --- | --- | | 0 (fully resident) | 15.88 tps / 63.0 ms | 15.91 tps / 62.85 ms | | 1443 MiB | 13.19 / 75.8 | 13.17 / 75.95 | | 2028 MiB | **9.05** / 110 | **12.23** / 81.79 | | 2828 MiB | **3.96** / 253 | **10.28** / 97.31 | | 7624 MiB | (already collapsed by 2.8 GiB) | 5.53 / 180.84 | **The fix converts launch-bound into link-bound.** BEFORE, the deep-tail slope is ~5.5 GB/s effective (2028→2828 MiB costs +142 ms) — i.e. even on bs2's 50 GB/s hardware the *serial per-tile launches* throttle spill to 3090-class bandwidth, and the curve cliffs at ~2 GiB. AFTER, the deep slope is **~56 GB/s** (2828→7624 MiB costs +83.5 ms for 4.68 GB) — matching bs2's physical link — so at a 2.8 GiB tail you keep **10.28 tps = 65% of resident** (vs 3.96 = 25% before), and a 7.6 GiB tail still decodes at 5.53 tps. This is the measured confirmation of the "~8× on bs2's link" projection in the #588 section, and it makes window mode a genuinely usable graceful-overflow tier on the production host, not just a non-cliff fallback. Data: `bs2:/srv/ml/opencoti-c1/tailcurve-out/` (BEFORE) and `tailcurve-after-out/` (AFTER); decode tps read from `slot print_timing … eval time` in the per-cell `tailN.log` server logs (the AFTER driver never aggregated a `run.log` beyond tail=200 — the numbers above are recovered from the server logs). **Two AFTER cells are truncated** (host tail 1638 MiB / 4426 MiB — run cut off; the log has prompt-processing but no eval line); they will be re-run to complete the curve when GPU1 frees from the bug-2121 window-mode correctness gate (which exercises this same spill path). The AFTER slope is already pinned by the 2828→7624 MiB segment, so the verdict does not depend on the two missing points. ### 2026-07-06 — multi-model spill sweep: production KV is compact at 128k (spill barely triggers) Extending the bs2 curve to four production targets (Gemma-4-31B Q6_K, Gemma-4-A4B-128e Q4_K_M, Qwen3.6-27B-Omnimerge Q6_K, Qwen3.6-35B-A3B Q6_K), q8_0-K/q4_0-V, ctx 131072, surfaced a physical finding **before** any bandwidth number: at 128k these models simply do not have enough KV to spill. Measured resident KV (fully on-GPU) and fully-resident decode tps: | model | attention | resident KV @131072 | resident decode tps | | --- | --- | --- | --- | | Qwen2.5-14B-1M Q8_0 | full | **9584 MiB** | 15.9 | | Qwen3.6-27B-Omnimerge Q6_K | full | 3328 MiB | 29.8 | | Qwen3.6-35B-A3B Q6_K | full (MoE) | 1040 MiB | 89.4 | | Gemma-4-31B Q6_K | **iSWA** | **488 MiB** (global only) | 25.6 | | Gemma-4-A4B-128e Q4_K_M | **iSWA** | (small global) | — | Two structural reasons the production models barely spill at 128k, both of which the 14B-1M (the original curve) side-steps: **(1) GQA compactness** — the Qwen production quants have few KV heads, so full-attention KV is only 1–3.3 GB at 128k vs the 14B-1M's 9.6 GB (Q8, more KV heads); **(2) iSWA** — Gemma-4 (both A4B and 31B) keeps only ~1/6 of layers global, so the cache that *can* spill is tiny (31B global = 488 MiB). On a 96 GB card none of these overflow at 128k, so a natural spill only appears at far longer context (256k–1M for the Qwen full-attn models; iSWA Gemma effectively never spills its global cache at usable context). A follow-up sweep (`tailcurve-multi2`) forced spill by capping `--vram-target` below (weights + KV), weights measured from `nvidia-smi`, tail ladder = 0/25/50/75/90 % of each model's KV. Results (bs2, DSO `8b23afbf`, q8_0-K/q4_0-V, ctx 131072, needle PRESENT in **every** cell): **14B-1M — the one clean forced-spill curve** (full-attn, 9584 MiB KV): | host tail | decode tps | ms/tok | | --- | --- | --- | | 1019 MiB | 13.16 | 75.97 | | 3515 MiB | 9.36 | 106.87 | | 6011 MiB | 6.46 | 154.83 | | 8507 MiB | 5.03 | 198.80 | | 9984 MiB (all KV on host) | 4.35 | 229.62 | Deep slope (6011→9984 MiB) = **51.9 GB/s**, i.e. the bs2 host link — the same result as the `tailcurve-after` curve above, reproduced by an independent forced-spill trigger. Window mode degrades linearly and holds retrieval to the fully-spilled floor. **The four production targets do not give a spill curve at 128k**, for two distinct reasons: - **iSWA (Gemma-31B, A4B):** the *global* cache — the only part window mode spills — is just **1536 cells (~5 MiB)**. Every forced-spill cell shows `256 / 1536 resident`, decode stays flat (31B ~21 tps, A4B ~38 tps), needle PRESENT. Spill is a non-event; the sliding-window layers never leave VRAM. **iSWA models effectively cannot be made to spill** at any usable context. - **Compact-GQA full-attn (Qwen-27B 3328 MiB, 35B-A3B 1040 MiB):** KV fits so easily that the calibrated low-VT still didn't drop below (weights+KV) — all cells stayed FULLY RESIDENT (flat 29.8 / 89.4 tps). Forcing them to spill would need a hardcoded sub-weights VT; at 128k they simply have no overflow. A genuine spill on these needs 256k–512k context (bigger KV). (A first attempt, `tailcurve-multi`, was discarded — its calibration read the high-VT budget as free-VRAM → negative targets → BOOT-FAILs; only the resident-tps column above survived it.) **Bottom line for the four requested models:** at 128k only the 14B-1M (large Q8 KV) spills, and it hits the ~52 GB/s link exactly like the reference curve. The production targets either *can't* spill (iSWA global cache is ~5 MiB) or *don't* (compact GQA KV fits) — so window mode is a no-op safety-net for them at this context, and real spill characterization would require 256k–512k runs. ## 2026-07-12 — #582-P2 SHIPPED: position-axis mixed KV (compressed tail, `-ctkt`/`-ctvt`) The window-mode spill above re-streams the tail every token at *the boot KV type*, so a big Q8 tail is both slow (bytes) and redundant (older tokens are attention-light). **P2 makes the spilled tail carry a MORE-compressed type than the resident window** — the position-axis complement to P0 (head-axis, #620/#643) and P1 (boot-tier auto-select, #621). Recent tokens in the window `[0,wc)` keep their high-fidelity type on-device; the tail `[wc,n_kv)` on pinned host is quantized further, so each decode step ships fewer H2D bytes without touching the part of the cache attention actually leans on. **Interface.** Two new CLI flags set the *tail* K/V type independently of the window: `-ctkt ` / `-ctvt ` (`--cache-type-k-tail` / `--cache-type-v-tail`). Unset → sentinel `GGML_TYPE_COUNT` → "tail == window", and every code path is byte-identical to the shipped uniform-window behaviour (regression leg of the gate, `REG==U8` proven byte-identical). Typical use: `-ctk q8_0 -ctv q8_0 -ctkt q4_0 -ctvt q4_0` (q8 window ⊕ q4 tail). **Where the win lands (decode).** The POSITION_WINDOW streaming FA op (`fattn.cu`) already reads window and tail *each in-register at its own type* (the #620/#629 reader path) — there is no whole-cache f16 materialise, so a q4 tail is genuinely ~half the H2D bytes of a q8 tail. **Two correctness-fallback fixes** were needed for the paths that *don't* go through the in-register decode op (both host-only, both byte-identical when tail type == window type): - **bug-2161** — prefill / graph-reserve (n_q > `WS2_NQ_MAX`) reassembles window+tail via `ggml_concat`, which asserts `a->type == b->type`. Fixed with a file-local `poswin_lift_to_f16` (mirror of `dca_lift_to_f16`: scalar quants → `ggml_cast` F32 → F16, byte-exact) applied to both region views only when their types differ. - **bug-2162** — an in-memory **context checkpoint** (llama.cpp PR16391, on by default) hits the Stage 3c-6 state-IO window paths, which had assumed the tail row size equals the window row size — so a mixed config asserted on the 2nd prompt. Fixed with a per-region `ggml_row_size(k_cpu->type / v_cpu->type, …)`, symmetric across writer and reader. **Gate (S7, bs2 Qwen2.5-14B-Instruct-1M-Q8_0, `-c 24576` `--vram-target 18000`, ~21k-token prompts → wc ≈ 13312 resident, ~11264-cell host tail; logit-equiv vs the model's own full-resident gold — NEVER greedy needle):** | config | tail type | correctness (real_frac) | advisory mean-TV vs gold | decode tps | | --- | --- | --- | --- | --- | | U8 (uniform q8 tail) | q8_0 | 0.0 | 0.0069 | 26.4 | | **MIX (q8 window ⊕ q4 tail)** | q4_0 | **0.0** | **0.0110** | **28.1** | | U4 (uniform q4) | q4_0 | 0.0 | 0.0118 | — | Mixed sits **strictly between** uniform-q8 and uniform-q4 on fidelity (closer to gold than uniform-q4) at **+6.4% decode tps** over uniform-q8 — the tail's fewer H2D bytes convert directly to throughput, and the recent window keeps quality above the uniform-q4 floor. The gate exercises checkpoints, so it doubles as a bug-2162 checkpoint round-trip test. Shipped as patch `0120-p2-mixed-kv` (marker `opencoti-hook: P2 mixed-kv`, 12 nested-llama.cpp files, additive; chain state → 66 patches). See UPSTREAM_SYNC.md registry entry for the full 5-point plumbing surface. ### 2026-07-12 — #653 P2 productized: TS adapter flags + P2-auto boot policy Two follow-ons turn the raw `-ctkt`/`-ctvt` mechanism into a first-class, opt-in opencoti feature. **Part 1 — adapter (`packages/opencoti-llamafile`).** The mixed-tail flags are now typed config fields, so a caller sets them like any other KV knob instead of hand-passing `extraArgs`: - `config.ts` — `ctypeKTail` / `ctypeVTail: string | undefined` (env `OPENCOTI_LLAMAFILE_CTYPE_K_TAIL` / `_V_TAIL`), defaulting to `undefined` = server default = "tail == window". - `launch.ts` — `buildServerArgs` emits `-ctkt`/`-ctvt` only when set, ordered right after `-ctk`/`-ctv` (the window types they refine), before `--neo-pipeline`. Unset → nothing emitted → byte-identical argv. - `test/launch-args.test.ts` — 4 new cases (omit-when-unset, symmetric q4_0 tail, independent side, ordering); suite 49 pass / 0 fail. **Part 2 — P2-auto boot policy (patch `0121-p2-auto-tail`).** Extends the P1 boot auto-tier so that when it is *forced to spill* (the auto-selected window still overflows the VRAM budget), it can auto-pick a compressed tail rather than re-streaming the whole tail at the window type. Policy is **"q4_0 floor when window > q4_0"**: per axis, a window higher-bit than q4_0 gets a q4_0 tail; a window already at/below q4_0 keeps tail == window. It fires only on the two spill exits of `opencoti_auto_select_kv_tier` (never on the fully-resident fast path), and an explicit `-ctkt`/`-ctvt` is always honored. Gate opt-in is **separate from P1** — `OPENCOTI_KV_AUTO_TIER_TAIL=1` on top of `OPENCOTI_KV_AUTO_TIER=1` — so a P1-only boot stays byte-identical: | leg | envs | tail line | correctness | | --- | --- | --- | --- | | P1AUTO | `OPENCOTI_KV_AUTO_TIER=1` | *(none — tail == window)* | `real_frac 0.0` | | **P2AUTO** | `+ OPENCOTI_KV_AUTO_TIER_TAIL=1` | `auto KV tail = q4_0/q4_0` (~189 MiB / 3072-cell spill) | `real_frac 0.0`, mean-TV 0.0087 | The window was `q5_1/q4_0` at the **second-pass binding budget** — the bug-1342 two-pass compute reserve runs the auto-tier twice; the first (looser) pass fits, the binding pass spills, and the tail policy fires there. HOST-ONLY (`src/llama-kv-cache.cpp` is libllama, not the CUDA DSO). Marker: reuses the 0116 `opencoti-hook: kv-auto-tier` site (like 0118, no new hook). Chain state → 67 patches. ## 2026-07-14 — bug-2172 SHIPPED (patch 0125): rolling-KV window on a HYBRID attention sub-cache Until now the position-window armed cleanly only on caches whose cell placement is **append-only** — dest cell idx == key position, monotone in batch order (every dense / iSWA / unified `llama_kv_cache`). A **hybrid** model (qwen35moe / qwen3next: gated-delta-net recurrent ⊕ full-attention ⊕ MoE ⊕ NextN) routes its attention sub-cache (`mem_attn` inside `llama_memory_hybrid`) through the **recurrent** batch pipeline (`split_seq` + `[TAG_RECURRENT_ROLLBACK_SPLITS]` + `find_slot` cell-reuse/wrap), which hands the scatter **non-monotone** destination idxs. Under tight VRAM / large ctx the window armed and the append-only scatter's monotone `GGML_ASSERT` in `cpy_k` fired at boot. **Root cause** — not a bandwidth or a residency bug (those were #2142 / #2144); the append-only window scatter *assumed* a contiguous window-prefix / tail-suffix partition of the batch, which a recurrent-rollback batch violates. **Fix (host-only, one TU + two one-liners; NO CUDA / NO DSO rebuild):** - New `llama_kv_cache::append_only` flag (default **true** — every existing cache is byte-identical). Set **false** only by `llama_memory_hybrid` on its `mem_attn`. - `!append_only && windowed`: window + tail tensors get **+1 trash sink cell**; the idx tensors become **2-row `[n_tokens,2]`** (row 0 window-route, row 1 tail-route, off-target → trash); `cpy_k`/`cpy_v` scatter the full per-stream K/V into **both** regions via two arbitrary-dest `ggml_set_rows` (drops the monotone assert). `ne[0]` stays `n_tokens` (graph guards hold); fixed 2-set_rows/stream topology (reserve == live). - **Read side unchanged** — `get_k/v_window` / `get_k/v_tail` already cover every active scattered cell, and the streaming-FA `KQ_mask` `-INF`-masks the rollback holes by absolute key position. **Mask contract verified**: `GGML_OP_STREAMING_FLASH_ATTN` is per-cell, absolute-key-indexed, softmax-denominator-normalized ⇒ no kernel change. - **Scope**: iSWA-hybrid out of scope (no current model is both hybrid AND SWA; the window never arms on `swa_type!=NONE`). **Validated** (solidPC host `make`): qwen35moe auto-spill boots with **no assert** + correct output; dense (Qwen) + iSWA (Gemma) regression shields **byte-identical** (append_only path untouched). Marker: `opencoti-hook: hybrid-window-orderagnostic` (3 sites). Chain state → 71 patches. **RYS × rolling-KV-window composition confirmed** (2026-07-14, solidPC 3090, Qwen_Qwen3.6-35B-A3B-Q3_K_M, `--repeat-layers 5,9 --kv-residency-mode window --vram-target 18000 -c 40960 -fa on`): RYS ENGAGED (40 → **44 effective layers**, +4 duplicated) **and** the window armed on the hybrid attention sub-cache (256 resident / 40704 host-tail = 874 MiB) **simultaneously**. A 946-token needle prompt (needle sitting in the ~690-cell spilled tail) → greedy-correct retrieval (`74-ALPHA-purple-9931`), **no assert / no NaN**, 16 graphs reused, clean slot release (pp 1452 tps / tg 57.5 tps). The two features that each stress the hybrid attention cache — RYS's extra duplicated KV layers and the order-agnostic window scatter — compose without interference.