docs/features/rolling_kv.md

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 (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/<bus>/{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_ASSERTed on the M2/NEO/M7 split FA naming __fattn___g-<il> (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 <type> / -ctvt <type> (--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.