From: opencoti Subject: [PATCH 0070] F5 M7 Rolling KV — position-windowed KV residency (shipped) opencoti F5 M7 "Rolling KV" (#296/#338). Cumulative slot capturing the whole W4 workstream: auto-residency + runtime tactic table + the GPU streaming op + the Stage-3 position-window forward + the Stage-4 `--kv-residency-mode` knob. Provenance: opencoti-original; design in docs/features/rolling_kv.md and docs/features/advanced_kv.md (M7). Applies after 0042; reverse-applies clean. The orthogonal cosmocc state-IO fix (bug-269) ships separately as 0071. WHAT SHIPS (the position-window tactic, default `--kv-residency-mode auto`): An iSWA model's overflowing GLOBAL KV layer carries a device-resident window K/V [0, wc) plus a pinned-host tail [wc, n_kv) merged by an online-softmax combine. `wc` is sized from the live VRAM budget (`--vram-target`, else all free VRAM minus a compute reserve); the tactic is auto-selected at cache construction from the eligibility predicate (offload && !v_trans && 0=x1 links; kept in-tree for sub-x1 PCIe). ggml-cpu/ops.{cpp,h}, iqk. * position-axis storage (device window + pinned-host tail tensors), write routing (cpy_k/cpy_v prefix/suffix split), window/tail accessors, POSITION_WINDOW tactic + build_rolling_kv_plan + the host-adaptive compute/bandwidth selector, and state-IO window/tail save+load split: src/llama-kv-cache{,-iswa}.{cpp,h}. * build_attn_mha_position_window + iSWA dispatch: src/llama-graph.{cpp,h}. * --kv-residency-mode / --vram-target / --pcie-bw-gbps threading: common/{arg,common}.{cpp,h} -> include/llama.h -> src/llama-cparams.h -> src/llama-context.cpp -> the kv-cache + hybrid + hybrid-iswa ctor sites (the hybrid sites pass 0 = auto; bug-362). server.cpp pcie write-back. * op-scoped host-tail copy suppression: ggml/src/ggml-backend.cpp. DEFAULT/OFF SAFETY: with KV fitting VRAM the window is empty (single resident FA, byte-identical to vanilla); `--kv-residency-mode head` runs the M2 split. All Stage-2 env-gated diagnostics were removed in S4-2 (load-bearing logic hardcoded to production defaults). Acceptance is logit-distribution equivalence vs vanilla (.opencoti/rolling-kv-equiv-gate.sh), NOT greedy byte-equality (tiled online-softmax is fp-non-associative; bug-263 resolved). diff --git a/llama.cpp/common/arg.cpp b/llama.cpp/common/arg.cpp --- a/llama.cpp/common/arg.cpp +++ b/llama.cpp/common/arg.cpp @@ -1375,91 +1375,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.slot_initial_ctx = value; } ).set_env("LLAMA_ARG_SLOT_INITIAL_CTX").set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_EMBEDDING, LLAMA_EXAMPLE_PARALLEL})); - // opencoti F5 M1 rest-kv-eviction — see docs/features/advanced_kv.md - add_opt(common_arg( - {"--rest-kv-eviction"}, - {"--no-rest-kv-eviction"}, - string_format("when context shift fires, evict the lowest-retention (KeyDiff key-similarity) contiguous window instead of the positional middle (default: %s)", params.rest_kv_eviction ? "enabled" : "disabled"), - [](common_params & params, bool value) { - params.rest_kv_eviction = value; - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_REST_KV_EVICTION")); - add_opt(common_arg( - {"--rest-kv-recent"}, "N", - string_format("rest-kv: protect the most-recent N tokens from eviction (default: %d)", params.rest_kv_recent), - [](common_params & params, int value) { - params.rest_kv_recent = value; - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_REST_KV_RECENT")); - add_opt(common_arg( - {"--rest-kv-layer"}, "N", - string_format("rest-kv: model layer whose keys score retention, -1 = auto/middle (default: %d)", params.rest_kv_layer), - [](common_params & params, int value) { - params.rest_kv_layer = value; - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_REST_KV_LAYER")); - // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md - add_opt(common_arg( - {"--headinfer-gpu-heads-frac"}, "F", - string_format("headinfer: fraction of KV heads kept GPU-resident; the rest are offloaded to host memory and streamed back per step (default: %.2f, 1.0 = off). Only effective for GPU-offloaded layers with flash-attention.", (double) params.headinfer_gpu_heads_frac), - [](common_params & params, const std::string & value) { - params.headinfer_gpu_heads_frac = std::stof(value); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_HEADINFER_GPU_HEADS_FRAC")); - // opencoti F5 M3 neo — see docs/features/advanced_kv.md - add_opt(common_arg( - {"--neo-pipeline"}, "MODE", - "neo (asymmetric GPU/CPU attention pipelining): off (default), on (engage when headinfer split is active for the layer), auto (reserved, same as on). Requires --headinfer-gpu-heads-frac < 1.0 to do anything.", - [](common_params & params, const std::string & value) { - if (value == "off") params.neo_pipeline_mode = 0; - else if (value == "on") params.neo_pipeline_mode = 1; - else if (value == "auto") params.neo_pipeline_mode = 2; - else throw std::invalid_argument("--neo-pipeline must be off|on|auto"); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_NEO_PIPELINE")); - // opencoti F5-opt W2 (#290) iqk CPU flash-attention — see docs/features/advanced_kv.md - add_opt(common_arg( - {"--iqk-flash-attn"}, "MODE", - "iqk CPU flash-attention (vendored ik_llama.cpp engine): off (default), on. " - "Speeds up the M2 CPU-half attention on the token-generation path; falls back " - "to the mainline kernel for unsupported (type, shape) cases. x86_64-only.", - [](common_params & params, const std::string & value) { - if (value == "off" || value == "false" || value == "0") params.iqk_flash_attn = false; - else if (value == "on" || value == "true" || value == "1") params.iqk_flash_attn = true; - else throw std::invalid_argument("--iqk-flash-attn must be on|off"); - ggml_iqk_flash_attn_set_enabled(params.iqk_flash_attn); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_IQK_FLASH_ATTN")); - // opencoti F5-opt W3 (#292) fused-MoE — see docs/features/advanced_kv.md - add_opt(common_arg( - {"--fused-moe-up-gate"}, "MODE", - "fused MoE up+gate+GLU (CUDA decode): off (default), on. Collapses the " - "separate-up/gate expert matmuls + GLU into one op at decode so the CUDA " - "backend runs the fused mmvq kernel without graph-pattern fusion. Only " - "engages for separate-up/gate MoE models (e.g. Gemma-4 A4B) at n_tokens==1.", - [](common_params & params, const std::string & value) { - if (value == "off" || value == "false" || value == "0") params.fused_moe_up_gate = false; - else if (value == "on" || value == "true" || value == "1") params.fused_moe_up_gate = true; - else throw std::invalid_argument("--fused-moe-up-gate must be on|off"); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_FUSED_MOE_UP_GATE")); - // opencoti F5-opt W1 (#293) pcie profile — see docs/features/rolling_kv.md - add_opt(common_arg( - {"--pcie-autodetect"}, "MODE", - "pcie profile: auto-detect host<->device bandwidth + PCIe link at boot (on (default), off). Reads the rebar-probe JSON / nvidia-smi; the result feeds M7 Rolling KV tile sizing.", - [](common_params & params, const std::string & value) { - if (value == "on" || value == "true" || value == "1") params.pcie_autodetect = true; - else if (value == "off" || value == "false" || value == "0") params.pcie_autodetect = false; - else throw std::invalid_argument("--pcie-autodetect must be on|off"); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_PCIE_AUTODETECT")); - add_opt(common_arg( - {"--pcie-bw-gbps"}, "F", - string_format("pcie profile: manual override of effective host->device bandwidth in GB/s (default: %.1f = autodetect). > 0 skips detection.", (double) params.pcie_bw_gbps), - [](common_params & params, const std::string & value) { - params.pcie_bw_gbps = std::stof(value); - } - ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_PCIE_BW_GBPS")); add_opt(common_arg( // opencoti F4 M3 Phase 3 — see docs/decisions/0001-lazy-slot-context.md {"--slot-shrink-idle-ms"}, "N", @@ -1511,13 +1426,43 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_REST_KV_LAYER")); // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md + // opencoti F5 M7 Rolling KV — Rung 0: accepts the literal "auto" which sets the + // sentinel -1, telling the KV-cache constructor to size residency from free VRAM + // (GPU_RESIDENT when KV fits, else a fit-sized CPU spill). A numeric value is the + // manual escape hatch (the upstream M2 fixed split). See docs/features/rolling_kv.md. add_opt(common_arg( - {"--headinfer-gpu-heads-frac"}, "F", - string_format("headinfer: fraction of KV heads kept GPU-resident; the rest are offloaded to host memory and streamed back per step (default: %.2f, 1.0 = off). Only effective for GPU-offloaded layers with flash-attention.", (double) params.headinfer_gpu_heads_frac), + {"--headinfer-gpu-heads-frac"}, "F|auto", + string_format("headinfer: fraction of KV heads kept GPU-resident; the rest are offloaded to host memory and streamed back per step (default: %.2f, 1.0 = off). \"auto\" sizes the fraction from free VRAM (with --vram-target). Only effective for GPU-offloaded layers with flash-attention.", (double) params.headinfer_gpu_heads_frac), [](common_params & params, const std::string & value) { - params.headinfer_gpu_heads_frac = std::stof(value); + if (value == "auto") { + params.headinfer_gpu_heads_frac = -1.0f; // opencoti M7 Rung 0 auto-residency sentinel + } else { + params.headinfer_gpu_heads_frac = std::stof(value); + } } ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_HEADINFER_GPU_HEADS_FRAC")); + // opencoti F5 M7 Rolling KV — Rung 0 auto residency — see docs/features/rolling_kv.md + add_opt(common_arg( + {"--vram-target"}, "MiB", + string_format("rolling-kv: VRAM budget cap in MiB for auto KV residency (used with --headinfer-gpu-heads-frac auto). 0 = use all free VRAM minus a compute reserve (maximize). (default: %d)", params.vram_target_mib), + [](common_params & params, int value) { + params.vram_target_mib = value; + } + ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_VRAM_TARGET")); + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob — see docs/features/rolling_kv.md + add_opt(common_arg( + {"--kv-residency-mode"}, "MODE", + "rolling-kv: KV residency tactic when the cache overflows VRAM — auto (default), head, window. " + "auto = position-window when the cache is eligible (offloaded, non-transposed-V, append-only " + "non-SWA cache with a real overflow), else the M2 head split. head = force the M2 head split " + "(legacy). window = force position-window (falls back with a warning if the cache is ineligible).", + [](common_params & params, const std::string & value) { + if (value == "auto") params.kv_residency_mode = 0; + else if (value == "head") params.kv_residency_mode = 1; + else if (value == "window") params.kv_residency_mode = 2; + else throw std::invalid_argument("--kv-residency-mode must be auto|head|window"); + } + ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_KV_RESIDENCY_MODE")); // opencoti F5 M3 neo — see docs/features/advanced_kv.md add_opt(common_arg( {"--neo-pipeline"}, "MODE", @@ -1542,6 +1487,19 @@ common_params_context common_params_parser_init(common_params & params, llama_ex ggml_iqk_flash_attn_set_enabled(params.iqk_flash_attn); } ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_IQK_FLASH_ATTN")); + // opencoti F5-opt W3 (#292) fused-MoE — see docs/features/advanced_kv.md + add_opt(common_arg( + {"--fused-moe-up-gate"}, "MODE", + "fused MoE up+gate+GLU (CUDA decode): off (default), on. Collapses the " + "separate-up/gate expert matmuls + GLU into one op at decode so the CUDA " + "backend runs the fused mmvq kernel without graph-pattern fusion. Only " + "engages for separate-up/gate MoE models (e.g. Gemma-4 A4B) at n_tokens==1.", + [](common_params & params, const std::string & value) { + if (value == "off" || value == "false" || value == "0") params.fused_moe_up_gate = false; + else if (value == "on" || value == "true" || value == "1") params.fused_moe_up_gate = true; + else throw std::invalid_argument("--fused-moe-up-gate must be on|off"); + } + ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_FUSED_MOE_UP_GATE")); // opencoti F5-opt W1 (#293) pcie profile — see docs/features/rolling_kv.md add_opt(common_arg( {"--pcie-autodetect"}, "MODE", diff --git a/llama.cpp/common/common.cpp b/llama.cpp/common/common.cpp --- a/llama.cpp/common/common.cpp +++ b/llama.cpp/common/common.cpp @@ -1634,6 +1634,13 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.slot_shrink_idle_ms = params.slot_shrink_idle_ms < 0 ? 0u : (uint32_t) params.slot_shrink_idle_ms; // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md cparams.headinfer_gpu_heads_frac = params.headinfer_gpu_heads_frac; + // opencoti F5 M7 Rolling KV — Rung 0 — see docs/features/rolling_kv.md + cparams.vram_target_mib = params.vram_target_mib < 0 ? 0u : (uint32_t) params.vram_target_mib; + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304). PCIe bandwidth for + // streaming-FA tile sizing, resolved by the boot-time probe (W1/0036). + cparams.pcie_bw_gbps = params.pcie_bw_gbps; + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency knob — see docs/features/rolling_kv.md + cparams.kv_residency_mode = params.kv_residency_mode; // opencoti F5 M3 neo — see docs/features/advanced_kv.md cparams.neo_pipeline_mode = params.neo_pipeline_mode; // opencoti F5 W3 #292 fused-MoE — see docs/features/advanced_kv.md diff --git a/llama.cpp/common/common.h b/llama.cpp/common/common.h --- a/llama.cpp/common/common.h +++ b/llama.cpp/common/common.h @@ -440,6 +440,16 @@ struct common_params { // memory and streamed back per step. 1.0 = off (no split). Only active for // GPU-offloaded layers with flash-attention (non-transposed V cache). float headinfer_gpu_heads_frac = 1.0f; + // opencoti F5 M7 Rolling KV — Rung 0 auto residency — see docs/features/rolling_kv.md + // VRAM budget cap (MiB) for auto KV residency (used with + // --headinfer-gpu-heads-frac auto). 0 = use all free VRAM minus a compute + // reserve (maximize). Negative is clamped to 0 when copied into cparams. + int32_t vram_target_mib = 0; + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob — see docs/features/rolling_kv.md + // 0 = auto (POSITION_WINDOW when the cache is eligible, else the M2 head split), + // 1 = head (force the M2 head split / never window), 2 = window (force POSITION_WINDOW, + // fall back with a warning when ineligible). + uint32_t kv_residency_mode = 0; // opencoti F5 M3 neo — see docs/features/advanced_kv.md // 0 = off (default; M2 concat path runs), 1 = on (engage when split // active for the layer), 2 = auto (reserved for future load-based @@ -460,37 +470,6 @@ struct common_params { // x86_64-only at runtime (the engine compiles only on the AVX2 x86_64 slice). bool iqk_flash_attn = false; float pcie_bw_gbps = 0.0f; - // opencoti F4 M3 Phase 2 — see docs/decisions/0001-lazy-slot-context.md - // Initial KV-cell allocation per sequence. 0 = full n_ctx_seq allocation - // (Phase 1 behavior — only zero-fill is deferred). Positive values cap - // the initial allocation and enable grow-on-demand up to n_ctx_seq. - int32_t slot_initial_ctx = 0; - // opencoti F4 M3 Phase 3 — see docs/decisions/0001-lazy-slot-context.md - // Idle threshold (ms) after which a grown soft cap is shrunk back to - // slot_initial_ctx and the over-cap pages are decommitted via - // posix_madvise. 0 disables shrink-on-idle (Phase 2 behavior). - int32_t slot_shrink_idle_ms = 0; - // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md - // Fraction of KV heads kept GPU-resident; the rest are offloaded to host - // memory and streamed back per step. 1.0 = off (no split). Only active for - // GPU-offloaded layers with flash-attention (non-transposed V cache). - float headinfer_gpu_heads_frac = 1.0f; - // opencoti F5 M3 neo — see docs/features/advanced_kv.md - // 0 = off (default; M2 concat path runs), 1 = on (engage when split - // active for the layer), 2 = auto (reserved for future load-based - // engagement; same as `on` today). - uint32_t neo_pipeline_mode = 0; - // opencoti F5-opt W1 (#293) pcie profile — see docs/features/rolling_kv.md - // Boot-time host<->device bandwidth + PCIe link detection, consumed by M7 - // Rolling KV (#296) tile sizing. autodetect reads the rebar-probe JSON / - // nvidia-smi; bw_gbps > 0 forces a manual override (0 = autodetect/default). - bool pcie_autodetect = true; - // opencoti F5-opt W2 (#290) iqk CPU flash-attention — see docs/features/advanced_kv.md - // Engage the vendored ik_llama.cpp CPU Flash-Attention engine on the M2 - // CPU-half (token-generation path). Off by default; flag-off is byte-identical. - // x86_64-only at runtime (the engine compiles only on the AVX2 x86_64 slice). - bool iqk_flash_attn = false; - float pcie_bw_gbps = 0.0f; int32_t n_batch = 2048; // logical batch size for prompt processing (must be >=32 to use BLAS) int32_t n_ubatch = 512; // physical batch size for prompt processing (must be >=32 to use BLAS) int32_t n_keep = 0; // number of tokens to keep from initial prompt diff --git a/llama.cpp/ggml/include/ggml-iqk-flash-attn.h b/llama.cpp/ggml/include/ggml-iqk-flash-attn.h --- a/llama.cpp/ggml/include/ggml-iqk-flash-attn.h +++ b/llama.cpp/ggml/include/ggml-iqk-flash-attn.h @@ -52,6 +52,21 @@ bool iqk_flash_attn_noalibi(int type_q, int type_mask, float max_bias, void * work_buffer, barrier_t barrier, void * barrier_data, int ith, int nth, int n_swa); +// opencoti-hook: f5-rolling-kv — W4 M7 S3 (#354). iqk-backed CPU tail flash-attention +// partial for the CPU_FA_TAIL rolling-KV tactic. Re-declared HERE (not in +// ggml/src/iqk/iqk_mul_mat.h) so the ggml-cpu tail-partial op (ggml-cpu/ops.cpp) can +// call it without the ARM-shim header (same reason as iqk_flash_attn_noalibi above). +// Computes ONE q row's online-softmax over the host-resident tail K/V (nq1=1) and writes +// the UN-normalized [M, S, VKQ] partial the CUDA streaming combine consumes (O=VKQ/S, +// lse=M+logS) — byte-compatible with ops.cpp _f16_one_chunk write_partials. The iqk +// engine uses f32 internal accumulation + the native Dk=Dv=512 blocked kernel (closer to +// the GPU MMA than the scalar _f16_one_chunk, which ceilings at niah~80). Returns false on +// unsupported (type, shape); the caller then falls back to the mainline _f16_one_chunk. +bool iqk_flash_attn_tail_partial(int type_k, int type_v, int Dk, int Dv, + int nk1, int stride_k, int stride_v, int stride_m, + const float * q, const void * k, const void * v, const void * mask, + float scale, float softcap, float * partial); + #ifdef __cplusplus } #endif diff --git a/llama.cpp/ggml/include/ggml.h b/llama.cpp/ggml/include/ggml.h --- a/llama.cpp/ggml/include/ggml.h +++ b/llama.cpp/ggml/include/ggml.h @@ -585,6 +585,10 @@ extern "C" { GGML_OP_MOE_FUSED_UP_GATE, // opencoti-hook: f5-opt-fmoe — W3 #292 + GGML_OP_STREAMING_FLASH_ATTN, // opencoti-hook: f5-rolling-kv — W4 M7-D #304 + + GGML_OP_FLASH_ATTN_TAIL_PARTIAL, // opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353 + GGML_OP_COUNT, }; @@ -2423,6 +2427,56 @@ extern "C" { float max_bias, float logit_softcap); + // opencoti F5 M7 Rolling KV (W4 M7-D #304) — streaming flash attention. + // Same contract as ggml_flash_attn_ext (q,k,v,mask,scale,max_bias, + // logit_softcap → F32 [n_embd_v, n_head_q, n_q, n_batch]) but the CUDA + // forward tiles over the key dimension and combines tiles with a + // cross-tile online-softmax (running max + denom in F32) so the full K/V + // need not fit one resident slot. CUDA-only (CPU forward aborts). The + // tile/slot sizing is decided inside the op from device free-VRAM; the + // graph builder passes cparams.pcie_bw_gbps via op_params for the + // double-buffer (R2-b). At R2-a the source K/V may still be GPU-resident + // — the op's only job is to prove tiled online-softmax == monolithic FA. + GGML_API struct ggml_tensor * ggml_streaming_flash_attn( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap); + + // opencoti F5 M7 Stage 3c-5 (#341) — position-window streaming FA: resident + // device window K/V (k_win/v_win) + pinned-host overflow tail K/V + // (k_tail/v_tail, may be NULL ⇒ single resident region). One online-softmax + // combine over both; mask spans the full key range [0, n_win + n_tail). + GGML_API struct ggml_tensor * ggml_streaming_flash_attn_window( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k_win, + struct ggml_tensor * v_win, + struct ggml_tensor * k_tail, + struct ggml_tensor * v_tail, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap); + + // opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353. CPU tail-partial emitter for + // the CPU_FA_TAIL tactic. dst = [2+DV, n_head, n_q, n_b] holding [M, S, VKQ] + // (un-normalized) per output row — the partial the CUDA online-softmax combine + // consumes (O=VKQ/S, lse=M+log S). See ggml.c for the full contract. + GGML_API struct ggml_tensor * ggml_flash_attn_ext_tail_partial( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap); + GGML_API void ggml_flash_attn_ext_set_prec( struct ggml_tensor * a, enum ggml_prec prec); diff --git a/llama.cpp/ggml/src/ggml-backend.cpp b/llama.cpp/ggml/src/ggml-backend.cpp --- a/llama.cpp/ggml/src/ggml-backend.cpp +++ b/llama.cpp/ggml/src/ggml-backend.cpp @@ -1247,6 +1247,23 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra for (int b = 0; b < sched->n_backends && *cur_backend_id == -1; b++) { ggml_backend_sched_set_if_supported(sched, node, b, cur_backend_id); } + // opencoti-hook: f5-rolling-kv — bug-259. STREAMING_FLASH_ATTN (and the + // MOE_FUSED_UP_GATE sibling) are CUDA-only ops with no CPU kernel. CUDA's + // supports_op gates on get_best_fattn_kernel, which returns NONE for the + // tiny -fit / empty-warmup TRIAL graphs — so the op is unassigned here and + // (with CPU now reporting it unsupported too) would hit the assert below. + // Pin it to the highest-priority GPU backend (CPU is the last index by this + // scheduler's own convention, lines ~1025). The -fit pass only reserves + // memory (no op execution) and the empty warmup run executes zero rows, so + // a conservative trial-graph kernel check is harmless; the real inference + // graph is assigned to CUDA normally (FA is supported there). The op's + // CUDA forward self-selects the per-tile kernel + resident fallback, and the + // S3c copy-suppression (~:1285) already handles its pinned-host K/V srcs. + if (*cur_backend_id == -1 && sched->n_backends > 1 && + (node->op == GGML_OP_STREAMING_FLASH_ATTN || node->op == GGML_OP_MOE_FUSED_UP_GATE)) { + *cur_backend_id = 0; + SET_CAUSE(node, "4.cuda-only"); + } GGML_ASSERT(*cur_backend_id != -1); } @@ -1357,7 +1374,31 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra } } - if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) { + // opencoti-hook: f5-rolling-kv — S3c scheduler bypass. + // GGML_OP_STREAMING_FLASH_ATTN reads its CPU-half K/V (src 1,2) + // directly from the pinned (cuda_host) buffer and DMAs tiles + // host->device itself on its copy_stream (fattn.cu). Suppress the + // scheduler's full-CPU-half device materialization for THIS op + // only; node->src[j] is left pointing at the host view so the cpy + // store->read edge on k_cpu_per_stream survives (correctness) and + // the op recovers the pinned-host pointer via src->data. Scoped to + // the distinct op type → inert for every other op. + // The host streaming path is the product default: the op reads its + // host-resident K/V in place and DMAs tiles itself, so suppress the + // scheduler's device materialization for any host-buffer src on this op. + // 3c-5: the POSITION_WINDOW variant carries the pinned-host overflow TAIL + // at src[5]/src[6] (the device window is src[1]/src[2]); the host-buffer + // guard leaves the device window untouched, so the single-region path + // (no src[5]) is byte-identical. + bool s3c_bypass = false; + if (node->op == GGML_OP_STREAMING_FLASH_ATTN && (j == 1 || j == 2 || j == 5 || j == 6)) { + ggml_backend_buffer_t sbuf = src->view_src ? src->view_src->buffer : src->buffer; + if (sbuf != NULL && ggml_backend_buffer_is_host(sbuf)) { + s3c_bypass = true; + } + } + + if (src_backend_id != cur_backend_id && !s3c_bypass && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) { // create a copy of the input in the split's backend if (tensor_id_copy(src_id, cur_backend_id, 0) == NULL) { ggml_backend_t backend = sched->backends[cur_backend_id]; diff --git a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c --- a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c +++ b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c @@ -2035,6 +2035,12 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_flash_attn_ext(params, tensor); } break; + case GGML_OP_FLASH_ATTN_TAIL_PARTIAL: + { + // opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353. CPU tail-partial + // emitter for the CPU_FA_TAIL tactic (host-resident tail K/V). + ggml_compute_forward_flash_attn_ext_tail_partial(params, tensor); + } break; case GGML_OP_FLASH_ATTN_BACK: { int32_t t = ggml_get_op_params_i32(tensor, 0); @@ -2073,6 +2079,16 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm // owns the FFN nodes, so the CPU scheduler never reaches here. GGML_ABORT("MOE_FUSED_UP_GATE has no CPU implementation (CUDA-only op)"); } break; + case GGML_OP_STREAMING_FLASH_ATTN: + { + // opencoti-hook: f5-rolling-kv — W4 M7-D #304. CUDA-only op. The + // build_attn_mha_streaming graph hook emits it only when the + // GPU_STREAM tactic is assigned (which requires the CUDA + // backend), so the CPU scheduler never reaches here. The CPU + // tactic for an over-budget layer is CPU_SPILL (the W2 iqk FA + // engine), a different path entirely. + GGML_ABORT("STREAMING_FLASH_ATTN has no CPU implementation (CUDA-only op)"); + } break; case GGML_OP_GET_REL_POS: { ggml_compute_forward_get_rel_pos(params, tensor); @@ -2346,6 +2362,12 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { // scheduled on CPU. Single task keeps the planner well-formed. n_tasks = 1; } break; + case GGML_OP_STREAMING_FLASH_ATTN: + { + // opencoti-hook: f5-rolling-kv — W4 M7-D #304. CUDA-only; never + // scheduled on CPU. Single task keeps the planner well-formed. + n_tasks = 1; + } break; case GGML_OP_SILU_BACK: case GGML_OP_MUL: case GGML_OP_DIV: @@ -2423,6 +2445,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_ARGSORT: case GGML_OP_TOP_K: case GGML_OP_FLASH_ATTN_EXT: + case GGML_OP_FLASH_ATTN_TAIL_PARTIAL: // opencoti-hook: f5-rolling-kv — S3-a #353 case GGML_OP_FLASH_ATTN_BACK: case GGML_OP_SSM_CONV: case GGML_OP_SSM_SCAN: @@ -3002,6 +3025,17 @@ struct ggml_cplan ggml_graph_plan( cur = MAX(cur, iqk_fa_work_buffer_size(node, n_tasks)); #endif } break; + case GGML_OP_FLASH_ATTN_TAIL_PARTIAL: + { + // opencoti-hook: f5-rolling-kv — S3-a #353. Per-thread + // _f16_one_chunk scratch only (the [M,S,VKQ] partial lands in + // dst, not wdata). Mirrors that fn's per-thread wdata stride + // (DK + 2*DV + cache-line pad); +64 floats covers the pad + // generously across cache-line sizes. + const int64_t DK = node->src[1]->ne[0]; + const int64_t DV = node->src[2]->ne[0]; + cur += sizeof(float)*n_tasks*(DK + 2*DV + 64); + } break; case GGML_OP_FLASH_ATTN_BACK: { const int64_t D = node->src[0]->ne[0]; diff --git a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp --- a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -468,6 +468,17 @@ static bool GGML_CALL ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev case GGML_OP_OUT_PROD: return (src0->type == GGML_TYPE_F32 || (ggml_is_quantized(src0->type) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; + // opencoti-hook: f5-rolling-kv / f5-opt-fmoe — bug-259. These are CUDA-only + // ops (no CPU forward; the dispatch aborts at ggml-cpu.c:2026/2036). Report + // them UNSUPPORTED on CPU so the scheduler always pins them to the CUDA + // backend. Without this they fall to `default: return true`, and the ggml + // offload heuristic can drift the op to CPU for a small/empty (warmup) batch + // whose K/V srcs are pinned CUDA_Host buffers → abort. (MOE_FUSED was saved + // only by graph-hook placement; STREAMING_FLASH_ATTN under M7 stream-by- + // default is not — both fixed here for robustness.) + case GGML_OP_STREAMING_FLASH_ATTN: + case GGML_OP_MOE_FUSED_UP_GATE: + return false; default: return true; } diff --git a/llama.cpp/ggml/src/ggml-cpu/ops.cpp b/llama.cpp/ggml/src/ggml-cpu/ops.cpp --- a/llama.cpp/ggml/src/ggml-cpu/ops.cpp +++ b/llama.cpp/ggml/src/ggml-cpu/ops.cpp @@ -8383,7 +8383,12 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk( const int64_t DV = nev0; const int64_t N = neq1; - GGML_ASSERT(ne0 == DV); + // opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353. In write_partials mode the + // output goes to `partials` (caller-owned), not dst->data, so dst's ne0 is the + // caller's layout: the stock FA split_kv_path keeps ne0==DV, while the tail- + // partial op uses ne0==2+DV ([M,S,VKQ]). Skip the ne0==DV dst-shape check when + // writing partials; the normalized-output path (FLASH_ATTN_EXT) is unchanged. + GGML_ASSERT(write_partials || ne0 == DV); GGML_ASSERT(ne2 == N); // input tensor rows must be contiguous @@ -9227,6 +9232,54 @@ void ggml_compute_forward_flash_attn_ext( } } +// opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353. CPU forward for +// GGML_OP_FLASH_ATTN_TAIL_PARTIAL. One online-softmax pass over the FULL (host- +// resident) tail K/V per output row, writing the un-normalized [M, S, VKQ] partial +// into dst (shape [2+DV, n_head, n_q, n_b]). Reuses the proven _f16_one_chunk in +// its write_partials mode (ic over [0, nek1), partial_stride = 2+DV); rows are +// partitioned across threads. No reduce/normalize here — the partial IS the output; +// the CUDA online-softmax combine normalizes (O=VKQ/S) and merges with the device +// window leg. Decode consumer (n_q==1): one_chunk's ir order is head-fastest then, +// matching the combine's row index. f16/f32 K/V only (mainline path; no iqk fast +// path — correctness first, the selector measures whatever this costs). +void ggml_compute_forward_flash_attn_ext_tail_partial( + const ggml_compute_params * params, + ggml_tensor * dst) { + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k = dst->src[1]; + const ggml_tensor * v = dst->src[2]; + + const int64_t DK = k->ne[0]; + const int64_t DV = v->ne[0]; + const int64_t nek1 = k->ne[1]; + + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(dst->ne[0] == 2 + DV); + + // rows = (head, q-token, batch) + const int64_t nr = q->ne[1] * q->ne[2] * q->ne[3]; + + const int ith = params->ith; + const int nth = params->nth; + + const int64_t dr = (nr + nth - 1) / nth; + const int64_t ir0 = dr * ith; + const int64_t ir1 = MIN(ir0 + dr, nr); + + if (ir0 >= ir1) { + return; + } + + // opencoti F5 M7 S3 (#354): CPU tail-partial. Single online-softmax pass over the + // host-resident tail K/V, emitting each row's [M, S, VKQ_unnorm] partial exactly where + // the streaming online-softmax combine expects it. Reuses the mainline _f16_one_chunk + // write_partials mode (zero recompute). Dormant unless the S2 selector picks CPU_FA_TAIL + // (sub-x1 PCIe link; gated out of auto-select per #354). + ggml_compute_forward_flash_attn_ext_f16_one_chunk( + params, dst, (int) ir0, (int) ir1, 0, nek1, + (float *) dst->data, 2 + DV); +} + // ggml_compute_forward_flash_attn_back static void ggml_compute_forward_flash_attn_back_f32( diff --git a/llama.cpp/ggml/src/ggml-cpu/ops.h b/llama.cpp/ggml/src/ggml-cpu/ops.h --- a/llama.cpp/ggml/src/ggml-cpu/ops.h +++ b/llama.cpp/ggml/src/ggml-cpu/ops.h @@ -87,6 +87,8 @@ void ggml_compute_forward_leaky_relu(const struct ggml_compute_params * params, void ggml_compute_forward_tri(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_fill(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_flash_attn_ext(const struct ggml_compute_params * params, struct ggml_tensor * dst); +// opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353. Emits the [M,S,VKQ] tail partial. +void ggml_compute_forward_flash_attn_ext_tail_partial(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_flash_attn_back( const struct ggml_compute_params * params, const bool masked, diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh --- a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh @@ -41,6 +41,13 @@ typedef void (* fattn_kernel_t)( const int32_t ne31, const int32_t ne32, const int32_t ne33, const int32_t nb31, const int32_t nb32, const int64_t nb33); +// opencoti-hook: rolling-kv M7 (Stage 3a) — consume-once side channel for requesting a per-row +// log-sum-exp output from the next launch_fattn, without threading dst_lse through the 5-level +// MMA dispatch (mma_f16 → switch_ncols2 → switch_ncols1 → case → launch_fattn). The windowed-FA +// op sets this immediately before dispatching one FA; launch_fattn drains it (sets it back to +// nullptr) so it never leaks to an unrelated FA. Defined in fattn.cu. See UPSTREAM_SYNC.md. +extern thread_local float * opencoti_fattn_dst_lse; + typedef float (*vec_dot_KQ_t)( const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8 , const void * __restrict__ Q_ds); @@ -679,6 +686,7 @@ template // D == head size __launch_bounds__(D, 1) static __global__ void flash_attn_stream_k_fixup_uniform( float * __restrict__ dst, + float * __restrict__ dst_lse, // opencoti-hook: rolling-kv M7 — optional per-row log-sum-exp out (nullptr = off) const float2 * __restrict__ dst_fixup, const int ne01, const int ne02, const int ne12, const int nblocks_stream_k, @@ -753,6 +761,12 @@ static __global__ void flash_attn_stream_k_fixup_uniform( // Write back final result: *dst = dst_val / rowsum; + + // opencoti-hook: rolling-kv M7 — per-row log-sum-exp for windowed-FA online-softmax combine (Stage 3a). See docs/protocols/UPSTREAM_SYNC.md. + if (dst_lse && tid == 0) { + const int64_t lse_row = int64_t(sequence)*ne02*ne01 + int64_t(jt)*ne02*ncols1 + zt_Q + (j*ne02 + c); + dst_lse[lse_row] = max_val + logf(rowsum); + } } // General fixup kernel for the case where the number of blocks per tile is not uniform across tiles @@ -761,6 +775,7 @@ template // D == head size __launch_bounds__(D, 1) static __global__ void flash_attn_stream_k_fixup_general( float * __restrict__ dst, + float * __restrict__ dst_lse, // opencoti-hook: rolling-kv M7 — optional per-row log-sum-exp out (nullptr = off) const float2 * __restrict__ dst_fixup, const int ne01, const int ne02, const int gqa_ratio, @@ -862,6 +877,12 @@ static __global__ void flash_attn_stream_k_fixup_general( // Write back final result: *dst = dst_val / rowsum; + + // opencoti-hook: rolling-kv M7 — per-row log-sum-exp for windowed-FA online-softmax combine (Stage 3a). See docs/protocols/UPSTREAM_SYNC.md. + if (dst_lse && tid == 0) { + const int64_t lse_row = int64_t(sequence)*ne02*ne01 + int64_t(jt)*ne02*ncols1 + zt_Q + (j*ne02 + c); + dst_lse[lse_row] = max_val + logf(rowsum); + } } template // D == head size @@ -870,6 +891,7 @@ static __global__ void flash_attn_combine_results( const float * __restrict__ VKQ_parts, const float2 * __restrict__ VKQ_meta, float * __restrict__ dst, + float * __restrict__ dst_lse, // opencoti-hook: rolling-kv M7 — optional per-row log-sum-exp out (nullptr = off) const int parallel_blocks) { ggml_cuda_pdl_lc(); // Dimension 0: threadIdx.x @@ -917,15 +939,27 @@ static __global__ void flash_attn_combine_results( } dst[tid] = VKQ_numerator / VKQ_denominator; + + // opencoti-hook: rolling-kv M7 — per-row log-sum-exp for windowed-FA online-softmax combine (Stage 3a). See docs/protocols/UPSTREAM_SYNC.md. + if (dst_lse && tid == 0) { + dst_lse[j_dst_unrolled] = kqmax + logf(VKQ_denominator); + } } template void launch_fattn( ggml_backend_cuda_context & ctx, ggml_tensor * dst, fattn_kernel_t fattn_kernel, const int nwarps, const size_t nbytes_shared, - const int nbatch_fa, const bool need_f16_K, const bool need_f16_V, const bool stream_k, const int warp_size = WARP_SIZE + const int nbatch_fa, const bool need_f16_K, const bool need_f16_V, const bool stream_k, const int warp_size = WARP_SIZE, + float * dst_lse = nullptr // opencoti-hook: rolling-kv M7 — per-row LSE out (Stage 3a); nullptr = byte-identical legacy path ) { constexpr int ncols = ncols1 * ncols2; + // opencoti-hook: rolling-kv M7 (Stage 3a) — drain the consume-once dst_lse side channel. + if (dst_lse == nullptr && opencoti_fattn_dst_lse != nullptr) { + dst_lse = opencoti_fattn_dst_lse; + } + opencoti_fattn_dst_lse = nullptr; + const ggml_tensor * Q = dst->src[0]; const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; @@ -1171,6 +1205,18 @@ void launch_fattn( ); CUDA_CHECK(cudaGetLastError()); + // opencoti-hook: rolling-kv M7 (Stage 3a) — dst_lse is only populated by the finalize kernels + // (uniform/general fixup, or the parallel-blocks combine). If the dispatch resolves to the + // in-kernel single-block write path, no finalize runs and dst_lse would be left garbage; assert + // loudly instead of silently corrupting the windowed-FA online-softmax combine. The in-kernel LSE + // emit (typedef-threaded MMA epilogue) is a separate prefill-overflow stage; on the decode path + // (stream_k always engaged) a fixup always runs. See docs/protocols/UPSTREAM_SYNC.md. + if (dst_lse) { + const bool fixup_ran = stream_k && (((int)blocks_num.x % ntiles_dst == 0 && (int)blocks_num.x > ntiles_dst) || (ntiles_dst % (int)blocks_num.x != 0)); + const bool combine_ran = !stream_k && parallel_blocks > 1; + GGML_ASSERT((fixup_ran || combine_ran) && "rolling-kv: dst_lse requested but FA resolved to the in-kernel write path (no finalize kernel); unsupported on this path"); + } + if (stream_k) { if ((int)blocks_num.x % ntiles_dst == 0 && (int)blocks_num.x > ntiles_dst) { // Optimized fixup: nblocks_stream_k is a multiple of ntiles_dst, launch one block per tile. @@ -1186,7 +1232,7 @@ void launch_fattn( const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks_num_combine, block_dim_combine, 0, main_stream); ggml_cuda_kernel_launch(flash_attn_stream_k_fixup_uniform, launch_params, - (float *) KQV->data, dst_tmp_meta.ptr, + (float *) KQV->data, dst_lse, dst_tmp_meta.ptr, Q->ne[1], Q->ne[2], K->ne[2], nblocks_sk, gqa_ratio, bpt, fd0, fd1, fd2); } else if (ntiles_dst % blocks_num.x != 0) { @@ -1203,7 +1249,7 @@ void launch_fattn( const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks_num_combine, block_dim_combine, 0, main_stream); ggml_cuda_kernel_launch(flash_attn_stream_k_fixup_general, launch_params, - (float *) KQV->data, dst_tmp_meta.ptr, + (float *) KQV->data, dst_lse, dst_tmp_meta.ptr, Q->ne[1], Q->ne[2], gqa_ratio, total_work, fd_k_j_z_ne12, fd_k_j_z, fd_k_j, fd_k); } @@ -1214,7 +1260,7 @@ void launch_fattn( const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks_num_combine, block_dim_combine, nbytes_shared_combine, main_stream); ggml_cuda_kernel_launch(flash_attn_combine_results, launch_params, - dst_tmp.ptr, dst_tmp_meta.ptr, (float *) KQV->data, parallel_blocks); + dst_tmp.ptr, dst_tmp_meta.ptr, (float *) KQV->data, dst_lse, parallel_blocks); } CUDA_CHECK(cudaGetLastError()); } diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn.cu b/llama.cpp/ggml/src/ggml-cuda/fattn.cu --- a/llama.cpp/ggml/src/ggml-cuda/fattn.cu +++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cu @@ -6,6 +6,11 @@ #include "fattn-wmma-f16.cuh" #include "fattn.cuh" +// opencoti-hook: rolling-kv M7 (Stage 3a) — definition of the consume-once dst_lse side channel +// declared in fattn-common.cuh. Set immediately before a single launch_fattn dispatch to request a +// per-row log-sum-exp output; launch_fattn drains it back to nullptr. See docs/protocols/UPSTREAM_SYNC.md. +thread_local float * opencoti_fattn_dst_lse = nullptr; + template static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; @@ -560,3 +565,841 @@ void ggml_cuda_flash_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst bool ggml_cuda_flash_attn_ext_supported(int device, const ggml_tensor * dst) { return ggml_cuda_get_best_fattn_kernel(device, dst) != BEST_FATTN_KERNEL_NONE; } + +// ============================================================================ +// opencoti-hook: f5-rolling-kv — W4 M7-D #304 / R2-b S2 (#311). +// Streaming flash-attention via key-axis tiling + LSE-weighted online-softmax. +// +// The host head-group's K/V (dst->src[1], src[2]) are processed along the key +// axis in tiles. Each tile runs the STOCK ggml_cuda_flash_attn_ext (untouched — +// the resident path stays byte-identical, every kernel family is covered) over +// a key-slice, giving that tile's normalised output O_t. A small reduction +// emits each tile's per-output-row log-sum-exp lse_t. Tiles merge in F32: +// O = Σ_t O_t·exp(lse_t − M) / Σ_t exp(lse_t − M), M = max_t lse_t +// which is the exact online-softmax (algebra verified). Single-tile +// short-circuits to plain FA (= R2-b S1, byte-identical, zero overhead). +// +// Tile count is AUTO (sized from n_kv); a single resolved tile over a resident +// source short-circuits to plain FA. The slot pool (default 2 slots) stages each +// tile's K/V through a round-robin GPU slot via a +// SINGLE-STREAM cudaMemcpyAsync(cudaMemcpyDefault) on the op's own ctx.stream() +// — no dedicated copy_stream and no events. UVA infers H2D (pinned-host spill) +// or D2D (resident) per pointer. ctx.stream() rides whatever stream the M2/M7 +// head-split's concurrent_events forked the op onto, so it cannot collide with +// that machinery (bug-253) and stays CUDA-graph-capturable. Bit-stable vs the +// resident path. The dedicated copy_stream + copy_done/compute_done events +// (the actual double-buffer overlap) are S3b's job; S3c adds the host→device +// streaming that is the memory win. +// +// Scoring coverage: scale + additive mask only — max_bias==0 (no ALiBi), +// logit_softcap==0, no sinks (src[4]==NULL). Otherwise this falls back to +// resident FA. Gate model qwen2.5-1.5b satisfies all three. +// ============================================================================ + +// One block per output row (head, q-token, batch); WARP_SIZE threads cooperate +// on each head_dim dot. Online (max, sumexp) over the tile's keys → lse[row]. +static __global__ void streaming_lse_kernel( + const char * __restrict__ Q, const char * __restrict__ K, + const char * __restrict__ mask, float * __restrict__ lse, + const float scale, const float logit_softcap, const int tile_kv, const int gqa_ratio, + const int head_dim, const int n_head, const int n_q, + const int64_t q_nb1, const int64_t q_nb2, const int64_t q_nb3, + const int64_t k_nb1, const int64_t k_nb2, const int64_t k_nb3, + const int64_t m_nb0, const int64_t m_nb1) { + const int row = blockIdx.x; // flattened (h, qi, b) + const int h = row % n_head; + const int qi = (row / n_head) % n_q; + const int b = row / (n_head * n_q); + const int hk = h / gqa_ratio; + const int lane = threadIdx.x; + + const float * Qp = (const float *)(Q + (int64_t)qi*q_nb1 + (int64_t)h*q_nb2 + (int64_t)b*q_nb3); + const char * Kb = K + (int64_t)hk*k_nb2 + (int64_t)b*k_nb3; + const char * Mb = mask ? (mask + (int64_t)qi*m_nb1) : nullptr; + + float m = -INFINITY; + float l = 0.0f; + for (int j = 0; j < tile_kv; ++j) { + const half * Kj = (const half *)(Kb + (int64_t)j*k_nb1); + float part = 0.0f; + for (int d = lane; d < head_dim; d += WARP_SIZE) { + part += Qp[d] * __half2float(Kj[d]); + } +#pragma unroll + for (int off = WARP_SIZE/2; off > 0; off >>= 1) { + part += __shfl_xor_sync(0xFFFFFFFF, part, off, WARP_SIZE); + } + // S3b softcap (#340): mirror plain FA — softcap*tanhf((scale/softcap)*qk) + // BEFORE the mask add (fattn-mma-f16.cuh:611-621 then :640), divide-by-zero + // guarded. softcap==0 is the byte-identical pre-S3b path. + float s = (logit_softcap != 0.0f) + ? logit_softcap * tanhf((scale / logit_softcap) * part) + : scale * part; + if (Mb) { + s += __half2float(*(const half *)(Mb + (int64_t)j*m_nb0)); + } + if (s == -INFINITY) { // fully-masked key — contributes nothing + continue; + } + const float m_new = fmaxf(m, s); + l = l * expf(m - m_new) + expf(s - m_new); // m=-inf first hit: expf(-inf)=0 + m = m_new; + } + if (lane == 0) { + lse[row] = (l > 0.0f) ? (m + logf(l)) : -INFINITY; + } +} + +// One thread per output element. dst layout [DV, n_head, n_q, n_b]; each tile's +// O block and the dst share that contiguous layout, so the row index used for +// lse is exactly elem/DV — identical to streaming_lse_kernel's `row`. +static __global__ void streaming_combine_kernel( + const float * __restrict__ O_tiles, const float * __restrict__ lse_tiles, + float * __restrict__ dst, const int n_tiles, const int dv, const int n_rows) { + const int64_t total = (int64_t)dv * n_rows; + const int64_t gid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total) { + return; + } + const int row = (int)(gid / dv); + float M = -INFINITY; + for (int t = 0; t < n_tiles; ++t) { + M = fmaxf(M, lse_tiles[(int64_t)t*n_rows + row]); + } + float num = 0.0f; + float den = 0.0f; + for (int t = 0; t < n_tiles; ++t) { + const float lse = lse_tiles[(int64_t)t*n_rows + row]; + const float w = (lse == -INFINITY || M == -INFINITY) ? 0.0f : expf(lse - M); + if (w != 0.0f) { + num += O_tiles[(int64_t)t*total + gid] * w; + den += w; + } + } + dst[gid] = (den > 0.0f) ? (num / den) : 0.0f; +} + +// opencoti-hook: f5-rolling-kv — S3-b (#353) CPU_FA_TAIL. Convert a precomputed +// CPU tail partial into ONE combine tile. The partial is the [2+DV, n_rows] +// output of ggml_flash_attn_ext_tail_partial: per output row r (contiguous at +// r*(2+DV)) it holds [M, S, VKQ_unnorm[0..DV)] — un-normalized online-softmax +// state over the host-resident tail keys. This kernel writes that row's: +// O_tile[row*DV + d] = VKQ_unnorm[d] / S (normalized O, DV-fastest — the +// exact streaming_combine_kernel layout) +// lse_tile[row] = M + logf(S) (the per-row log-sum-exp) +// so the tail merges through the UNCHANGED streaming_combine_kernel alongside the +// device window tile(s). Only the tiny partial (2+DV floats/row) ever crossed +// PCIe — never the tail K/V (the CPU_FA_TAIL win on slow links). One thread per +// (row, d) output element. S<=0 (fully-masked / empty tail) → O=0, lse=-inf +// (the combine ignores a -inf tile). +static __global__ void tail_partial_to_tile_kernel( + const float * __restrict__ partial, float * __restrict__ O_tile, + float * __restrict__ lse_tile, const int dv, const int n_rows) { + const int64_t total = (int64_t)dv * n_rows; + const int64_t gid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total) { + return; + } + const int row = (int)(gid / dv); + const int d = (int)(gid % dv); + const float * pr = partial + (int64_t)row * (2 + dv); + const float M = pr[0]; + const float S = pr[1]; + O_tile[gid] = (S > 0.0f) ? (pr[2 + d] / S) : 0.0f; + if (d == 0) { + lse_tile[row] = (S > 0.0f) ? (M + logf(S)) : -INFINITY; + } +} + +void ggml_cuda_streaming_flash_attn(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + // opencoti F5 M7: the tile count is AUTO — sized from n_kv below (~the default + // keys/tile). The single-tile short-circuit to plain FA is deferred below to + // AFTER the host detection — a host (bypass) source must still take the slot lift + // even at one tile (plain FA would read host memory as device); only a + // device/resident single tile is byte-identically plain FA (the S1 anchor). + + float scale = 1.0f, max_bias = 0.0f, logit_softcap = 0.0f; + memcpy(&scale, (const float *) dst->op_params + 0, sizeof(float)); + memcpy(&max_bias, (const float *) dst->op_params + 1, sizeof(float)); + memcpy(&logit_softcap, (const float *) dst->op_params + 2, sizeof(float)); + + // Unsupported scoring → keep this layer resident (correctness over relief). + // opencoti-hook: f5-rolling-kv — S3b (#340) softcap relief. Plain FA handles + // softcap natively, so drop the softcap bail for head_dim>256 (MMA-only: VEC is + // excluded at >256, so the per-tile FA always reaches a finalize kernel — see + // the dst_lse decode channel below). head_dim<=256 keeps the softcap bail (VEC + // could be selected → no finalize → the dst_lse safety assert). ALiBi (max_bias) + // and sinks (src[4]) bails stay unconditional (the combine folds neither yet). + // dst->src[1]==K, but K is not fetched until below, so read ne[0] off src[1]. + if (max_bias != 0.0f || dst->src[4] != nullptr || + (logit_softcap != 0.0f && dst->src[1]->ne[0] <= 256)) { + ggml_cuda_flash_attn_ext(ctx, dst); + return; + } + + const ggml_tensor * Q = dst->src[0]; + const ggml_tensor * K = dst->src[1]; + const ggml_tensor * V = dst->src[2]; + const ggml_tensor * mask = dst->src[3]; + + // opencoti-hook: f5-rolling-kv — Stage 3c-5 (#346) two-region (position-window) + // forward. When src[5]/src[6] (k_tail/v_tail) are present the op fuses a + // device-resident WINDOW [0,n_win) (src[1]/src[2], the bulk) and a pinned-host + // overflow TAIL [n_win,n_kv) (src[5]/src[6], local idx = global−n_win) into ONE + // online-softmax combine: window tiles read the device source, tail tiles read + // the host source, every partial merges through the unchanged + // streaming_combine_kernel. src[5]==null ⇒ single region (the established + // S2/S3a/S3b path), byte-identical. The window and tail share K/V type, head_dim + // and n_head_kv (asserted below); only the key-position count (ne[1]) differs, so + // the resolve() mapping below treats them as one logical [0,n_kv) key space split + // at n_win. + // opencoti-hook: f5-rolling-kv — S3-b (#353) CPU_FA_TAIL. op_params[4]==1 marks + // "tail is a precomputed CPU partial" mode: src[5] is NOT k_tail but the + // [2+DV, n_head, n_q, n_b] partial [M,S,VKQ_unnorm] produced on the CPU by + // ggml_flash_attn_ext_tail_partial. In this mode the op runs ordinary + // window-only tiling over the device-resident bulk [0,n_win) (src[1]/src[2]) + // and merges the converted partial (O=VKQ/S, lse=M+logS) as ONE extra combine + // tile — only the tiny partial crossed PCIe, never the tail K/V. src[6] unused. + // op_params is zero-initialised (ggml.c new_tensor), so every legacy/two-region + // op has op_params[4]==0 ⇒ this branch never fires for them. + const bool tail_is_partial = (ggml_get_op_params_i32(dst, 4) == 1); + const ggml_tensor * tail_partial = tail_is_partial ? dst->src[5] : nullptr; + const ggml_tensor * K_tail = tail_is_partial ? nullptr : dst->src[5]; + const ggml_tensor * V_tail = tail_is_partial ? nullptr : dst->src[6]; + const bool two_region = (K_tail != nullptr); + if (two_region) { + GGML_ASSERT(V_tail != nullptr); + GGML_ASSERT(K_tail->type == K->type && V_tail->type == V->type); + GGML_ASSERT(K_tail->ne[0] == K->ne[0] && V_tail->ne[0] == V->ne[0]); + GGML_ASSERT(K_tail->ne[2] == K->ne[2]); + } + + // opencoti-hook: f5-rolling-kv — S3d dequant-on-lift. The streaming slot is + // ALWAYS f16 (streaming_lse_kernel reads it via __half2float); a quantized + // K/V is dequantized into the slot on the lift using ggml_get_to_fp16_nc_cuda + // — the exact primitive fattn-common.cuh:965-1025 uses to convert strided K/V + // to f16 (strides in elements = nb/type_size, packed f16 output). This + // REPLACES the bug-254 f16 stop-gap and unblocks bug-226 (q8_0 @ 262144). + // Two safety guards: (1) an unsupported KV type (no nc converter) stays + // resident; (2) a quantized type that ends up with no slot to dequant into + // (n_slots==0) stays resident — both enforced below. The f16 path resolves + // to nullptr converters and is byte-identical to pre-S3d. + const bool all_f16 = (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16); + to_fp16_nc_cuda_t to_fp16_k = all_f16 ? nullptr : ggml_get_to_fp16_nc_cuda(K->type); + to_fp16_nc_cuda_t to_fp16_v = all_f16 ? nullptr : ggml_get_to_fp16_nc_cuda(V->type); + if (!all_f16 && (to_fp16_k == nullptr || to_fp16_v == nullptr)) { + ggml_cuda_flash_attn_ext(ctx, dst); // unsupported quant type → resident + return; + } + + const int head_dim = K->ne[0]; + // 3c-5: n_kv is the TOTAL logical key span — window (src[1]) + tail (src[5]). + // Single region keeps n_kv == K->ne[1] (n_tail==0) → byte-identical downstream. + const int n_win = K->ne[1]; + const int n_tail = two_region ? (int) K_tail->ne[1] : 0; + const int n_kv = n_win + n_tail; + const int n_head_kv = K->ne[2]; + const int n_head = Q->ne[2]; + const int n_q = Q->ne[1]; + const int n_b = Q->ne[3]; + const int gqa = n_head / n_head_kv; + const int dv = V->ne[0]; + const int n_rows = n_head * n_q * n_b; // == dst->ne[1]*ne[2]*ne[3] + + // Tile boundaries aligned to FATTN_KQ_STRIDE so every slice is a valid FA + // key length; the final tile carries the remainder. + const int n_blk = (n_kv + FATTN_KQ_STRIDE - 1) / FATTN_KQ_STRIDE; + // Resolve the auto tile count. Target ~DEFAULT_TILE_KEYS keys per tile so a slot + // holds far less than the full KV (the streaming memory win). Clamp to [1, n_blk]. + // + // opencoti-hook: f5-rolling-kv — M7-P1 (#335) launch-slack kill. At 2048 + // keys/tile a 256k spill becomes 128 tiles, each firing converter×2 + + // per-tile FA + lse (+ one final combine) ≈ hundreds of kernel launches per + // token — pure overhead, since total attention work is O(n_kv) regardless of + // how it is tiled. Coarsen to 32768 keys/tile so n_tiles at 256k drops + // 128 → ~8: the per-tile FA/lse simply process a longer key slice (no extra + // compute), the combine reduces ~8 partials instead of 128, and the + // inter-tile online-softmax stays exact (associative reduction; cosine- and + // RULER-gated). The win is launch overhead only. A VRAM cap keeps the f16 + // double-buffer slot pool bounded: each streamed key costs + // (head_dim+dv)·sizeof(half)·n_head_kv·n_b across K+V, so a coarse tile that + // would push the worst-case slot pool past the slot-pool cap (512 MiB) raises + // n_tiles back up — coarsening never OOMs. + const int DEFAULT_TILE_KEYS = 32768; + int n_tiles = (n_kv + DEFAULT_TILE_KEYS - 1) / DEFAULT_TILE_KEYS; + // VRAM clamp so a too-coarse tile cannot OOM. per_key_slot = f16 K+V bytes a + // single streamed key occupies in one slot; the pool is n_slots (≤4 below) + // copies of tile_kv_full keys. + { + const size_t per_key_slot = + ((size_t)head_dim + (size_t)dv) * sizeof(half) * (size_t)n_head_kv * (size_t)n_b; + const size_t cap_mib = 512; + const size_t denom = per_key_slot * 4; // worst-case n_slots (S3a cap) + size_t mtk = (denom == 0) ? (size_t)n_kv : (cap_mib * 1024 * 1024) / denom; + if (mtk < (size_t)FATTN_KQ_STRIDE) mtk = (size_t)FATTN_KQ_STRIDE; + const int n_tiles_vram = (n_kv + (int)mtk - 1) / (int)mtk; + if (n_tiles < n_tiles_vram) n_tiles = n_tiles_vram; + } + if (n_tiles < 1) n_tiles = 1; + if (n_tiles > n_blk) n_tiles = n_blk; + const int blk_per_tile = (n_blk + n_tiles - 1) / n_tiles; + const int tile_kv_full = blk_per_tile * FATTN_KQ_STRIDE; + + // 3c-5: re-tile the WINDOW and TAIL SEPARATELY at the resolved tile_kv_full so no + // tile straddles the n_win boundary (a straddle tile would need both a device and + // a host source). n_tiles grows by ≤1 partial vs the unified count; the pools and + // combine below size for the updated n_tiles. resolve(t) maps t O_tiles (pool, (size_t)n_combine_tiles * nel); + ggml_cuda_pool_alloc lse_tiles(pool, (size_t)n_combine_tiles * n_rows); + + // Slot-pool staging: each tile's K/V is routed through a round-robin GPU slot, + // enabling the copy/compute double buffer below. A 2-slot double buffer is the + // product default. Slot lift requires a contiguous resident K/V; otherwise the + // op falls back to the in-place view path. The lse kernel reads K as half, so + // the slot path assumes f16 K. + int n_slots = 2; + // 3c-5: the two-region forward REQUIRES the double-buffer (window tiles on the + // device source interleave with host-tail DMA) — never honour a slot-pool + // disable for it. + if (two_region && n_slots < 2) n_slots = 2; + // opencoti-hook: f5-rolling-kv — S3c scheduler-bypass host streaming. The + // scheduler's device copy is suppressed for this op (ggml-backend.cpp), so its + // K/V point STRAIGHT at the pinned-host CPU-half: a + // permuted (non-contiguous) view whose key rows are strided by + // nb[1]=n_head_cpu*row(head_dim), NOT row(head_dim). The packed 1-D lift + // below assumes contiguous key rows (the device-dup layout), so a host source + // MUST be lifted with a strided cudaMemcpy2DAsync (src pitch = nb[1]); and the + // slot path MUST be forced for it (the S2 in-place path would hand a host + // pointer to the FA kernel as device → wrong). Device sources keep the exact + // S3a/S3b 1-D lift → byte-identical. Detect the host source by pointer type. + bool host_src = false; + { + cudaPointerAttributes kattr = {}; + cudaPointerAttributes vattr = {}; + const cudaError_t ek = cudaPointerGetAttributes(&kattr, K->data); + const cudaError_t ev = cudaPointerGetAttributes(&vattr, V->data); + if (ek == cudaSuccess && ev == cudaSuccess && + kattr.type == cudaMemoryTypeHost && vattr.type == cudaMemoryTypeHost) { + host_src = true; + } + cudaGetLastError(); // clear any sticky error from probing an unregistered ptr + } + + // 3c-5: probe the TAIL source independently (the window is device-resident → + // host_src false; the overflow tail is pinned-host → host_tail true). resolve(t) + // hands each path its own host flag so window tiles take the 1-D device lift and + // tail tiles take the 2-D pinned-host lift. + bool host_tail = false; + if (two_region) { + cudaPointerAttributes ta = {}; + if (cudaPointerGetAttributes(&ta, K_tail->data) == cudaSuccess && + ta.type == cudaMemoryTypeHost) { + host_tail = true; + } + cudaGetLastError(); + } + + // 3c-5: per-tile region descriptor. Single region returns exactly the legacy + // (kv0 = t·tile_kv_full) mapping → byte-identical. Two region: tiles [0,n_win_tiles) + // cover the device window at local==global offset; tiles [n_win_tiles,n_tiles) + // cover the host tail at local = (t−n_win_tiles)·tile_kv_full, global = n_win+local. + // The mask is sliced by the GLOBAL offset; the source is read at the LOCAL offset. + struct tile_region_t { + const ggml_tensor * K; + const ggml_tensor * V; + bool host; + int local_kv0; + int global_kv0; + int this_kv; + }; + auto resolve = [&](int t) -> tile_region_t { + if (!two_region) { + const int kv0 = t * tile_kv_full; + const int tk = kv0 >= n_kv ? 0 : min(tile_kv_full, n_kv - kv0); + return { K, V, host_src, kv0, kv0, tk }; + } + if (t < n_win_tiles) { + const int loc = t * tile_kv_full; + const int tk = loc >= n_win ? 0 : min(tile_kv_full, n_win - loc); + return { K, V, host_src, loc, loc, tk }; + } + const int j = t - n_win_tiles; + const int loc = j * tile_kv_full; + const int tk = loc >= n_tail ? 0 : min(tile_kv_full, n_tail - loc); + return { K_tail, V_tail, host_tail, loc, n_win + loc, tk }; + }; + + // A single resolved tile over a device/resident source is exactly + // plain FA — the byte-identical S1 anchor, and the correct resident path for a + // quantized source too. A host (bypass) source must take the slot lift even at + // one tile (plain FA would read host memory as device), so guard on !host_src. + // S3-b (#353): never short-circuit in CPU_FA_TAIL mode — a lone window tile is + // plain FA over the BULK only; the appended partial tile (the overflow) must + // still be converted and combined. Fall through to the tiling+combine path. + if (n_tiles == 1 && !host_src && !tail_is_partial) { + ggml_cuda_flash_attn_ext(ctx, dst); + return; + } + + // S3d: a quantized K/V MUST go through a slot (the S2 in-place path and the + // f16 lse kernel read K->data as f16 → wrong for quant), so !all_f16 forces + // the slot path; with no slot to dequant into it falls back to resident. + const bool use_slots = n_slots > 0 && + (two_region || host_src || !all_f16 || (ggml_is_contiguous(K) && ggml_is_contiguous(V))); + if (!all_f16 && !use_slots) { + ggml_cuda_flash_attn_ext(ctx, dst); // quant but no slot → resident + return; + } + + if (use_slots) { + if (n_slots > n_tiles) n_slots = n_tiles; + if (n_slots > 4) n_slots = 4; // S3a pool cap + + // S3d: the slot is ALWAYS f16 (dequant-on-lift), so it is sized for f16 + // regardless of K/V storage type. For f16 source this equals the prior + // ggml_row_size(F16,·) → byte-identical anchor preserved. + const size_t k_row = (size_t)head_dim * sizeof(half); + const size_t v_row = (size_t)dv * sizeof(half); + const size_t k_slot = k_row * (size_t)tile_kv_full * n_head_kv * n_b; + const size_t v_slot = v_row * (size_t)tile_kv_full * n_head_kv * n_b; + // S3d: type-size of the SOURCE (storage) type → nc-converter strides are + // s = nb/ts (blocks-per-row for quant, elements for f16). Unused on the + // all_f16 fast path (the memcpy lift) but cheap to compute. + const size_t ts_k = ggml_type_size(K->type); + const size_t ts_v = ggml_type_size(V->type); + ggml_cuda_pool_alloc slotK(pool, (size_t)n_slots * k_slot); + ggml_cuda_pool_alloc slotV(pool, (size_t)n_slots * v_slot); + + // R2-b S3b (#312): with >=2 slots the op runs a double-buffer ping-pong — the + // per-tile lift is issued on a dedicated copy_stream while FA/lse for the + // current tile run on ctx.stream(), so tile t+1's DMA overlaps tile t's + // compute. copy_done/compute_done events (pre-created, record/wait only) order + // the two streams; we never cudaEventCreate inside the op so the path stays + // cuda_graph-capturable (bug-251 was event *creation* during capture). + // copy_stream is the fixed top index (GGML_CUDA_MAX_STREAMS-1), clear of the + // QKV concurrent-event fan-out (streams 1-3, which join before this op) and + // NEO's stream 1 — the collision that was bug-253. <2 slots → S3a single-stream. + // 3c-5: two-region always runs the overlap path (the S3a/S2 single-stream + // branches use the uniform kv0=t·tile_kv_full mapping, which would straddle + // the window/tail boundary). n_slots≥2 is already forced above. + const bool use_overlap = two_region || (n_slots >= 2); + + if (use_overlap) { + const int cs_idx = GGML_CUDA_MAX_STREAMS - 1; // fixed, collision-free + GGML_ASSERT(cs_idx >= 1); + cudaStream_t cs = ctx.stream(ctx.device, cs_idx); + + // Pre-created event pool (>= S3a slot cap 4). cudaEventDisableTiming; + // created ONCE during the eager warmup (op dispatch is single-thread, + // the first eval predates capture) → never a capture-time cudaEventCreate. + static cudaEvent_t copy_done[4] = {}; + static cudaEvent_t compute_done[4] = {}; + static bool ev_init = false; + if (!ev_init) { + for (int s = 0; s < 4; ++s) { + CUDA_CHECK(cudaEventCreateWithFlags(©_done[s], cudaEventDisableTiming)); + CUDA_CHECK(cudaEventCreateWithFlags(&compute_done[s], cudaEventDisableTiming)); + } + ev_init = true; + } + + // bug-255: the slot pool buffers are reclaimed/re-handed across op + // invocations (every layer reuses the same pool addresses). In S3a + // copy+FA shared one stream so layer L+1's copy was implicitly ordered + // after layer L's FA. With the copy moved to copy_stream, layer L+1's + // lift could overwrite a slot while layer L's FA (on the compute + // stream) is still reading it → cross-op WAR. Order copy_stream after + // all prior compute-stream work so a reused slot is never clobbered + // mid-read. (One event per op entry; intra-op order is the ping-pong.) + static cudaEvent_t cs_sync = nullptr; + if (cs_sync == nullptr) { + CUDA_CHECK(cudaEventCreateWithFlags(&cs_sync, cudaEventDisableTiming)); + } + CUDA_CHECK(cudaEventRecord(cs_sync, stream)); + CUDA_CHECK(cudaStreamWaitEvent(cs, cs_sync, 0)); + + // Lift tile t's K/V into its slot on the copy_stream (cudaMemcpyDefault + // → UVA H2D/D2D). Same packed 1-D-per-(b,h) layout as the S3a path. + auto lift = [&](int t) { + const tile_region_t tr = resolve(t); + const int this_kv = tr.this_kv; + if (this_kv <= 0) return; + const ggml_tensor * tK = tr.K; // 3c-5: window (device) | tail (host) + const ggml_tensor * tV = tr.V; + const int loc = tr.local_kv0; // region-local source offset + // 3c-5: the all-f16 lift packs head hh's keys into a contiguous slot. + // The packed 1-D copy is valid ONLY when the source already has each + // head's keys contiguous (nb[1]==row). The position-window device view + // (permuted, ALL heads per cell → nb[1]==n_embd_gqa*ts != row) is NOT, + // so it MUST take the 2-D strided lift like the host tail. Branch on + // contiguity for two_region; single region keeps the exact host_src + // predicate (byte-identical: there a host src is always the strided + // CPU-half and a device src is always the contiguous scheduler-dup). + const bool use_2d = two_region + ? ((size_t) tK->nb[1] != k_row || (size_t) tV->nb[1] != v_row) + : tr.host; + const int slot = t % n_slots; + const size_t k_nb2 = (size_t)this_kv * k_row; + const size_t v_nb2 = (size_t)this_kv * v_row; + char * sK = slotK.ptr + (size_t)slot * k_slot; + char * sV = slotV.ptr + (size_t)slot * v_slot; + if (!all_f16) { + // S3d dequant-on-lift: ONE nc-converter launch per tile reads the + // (head_dim × this_kv × n_head_kv × n_b) source sub-block at its + // native strides (s = nb/ts: blocks-per-row for quant, elements + // for f16) and writes the PACKED f16 slot. The converter's packed + // [ne00,ne01,ne02,ne03] output lands head/batch blocks at exactly + // dstoff*k_nb2 (dstoff = bb*n_head_kv+hh), so the FA/lse readers + // below are unchanged. Unifies device + pinned-host (S3c) sources + // (the kernel reads host via UVA). Slot is f16 → lse half-cast OK. + // 3c-5: tK/tV + loc select the window (device) or tail (host) source. + to_fp16_k((const char *)tK->data + (size_t)loc*tK->nb[1], (half *)sK, + head_dim, this_kv, n_head_kv, n_b, + (int64_t)(tK->nb[1]/ts_k), (int64_t)(tK->nb[2]/ts_k), (int64_t)(tK->nb[3]/ts_k), cs); + to_fp16_v((const char *)tV->data + (size_t)loc*tV->nb[1], (half *)sV, + dv, this_kv, n_head_kv, n_b, + (int64_t)(tV->nb[1]/ts_v), (int64_t)(tV->nb[2]/ts_v), (int64_t)(tV->nb[3]/ts_v), cs); + } else + for (int bb = 0; bb < n_b; ++bb) { + for (int hh = 0; hh < n_head_kv; ++hh) { + const size_t dstoff = (size_t)bb*n_head_kv + hh; + const char * srcK = (const char *)tK->data + (size_t)bb*tK->nb[3] + (size_t)hh*tK->nb[2] + (size_t)loc*tK->nb[1]; + const char * srcV = (const char *)tV->data + (size_t)bb*tV->nb[3] + (size_t)hh*tV->nb[2] + (size_t)loc*tV->nb[1]; + if (use_2d) { + // S3c: strided source (pinned-host CPU-half, OR the 3c-5 + // permuted device window — all heads per cell). Key rows + // strided by nb[1]; the 2D copy packs them into the + // contiguous slot (dst pitch == width == row), byte- + // equivalent to the device 1-D path. cudaMemcpyDefault + // (UVA) reads host or device transparently. + CUDA_CHECK(cudaMemcpy2DAsync( + sK + dstoff*k_nb2, k_row, srcK, tK->nb[1], k_row, this_kv, cudaMemcpyDefault, cs)); + CUDA_CHECK(cudaMemcpy2DAsync( + sV + dstoff*v_nb2, v_row, srcV, tV->nb[1], v_row, this_kv, cudaMemcpyDefault, cs)); + } else { + CUDA_CHECK(cudaMemcpyAsync(sK + dstoff*k_nb2, srcK, k_nb2, cudaMemcpyDefault, cs)); + CUDA_CHECK(cudaMemcpyAsync(sV + dstoff*v_nb2, srcV, v_nb2, cudaMemcpyDefault, cs)); + } + } + } + }; + + // Prologue: stage tile 0 so its compute can start as soon as it lands. + lift(0); + CUDA_CHECK(cudaEventRecord(copy_done[0], cs)); + + for (int t = 0; t < n_tiles; ++t) { + const int slot = t % n_slots; + + // Prefetch tile t+1 on copy_stream (overlaps tile t's compute). If + // its slot is being reused, first wait for the prior occupant's + // compute to release it (compute_done was recorded n_slots tiles ago). + if (t + 1 < n_tiles) { + const int nslot = (t + 1) % n_slots; + if (t + 1 >= n_slots) { + CUDA_CHECK(cudaStreamWaitEvent(cs, compute_done[nslot], 0)); + } + lift(t + 1); + CUDA_CHECK(cudaEventRecord(copy_done[nslot], cs)); + } + + // Compute tile t on ctx.stream(): wait for its copy, then FA + lse. + // 3c-5: kv0 is the GLOBAL key position (window tile: t·tile_kv_full; + // tail tile: n_win + local) → the mask slice below lands on the right + // logical column. The slot (sK/sV) was already packed by lift(t) from + // the region source, so the FA reads it region-agnostically. + const tile_region_t trc = resolve(t); + const int kv0 = trc.global_kv0; + const int this_kv = trc.this_kv; + float * lse_t = lse_tiles.ptr + (size_t)t * n_rows; + char * sK = slotK.ptr + (size_t)slot * k_slot; + char * sV = slotV.ptr + (size_t)slot * v_slot; + CUDA_CHECK(cudaStreamWaitEvent(stream, copy_done[slot], 0)); + + if (this_kv > 0) { + const size_t k_nb2 = (size_t)this_kv * k_row; + const size_t v_nb2 = (size_t)this_kv * v_row; + // S3d: the slot is always f16 and packed at k_row/v_row stride. + // Tag type=F16 (per-tile FA → (F16,F16) case) and set nb[0..1] + // to the packed-f16 layout. nb[1]=k_row also fixes the + // n_head_cpu>1 stride bug (the slot stride is k_row, not K->nb[1]); + // value-preserving on the f16/device path (there K->nb[1]==k_row). + // VALIDATED on Gemma-4 A4B (8 KV heads → frac 0.5 = 4 GPU/4 CPU, + // n_head_cpu=4, 25/30 layers streaming): RULER vt@4k q8_0 = 100%, + // byte-equal in score to resident-q8_0 ground truth (2026-06-01). + ggml_tensor Kt = *K; + Kt.type = GGML_TYPE_F16; + Kt.ne[1] = this_kv; Kt.data = sK; + Kt.nb[0] = sizeof(half); Kt.nb[1] = k_row; + Kt.nb[2] = k_nb2; Kt.nb[3] = k_nb2 * n_head_kv; + ggml_tensor Vt = *V; + Vt.type = GGML_TYPE_F16; + Vt.ne[1] = this_kv; Vt.data = sV; + Vt.nb[0] = sizeof(half); Vt.nb[1] = v_row; + Vt.nb[2] = v_nb2; Vt.nb[3] = v_nb2 * n_head_kv; + ggml_tensor Mt; + ggml_tensor * Mtp = nullptr; + if (mask) { + Mt = *mask; Mt.ne[0] = this_kv; + Mt.data = (char *)mask->data + (size_t)kv0*mask->nb[0]; + Mtp = &Mt; + } + ggml_tensor dst_t = *dst; + dst_t.op = GGML_OP_FLASH_ATTN_EXT; + dst_t.src[1] = &Kt; + dst_t.src[2] = &Vt; + dst_t.src[3] = Mtp; + dst_t.src[4] = nullptr; + dst_t.data = O_tiles.ptr + (size_t)t * nel; + // S3b HYBRID lse sourcing (opencoti-hook: f5-rolling-kv #340). + // DECODE on a finalize-guaranteed tile (n_q==1 && head_dim>256 ⇒ + // MMA; stream_k uniform fixup writes dst_lse for EVERY head-row): + // arm the consume-once channel so the per-tile FA emits lse_t + // itself, killing the recompute launch (the M7 launch-slack win). + // Else (prefill, head<=256/Qwen → VEC has no finalize) recompute + // with the softcap-aware streaming_lse_kernel. + // + // bug-263 (M7 Stage-3c spike, Gemma) — UNRESOLVED, see .wolf/buglog.json. + // WARNING: the `> 2*FATTN_KQ_STRIDE` boundary below is NOT a fix. + // A 9-config sweep proved decode (n_q==1, head_dim 512, GQA) with + // tile_kv_full == EXACTLY 512 silently corrupts output (needle lost, + // NO assert); 256/768/1024 all pass. The defect is INDEPENDENT of lse + // source — the original `>=` armed finalize at 512 (broke), and this + // `>` instead routes the 512 tile to streaming_lse_kernel recompute + // (ALSO breaks). Neither path is proven correct at this_kv∈{256,512} + // for multi-head GQA decode. The earlier "ntiles_KV==2 / finalize lse + // wrong" rationale was WRONG (Ampere head 512 uses nbatch_fa=32 → + // ntiles_KV=16 at 512, not 2). `>` is kept only because it is no worse + // than `>=`; the real mechanism is unisolated. Real Stage-3c tail tiles + // can be ≤512 at decode → MUST resolve before relying on this path. + const bool decode_lse = (n_q == 1 && head_dim > 256 && this_kv > 2*FATTN_KQ_STRIDE); + if (decode_lse) opencoti_fattn_dst_lse = lse_t; + ggml_cuda_flash_attn_ext(ctx, &dst_t); + + if (!decode_lse) { + streaming_lse_kernel<<>>( + (const char *)Q->data, sK, + mask ? (const char *)mask->data + (size_t)kv0*mask->nb[0] : nullptr, + lse_t, scale, logit_softcap, this_kv, gqa, head_dim, n_head, n_q, + Q->nb[1], Q->nb[2], Q->nb[3], k_row, k_nb2, k_nb2 * n_head_kv, + mask ? mask->nb[0] : 0, mask ? mask->nb[1] : 0); + CUDA_CHECK(cudaGetLastError()); + } + } else { + // Empty tile (this_kv==0): the kernel writes lse=-inf and never + // dereferences the slot, so the K strides are unused → pass 0. + streaming_lse_kernel<<>>( + (const char *)Q->data, sK, nullptr, + lse_t, scale, logit_softcap, 0, gqa, head_dim, n_head, n_q, + Q->nb[1], Q->nb[2], Q->nb[3], k_row, 0, 0, 0, 0); + CUDA_CHECK(cudaGetLastError()); + } + CUDA_CHECK(cudaEventRecord(compute_done[slot], stream)); + } + } else + // S3a: single compute stream. The per-tile slot lift, the FA over the + // slot, and the lse all run on ctx.stream() — stream-ordered so the FA + // reads the freshly-copied slot, fully capturable under cuda_graph, and + // free of cross-stream collisions with M2's concurrent-event streams + // (which forced the earlier garbage). The overlapping double-buffer + // (dedicated copy_stream + events) is the S3b branch above. + for (int t = 0; t < n_tiles; ++t) { + const int slot = t % n_slots; + const int kv0 = t * tile_kv_full; + const int this_kv = kv0 >= n_kv ? 0 : min(tile_kv_full, n_kv - kv0); + float * lse_t = lse_tiles.ptr + (size_t)t * n_rows; + char * sK = slotK.ptr + (size_t)slot * k_slot; + char * sV = slotV.ptr + (size_t)slot * v_slot; + + if (this_kv > 0) { + const size_t k_nb2 = (size_t)this_kv * k_row; + const size_t v_nb2 = (size_t)this_kv * v_row; + // Lift the tile into the packed slot. For a DEVICE source the + // scheduler's device-dup makes each (b,h)'s this_kv key rows + // contiguous (nb[1]==k_row) → one packed 1-D copy each (sidesteps + // cudaMemcpy2D pitch limits; nb[2] of the source is < this_kv*k_row + // when ne[2]==1). For a HOST source (S3c bypass — the permuted + // pinned-host view, nb[1]==row(n_embd_cpu) != k_row), the rows are + // strided by the full per-token n_embd; a packed 1-D copy would + // scoop adjacent heads → 2-D lift (spitch=nb[1], width=k_row). + // cudaMemcpyDefault: UVA infers H2D (pinned spill) vs D2D. + // opencoti-hook: f5-rolling-kv — S3c scheduler-bypass host lift; + // S3d dequant-on-lift (the !all_f16 branch). + if (!all_f16) { + // S3d: one nc-converter launch per tile → packed f16 slot (see + // the overlap path for the layout proof). Source strides s=nb/ts. + to_fp16_k((const char *)K->data + (size_t)kv0*K->nb[1], (half *)sK, + head_dim, this_kv, n_head_kv, n_b, + (int64_t)(K->nb[1]/ts_k), (int64_t)(K->nb[2]/ts_k), (int64_t)(K->nb[3]/ts_k), stream); + to_fp16_v((const char *)V->data + (size_t)kv0*V->nb[1], (half *)sV, + dv, this_kv, n_head_kv, n_b, + (int64_t)(V->nb[1]/ts_v), (int64_t)(V->nb[2]/ts_v), (int64_t)(V->nb[3]/ts_v), stream); + } else + for (int bb = 0; bb < n_b; ++bb) { + for (int hh = 0; hh < n_head_kv; ++hh) { + const size_t dstoff = (size_t)bb*n_head_kv + hh; + const char * srcK = (const char *)K->data + (size_t)bb*K->nb[3] + (size_t)hh*K->nb[2] + (size_t)kv0*K->nb[1]; + const char * srcV = (const char *)V->data + (size_t)bb*V->nb[3] + (size_t)hh*V->nb[2] + (size_t)kv0*V->nb[1]; + if (host_src) { + CUDA_CHECK(cudaMemcpy2DAsync( + sK + dstoff*k_nb2, k_row, srcK, K->nb[1], k_row, this_kv, cudaMemcpyDefault, stream)); + CUDA_CHECK(cudaMemcpy2DAsync( + sV + dstoff*v_nb2, v_row, srcV, V->nb[1], v_row, this_kv, cudaMemcpyDefault, stream)); + } else { + CUDA_CHECK(cudaMemcpyAsync(sK + dstoff*k_nb2, srcK, k_nb2, cudaMemcpyDefault, stream)); + CUDA_CHECK(cudaMemcpyAsync(sV + dstoff*v_nb2, srcV, v_nb2, cudaMemcpyDefault, stream)); + } + } + } + + // FA over the packed slot. S3d: slot is always f16 at k_row/v_row + // stride → tag type=F16, nb[0..1] to the packed-f16 layout (nb[1]= + // k_row also fixes the n_head_cpu>1 stride bug — VALIDATED on + // Gemma-4 n_head_cpu=4, see overlap path above; value-preserving + // on f16/device where K->nb[1]==k_row). + ggml_tensor Kt = *K; + Kt.type = GGML_TYPE_F16; + Kt.ne[1] = this_kv; Kt.data = sK; + Kt.nb[0] = sizeof(half); Kt.nb[1] = k_row; + Kt.nb[2] = k_nb2; Kt.nb[3] = k_nb2 * n_head_kv; + ggml_tensor Vt = *V; + Vt.type = GGML_TYPE_F16; + Vt.ne[1] = this_kv; Vt.data = sV; + Vt.nb[0] = sizeof(half); Vt.nb[1] = v_row; + Vt.nb[2] = v_nb2; Vt.nb[3] = v_nb2 * n_head_kv; + ggml_tensor Mt; + ggml_tensor * Mtp = nullptr; + if (mask) { + Mt = *mask; Mt.ne[0] = this_kv; + Mt.data = (char *)mask->data + (size_t)kv0*mask->nb[0]; + Mtp = &Mt; + } + ggml_tensor dst_t = *dst; + dst_t.op = GGML_OP_FLASH_ATTN_EXT; + dst_t.src[1] = &Kt; + dst_t.src[2] = &Vt; + dst_t.src[3] = Mtp; + dst_t.src[4] = nullptr; + dst_t.data = O_tiles.ptr + (size_t)t * nel; + // S3a HYBRID lse sourcing (opencoti-hook: f5-rolling-kv #340) — see + // the S3b overlap path for the full bug-263 note. DECODE on a + // finalize tile arms the channel so the FA emits lse_t; else recompute + // (softcap-aware). The `> 2*FATTN_KQ_STRIDE` boundary is NOT a fix for + // bug-263 (UNRESOLVED: decode tile_kv_full==512 corrupts on either lse + // path) — kept only as no-worse-than-`>=`. See .wolf/buglog.json. + const bool decode_lse = (n_q == 1 && head_dim > 256 && this_kv > 2*FATTN_KQ_STRIDE); + if (decode_lse) opencoti_fattn_dst_lse = lse_t; + ggml_cuda_flash_attn_ext(ctx, &dst_t); + + if (!decode_lse) { + streaming_lse_kernel<<>>( + (const char *)Q->data, sK, + mask ? (const char *)mask->data + (size_t)kv0*mask->nb[0] : nullptr, + lse_t, scale, logit_softcap, this_kv, gqa, head_dim, n_head, n_q, + Q->nb[1], Q->nb[2], Q->nb[3], k_row, k_nb2, k_nb2 * n_head_kv, + mask ? mask->nb[0] : 0, mask ? mask->nb[1] : 0); + CUDA_CHECK(cudaGetLastError()); + } + } else { + // Empty trailing tile: lse=-inf, combine skips it (no copy/FA). The + // kernel never dereferences the slot, so K strides are unused → 0. + streaming_lse_kernel<<>>( + (const char *)Q->data, sK, nullptr, + lse_t, scale, logit_softcap, 0, gqa, head_dim, n_head, n_q, + Q->nb[1], Q->nb[2], Q->nb[3], k_row, 0, 0, 0, 0); + CUDA_CHECK(cudaGetLastError()); + } + } + } else { + // S2 resident-view path (default; byte-identical gate). + for (int t = 0; t < n_tiles; ++t) { + const int kv0 = t * tile_kv_full; + float * lse_t = lse_tiles.ptr + (size_t)t * n_rows; + const int this_kv = kv0 >= n_kv ? 0 : min(tile_kv_full, n_kv - kv0); + + // S2 HYBRID lse sourcing (opencoti-hook: f5-rolling-kv #340): DECODE on a + // finalize tile arms the channel so the FA emits lse_t; the recompute below + // is then skipped. Prefill / head<=256 (VEC, no finalize) / empty tiles fall + // to the recompute (which also writes the -inf fill). The `> 2*FATTN_KQ_STRIDE` + // boundary is NOT a fix for bug-263 (UNRESOLVED: decode tile_kv_full==512 + // corrupts on either lse path) — kept only as no-worse-than-`>=`. See the + // S3b overlap path's full note + .wolf/buglog.json. + const bool decode_lse = (n_q == 1 && head_dim > 256 && this_kv > 2*FATTN_KQ_STRIDE); + if (this_kv > 0) { + // FA over the tile's key-slice → O_t (stock kernel, untouched). + ggml_tensor Kt = *K; Kt.ne[1] = this_kv; Kt.data = (char *)K->data + (size_t)kv0*K->nb[1]; + ggml_tensor Vt = *V; Vt.ne[1] = this_kv; Vt.data = (char *)V->data + (size_t)kv0*V->nb[1]; + ggml_tensor Mt; + ggml_tensor * Mtp = nullptr; + if (mask) { + Mt = *mask; Mt.ne[0] = this_kv; Mt.data = (char *)mask->data + (size_t)kv0*mask->nb[0]; + Mtp = &Mt; + } + ggml_tensor dst_t = *dst; + dst_t.op = GGML_OP_FLASH_ATTN_EXT; // bona-fide FA op for the sub-call + dst_t.src[1] = &Kt; + dst_t.src[2] = &Vt; + dst_t.src[3] = Mtp; + dst_t.src[4] = nullptr; + dst_t.data = O_tiles.ptr + (size_t)t * nel; + if (decode_lse) opencoti_fattn_dst_lse = lse_t; + ggml_cuda_flash_attn_ext(ctx, &dst_t); + } + + // lse_t (this_kv==0 → empty tile: kernel writes -inf, combine ignores it). + if (!decode_lse) { + streaming_lse_kernel<<>>( + (const char *)Q->data, + (const char *)K->data + (size_t)kv0*K->nb[1], + mask ? (const char *)mask->data + (size_t)kv0*mask->nb[0] : nullptr, + lse_t, scale, logit_softcap, this_kv, gqa, head_dim, n_head, n_q, + Q->nb[1], Q->nb[2], Q->nb[3], K->nb[1], K->nb[2], K->nb[3], + mask ? mask->nb[0] : 0, mask ? mask->nb[1] : 0); + CUDA_CHECK(cudaGetLastError()); + } + } + } + + const int64_t total = (int64_t)dv * n_rows; + const int threads = 256; + const int64_t blocks = (total + threads - 1) / threads; + + // S3-b (#353) CPU_FA_TAIL: after the window tiles [0,n_tiles), convert the + // precomputed host partial into the final combine tile [n_tiles]. The partial + // (src[5] = [2+DV, n_rows] [M,S,VKQ]) is staged to device with a single UVA + // cudaMemcpyAsync — only this tiny buffer (≪ the tail K/V) crosses PCIe. The + // window leg already filled tiles [0,n_tiles) above; the combine below then + // merges window ⊕ tail through the unchanged online-softmax reduction. + if (tail_is_partial) { + const size_t pcount = (size_t)(2 + dv) * (size_t)n_rows; + ggml_cuda_pool_alloc p_dev(pool, pcount); + CUDA_CHECK(cudaMemcpyAsync(p_dev.ptr, tail_partial->data, pcount * sizeof(float), + cudaMemcpyDefault, stream)); + tail_partial_to_tile_kernel<<<(unsigned)blocks, threads, 0, stream>>>( + p_dev.ptr, O_tiles.ptr + (size_t)n_tiles * nel, + lse_tiles.ptr + (size_t)n_tiles * n_rows, dv, n_rows); + CUDA_CHECK(cudaGetLastError()); + } + + streaming_combine_kernel<<<(unsigned)blocks, threads, 0, stream>>>( + O_tiles.ptr, lse_tiles.ptr, (float *)dst->data, n_combine_tiles, dv, n_rows); + CUDA_CHECK(cudaGetLastError()); +} diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn.cuh --- a/llama.cpp/ggml/src/ggml-cuda/fattn.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cuh @@ -3,3 +3,9 @@ void ggml_cuda_flash_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst); bool ggml_cuda_flash_attn_ext_supported(int device, const ggml_tensor * dst); + +// opencoti-hook: f5-rolling-kv — W4 M7-D #304 / R2-b S2 (#311). +// Streaming flash-attention: key-axis tiling + LSE-weighted online-softmax +// combine over the stock FA kernels (which stay untouched). Falls back to +// plain FA for the single-tile / unsupported-scoring cases. +void ggml_cuda_streaming_flash_attn(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu --- a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -103,6 +103,11 @@ void ggml_cuda_error(const char * stmt, const char * func, const char * file, in GGML_LOG_ERROR(GGML_CUDA_NAME " error: %s\n", msg); GGML_LOG_ERROR(" current device: %d, in function %s at %s:%d\n", id, func, file, line); GGML_LOG_ERROR(" %s\n", stmt); + // opencoti: the llamafile log callback swallows GGML_LOG_ERROR on the abort + // path, so the failing stmt/msg never reaches the server log. Mirror it to + // stderr directly (flushed) so CUDA faults are diagnosable. Harmless. + fprintf(stderr, "[opencoti CUDA] error: %s | stmt: %s | %s:%d (dev %d)\n", msg, stmt, file, line, id); + fflush(stderr); // abort with GGML_ABORT to get a stack trace GGML_ABORT(GGML_CUDA_NAME " error"); } @@ -3024,6 +3029,16 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg ggml_cuda_mul_mat_vec_q(ctx, up_exps, cur, ids, dst, &fusion_data); } break; + case GGML_OP_STREAMING_FLASH_ATTN: + // opencoti-hook: f5-rolling-kv — W4 M7-D #304 / R2-b S2 (#311). + // Key-axis tiling + LSE-weighted online-softmax over the STOCK FA + // kernels (left untouched). Single-tile short-circuits to plain FA + // (= R2-b S1, byte-identical). NB: the mma kernel our 3090 selects + // normalises in-kernel and combines via stream_k (NOT + // flash_attn_combine_results), so the cross-tile combine is owned + // here in F32, not delegated to the kernel. See fattn.cu. + ggml_cuda_streaming_flash_attn(ctx, dst); + break; case GGML_OP_OUT_PROD: ggml_cuda_out_prod(ctx, dst); break; @@ -5579,6 +5594,35 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de #endif // GGML_USE_MUSA case GGML_OP_FLASH_ATTN_EXT: return ggml_cuda_flash_attn_ext_supported(dev_ctx->device, op); + case GGML_OP_STREAMING_FLASH_ATTN: { + // opencoti-hook: f5-rolling-kv — bug-259. The streaming forward repacks + // its (possibly non-contiguous / pinned-host) K/V into CONTIGUOUS f16 + // slots before each tile's stock FA, so delegating to + // ggml_cuda_flash_attn_ext_supported is WRONG: that check rejects any + // src whose nb[i] is not 16-aligned (true for the S3c host-permuted + // views and for some -fit/reserve trial graphs), which made the op + // report unsupported on CUDA at boot → it fell to the CPU backend (no + // kernel → abort) or to no backend (assignment assert). Validate only + // what the repacked slot FA actually needs — a supported head_dim with + // a matching V width; K/V storage type is irrelevant (dequantized to + // f16 on the lift). Scoring (max_bias/softcap/sinks) is handled by the + // op's own resident fallback at run time. + const ggml_tensor * K = op->src[1]; + const ggml_tensor * V = op->src[2]; + const int64_t dk = K->ne[0]; + const int64_t dv = V->ne[0]; + switch (dk) { + case 40: case 64: case 72: case 80: case 96: + case 112: case 128: case 256: + return dv == dk; + case 512: + return dv == dk; + case 576: + return dv == 512; + default: + return false; + } + } case GGML_OP_CROSS_ENTROPY_LOSS: case GGML_OP_CROSS_ENTROPY_LOSS_BACK: case GGML_OP_OPT_STEP_ADAMW: diff --git a/llama.cpp/ggml/src/ggml.c b/llama.cpp/ggml/src/ggml.c --- a/llama.cpp/ggml/src/ggml.c +++ b/llama.cpp/ggml/src/ggml.c @@ -1079,9 +1079,11 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "GLU", "MOE_FUSED_UP_GATE", + "STREAMING_FLASH_ATTN", + "FLASH_ATTN_TAIL_PARTIAL", }; -static_assert(GGML_OP_COUNT == 97, "GGML_OP_COUNT != 97"); +static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1190,9 +1192,11 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "glu(x)", "moe_up_gate(up,gate,x)", + "streaming_flash_attn(q,k,v)", + "flash_attn_tail_partial(q,k,v)", }; -static_assert(GGML_OP_COUNT == 97, "GGML_OP_COUNT != 97"); +static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -5419,10 +5423,166 @@ struct ggml_tensor * ggml_flash_attn_ext( return result; } +// opencoti F5 M7 Rolling KV (W4 M7-D #304) — streaming flash attention. +// Construction mirrors ggml_flash_attn_ext exactly (same src layout, same +// op_params {scale, max_bias, logit_softcap}, same F32 permuted output) so +// the CUDA forward can reuse the proven FA pre/post-processing and only +// differ in how it sweeps the key dimension (tiled + online-softmax). The +// op carries no extra params at R2-a; tile/slot sizing is decided inside the +// CUDA forward from device free-VRAM. Precision is forced to F32 by the +// graph builder via ggml_flash_attn_ext_set_prec (shared setter). +struct ggml_tensor * ggml_streaming_flash_attn( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap) { + GGML_ASSERT(ggml_can_mul_mat(k, q)); + + GGML_ASSERT(q->ne[3] == k->ne[3]); + GGML_ASSERT(q->ne[3] == v->ne[3]); + + if (mask) { + GGML_ASSERT(mask->type == GGML_TYPE_F16); + GGML_ASSERT(ggml_is_contiguous(mask)); + GGML_ASSERT(q->ne[2] % mask->ne[2] == 0); + GGML_ASSERT(q->ne[3] % mask->ne[3] == 0); + } + + if (max_bias > 0.0f) { + GGML_ASSERT(mask); + } + + // permute(0, 2, 1, 3) — identical to flash_attn_ext + int64_t ne[4] = { v->ne[0], q->ne[2], q->ne[1], q->ne[3] }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + + float params[] = { scale, max_bias, logit_softcap }; + ggml_set_op_params(result, params, sizeof(params)); + + result->op = GGML_OP_STREAMING_FLASH_ATTN; + result->src[0] = q; + result->src[1] = k; + result->src[2] = v; + result->src[3] = mask; + + return result; +} + +// opencoti F5 M7 Stage 3c-5 (#341) — position-window streaming FA. Same op as +// ggml_streaming_flash_attn but with the resident device window K/V in src[1]/ +// src[2] and the pinned-host overflow tail K/V in src[5]/src[6]. The CUDA forward +// tiles the window (device, in-place/D2D) AND the tail (host, H2D-streamed) into +// ONE online-softmax combine, so the bulk stays resident while only the overflow +// streams. mask spans the full logical key range [0, n_win + n_tail). tail null ⇒ +// identical to ggml_streaming_flash_attn (single resident region). +struct ggml_tensor * ggml_streaming_flash_attn_window( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k_win, + struct ggml_tensor * v_win, + struct ggml_tensor * k_tail, + struct ggml_tensor * v_tail, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap) { + GGML_ASSERT(ggml_can_mul_mat(k_win, q)); + GGML_ASSERT(q->ne[3] == k_win->ne[3]); + GGML_ASSERT(q->ne[3] == v_win->ne[3]); + if (k_tail) { + GGML_ASSERT(v_tail); + GGML_ASSERT(k_tail->ne[0] == k_win->ne[0]); + GGML_ASSERT(v_tail->ne[0] == v_win->ne[0]); + GGML_ASSERT(k_tail->ne[2] == k_win->ne[2]); + } + if (mask) { + GGML_ASSERT(mask->type == GGML_TYPE_F16); + GGML_ASSERT(ggml_is_contiguous(mask)); + GGML_ASSERT(q->ne[2] % mask->ne[2] == 0); + GGML_ASSERT(q->ne[3] % mask->ne[3] == 0); + } + if (max_bias > 0.0f) { + GGML_ASSERT(mask); + } + + int64_t ne[4] = { v_win->ne[0], q->ne[2], q->ne[1], q->ne[3] }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + + float params[] = { scale, max_bias, logit_softcap }; + ggml_set_op_params(result, params, sizeof(params)); + + result->op = GGML_OP_STREAMING_FLASH_ATTN; + result->src[0] = q; + result->src[1] = k_win; + result->src[2] = v_win; + result->src[3] = mask; + // src[4] reserved for sinks (ggml_flash_attn_ext_add_sinks) + result->src[5] = k_tail; + result->src[6] = v_tail; + + return result; +} + +// opencoti F5 M7 Stage 3d-S3 (#353) — CPU tail-partial emitter for the CPU_FA_TAIL +// tactic. Runs flash-attention over a (host-resident) tail K/V in ONE online-softmax +// pass and writes the UN-normalized partial [M, S, VKQ] per output row; dst shape is +// [2+DV, n_head, n_q, n_b]. That is exactly the partial the CUDA online-softmax combine +// (fattn.cu streaming_combine_kernel) consumes after a trivial per-row normalize: +// O_row = VKQ_row / S_row , lse_row = M_row + logf(S_row). +// The window leg's (O,lse) is computed on-device by the streaming op; only this tiny +// partial (2+DV floats per head) crosses PCIe instead of the whole tail K/V — the +// CPU_FA_TAIL memory win on slow links. NO sinks here (src[4] stays null): sinks are +// single-counted on the window partial only. Decode-only consumer (n_q==1): the CPU +// per-row order (q-fastest within head) coincides with the combine's head-fastest row +// only when n_q==1; the dispatch routes CPU_FA_TAIL for decode and falls back otherwise. +struct ggml_tensor * ggml_flash_attn_ext_tail_partial( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k, + struct ggml_tensor * v, + struct ggml_tensor * mask, + float scale, + float max_bias, + float logit_softcap) { + GGML_ASSERT(ggml_can_mul_mat(k, q)); + GGML_ASSERT(q->ne[3] == k->ne[3]); + GGML_ASSERT(q->ne[3] == v->ne[3]); + if (mask) { + GGML_ASSERT(mask->type == GGML_TYPE_F16); + GGML_ASSERT(ggml_is_contiguous(mask)); + GGML_ASSERT(q->ne[2] % mask->ne[2] == 0); + GGML_ASSERT(q->ne[3] % mask->ne[3] == 0); + } + if (max_bias > 0.0f) { + GGML_ASSERT(mask); + } + + // dst rows = (head, q-token, batch); the 2 leading slots hold (M, S) ahead of VKQ. + int64_t ne[4] = { 2 + v->ne[0], q->ne[2], q->ne[1], q->ne[3] }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + + float params[] = { scale, max_bias, logit_softcap }; + ggml_set_op_params(result, params, sizeof(params)); + + result->op = GGML_OP_FLASH_ATTN_TAIL_PARTIAL; + result->src[0] = q; + result->src[1] = k; + result->src[2] = v; + result->src[3] = mask; + + return result; +} + void ggml_flash_attn_ext_set_prec( struct ggml_tensor * a, enum ggml_prec prec) { - GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT); + // opencoti F5 M7 Rolling KV (W4 M7-D #304) — streaming-FA shares FA's + // op_params layout (prec at index 3), so this setter serves both. + GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT || a->op == GGML_OP_STREAMING_FLASH_ATTN); const int32_t prec_i32 = (int32_t) prec; @@ -5431,7 +5591,7 @@ void ggml_flash_attn_ext_set_prec( enum ggml_prec ggml_flash_attn_ext_get_prec( const struct ggml_tensor * a) { - GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT); + GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT || a->op == GGML_OP_STREAMING_FLASH_ATTN); const int32_t prec_i32 = ggml_get_op_params_i32(a, 3); @@ -5446,7 +5606,7 @@ void ggml_flash_attn_ext_add_sinks( return; } - GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT); + GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT || a->op == GGML_OP_STREAMING_FLASH_ATTN); GGML_ASSERT(a->src[4] == NULL); GGML_ASSERT(a->src[0]->ne[2] == sinks->ne[0]); GGML_ASSERT(sinks->type == GGML_TYPE_F32); diff --git a/llama.cpp/ggml/src/iqk/iqk_flash_attn.cpp b/llama.cpp/ggml/src/iqk/iqk_flash_attn.cpp --- a/llama.cpp/ggml/src/iqk/iqk_flash_attn.cpp +++ b/llama.cpp/ggml/src/iqk/iqk_flash_attn.cpp @@ -138,6 +138,60 @@ static inline bool are_kv_types_supported(ggml_type type_k, ggml_type type_v) { return it_k != supported.end() && it_v != supported.end(); } +// opencoti-hook: f5-rolling-kv — W4 M7 S3 (#354). iqk-backed CPU tail flash-attention +// partial for the CPU_FA_TAIL rolling-KV tactic. See ggml-include/ggml-iqk-flash-attn.h. +// One q row (nq1=1) over the host tail; emits the UN-normalized [M, S, VKQ] partial the +// CUDA streaming combine consumes. CONVENTION (verified fa/iqk_fa_templates.h:1185 +// normalize_and_store): when M AND S are non-null, iqk_flash_attn_impl returns the +// UN-normalized qkv directly (memcpy of qkv_cache = Sum exp(s-M)*V, NO divide by S) plus +// (M=max logit, S=sum exp(s-M)) — exactly the [M,S,VKQ] the combine wants (O=VKQ/S, +// lse=M+logS), byte-compatible with ops.cpp _f16_one_chunk write_partials. So we copy qkv +// straight through; do NOT multiply by S (that double-counts and produces garbage — the +// bug-354 niah=0.0). softcap is folded by the +// caller (scale/=softcap when softcap!=0) exactly as the iqk_flash_attn_noalibi call site; +// for Gemma-4 global layers softcap==0. No sinks here (single-counted on the GPU window +// leg). Returns false on unsupported (type, shape) so the caller falls back to +// _f16_one_chunk. The ggml thread partition (one row per call) replaces iqk's internal +// threading: iqk_flash_attn_impl is the single-call kernel (no barrier / wdata). +extern "C" IQK_API bool iqk_flash_attn_tail_partial(int type_k, int type_v, int Dk, int Dv, + int nk1, int stride_k, int stride_v, int stride_m, + const float * q, const void * k, const void * v, const void * mask, + float scale, float softcap, float * partial) { + if (!mask || nk1 % 32 != 0) { + return false; // iqk assumes a non-null mask and nk multiple of 32 + } + if (Dv > 512) { + return false; // qkv scratch bound below + } + if (!are_kv_types_supported((ggml_type) type_k, (ggml_type) type_v)) { + return false; + } + + float qkv[512]; // UN-normalized Sum exp(s-M)*V for the single row (M,S non-null mode) + float M = -INFINITY; // running max logit (scaled+masked) + float S = 0.0f; // sum exp(logit - M) + + const bool ok = iqk_flash_attn_impl( + type_k, type_v, Dk, Dv, + /*nq1=*/1, nk1, + /*stride_q=*/Dk * (int) sizeof(float), stride_k, stride_v, stride_m, + /*stride_qkv=*/Dv, + q, k, v, mask, /*sinksf=*/nullptr, /*nsinks=*/0, + scale, softcap, qkv, &M, &S); + if (!ok) { + return false; + } + + partial[0] = M; + partial[1] = S; + // qkv is ALREADY the un-normalized VKQ (M,S non-null path of normalize_and_store) — copy + // straight through. (The previous `qkv[d] * S` double-counted S → niah=0.0, bug-354.) + for (int d = 0; d < Dv; ++d) { + partial[2 + d] = qkv[d]; + } + return true; +} + // TODO: get the ggml_type enum here without polution // extern "C" IQK_API bool iqk_flash_attn_noalibi(int type_q, int type_mask, float max_bias, diff --git a/llama.cpp/include/llama.h b/llama.cpp/include/llama.h --- a/llama.cpp/include/llama.h +++ b/llama.cpp/include/llama.h @@ -358,6 +358,17 @@ extern "C" { // streamed back per step). 1.0 = off. Only effective for offloaded // layers with flash-attention (non-transposed V cache). float headinfer_gpu_heads_frac; + // opencoti F5 M7 Rolling KV — Rung 0 auto residency — see docs/features/rolling_kv.md + // VRAM budget cap (MiB) for auto KV residency. 0 = use all free VRAM minus a + // compute reserve (maximize). Consulted only when headinfer_gpu_heads_frac < 0. + uint32_t vram_target_mib; + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304). Effective H2D PCIe + // bandwidth (GB/s) for streaming-FA tile sizing. 0 = unknown. + float pcie_bw_gbps; + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob. + // 0 = auto (POSITION_WINDOW when eligible, else M2 head split), 1 = head + // (force M2 split), 2 = window (force POSITION_WINDOW, warn+fallback if ineligible). + uint32_t kv_residency_mode; // opencoti F5 M3 neo — see docs/features/advanced_kv.md // 0 = off (default), 1 = on, 2 = auto. Effective only when the // M2 head split is active for a given layer. diff --git a/llama.cpp/src/llama-context.cpp b/llama.cpp/src/llama-context.cpp --- a/llama.cpp/src/llama-context.cpp +++ b/llama.cpp/src/llama-context.cpp @@ -64,6 +64,12 @@ llama_context::llama_context( cparams.slot_shrink_idle_ms = params.slot_shrink_idle_ms; // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md cparams.headinfer_gpu_heads_frac = params.headinfer_gpu_heads_frac; + // opencoti F5 M7 Rolling KV — Rung 0 — see docs/features/rolling_kv.md + cparams.vram_target_mib = params.vram_target_mib; + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304) — see docs/features/rolling_kv.md + cparams.pcie_bw_gbps = params.pcie_bw_gbps; + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency knob — see docs/features/rolling_kv.md + cparams.kv_residency_mode = params.kv_residency_mode; // opencoti F5 M3 neo — see docs/features/advanced_kv.md cparams.neo_pipeline_mode = params.neo_pipeline_mode; // opencoti F5 W3 #292 fused-MoE — see docs/features/advanced_kv.md @@ -485,7 +491,21 @@ void llama_context::sched_reserve() { ggml_backend_dev_t device_fa = ggml_backend_get_device(ggml_backend_sched_get_tensor_backend(sched.get(), n)); // TODO: instead of the tensor names, use a map to keep track of which (FA) tensors belong to which layer - GGML_ASSERT(strncmp(n->name, LLAMA_TENSOR_NAME_FATTN "-", prefix_len) == 0); + // opencoti-hook: f5-rolling-kv — see docs/features/rolling_kv.md. + // The M2 head-split (0020), M3 NEO (0030), and M7 GPU_STREAM (0070) + // paths emit per-layer split FA ops named "__fattn___g-" / + // "__fattn___c-" — two FA ops per split layer, intentionally on + // DIFFERENT devices (the CPU half runs the iqk CPU-FA engine). + // Upstream's auto-FA resolver assumes exactly one "__fattn__-" + // op per layer on that layer's device, so a split op both fails this + // name parse (the original GGML_ASSERT crashed at boot under + // --flash-attn auto) and would trip a false device-mismatch that + // disables FA. Skip any FA op that is not the plain "__fattn__-" + // form; the split ops are validated by their own construction gates, + // which require FA. Non-split layers still get the device check. + if (strncmp(n->name, LLAMA_TENSOR_NAME_FATTN "-", prefix_len) != 0) { + continue; + } const int il = std::stoi(n->name + prefix_len); ggml_backend_dev_t device_kv = model.dev_layer(il); if (device_fa != device_kv) { @@ -3360,6 +3380,9 @@ llama_context_params llama_context_default_params() { /*.slot_initial_ctx =*/ 0, /*.slot_shrink_idle_ms =*/ 0, /*.headinfer_gpu_heads_frac =*/ 1.0f, + /*.vram_target_mib =*/ 0, // opencoti F5 M7 Rolling KV — 0 = maximize + /*.pcie_bw_gbps =*/ 0.0f, // opencoti F5 M7 Rolling KV — 0 = unknown + /*.kv_residency_mode =*/ 0, // opencoti F5 M7 Rolling KV — 0 = auto /*.neo_pipeline_mode =*/ 0, // opencoti F5 M3 neo — off by default /*.fused_moe_up_gate =*/ false, // opencoti F5 W3 #292 — off by default /*.n_rs_seq =*/ 0, diff --git a/llama.cpp/src/llama-cparams.h b/llama.cpp/src/llama-cparams.h --- a/llama.cpp/src/llama-cparams.h +++ b/llama.cpp/src/llama-cparams.h @@ -23,6 +23,18 @@ struct llama_cparams { uint32_t slot_shrink_idle_ms; // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md float headinfer_gpu_heads_frac; + // opencoti F5 M7 Rolling KV — Rung 0 auto residency — see docs/features/rolling_kv.md + // VRAM budget cap (MiB) for auto KV residency. 0 = use all free VRAM minus a + // compute reserve (maximize). Consulted only when headinfer_gpu_heads_frac < 0. + uint32_t vram_target_mib; + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304). Effective H2D PCIe + // bandwidth (GB/s), resolved at boot from the ReBAR probe / sysfs link + // (common/pcie-profile). Read at the streaming-FA dispatch to size tiles + // and the double-buffer; 0 = unknown (kernel uses a conservative default). + float pcie_bw_gbps; + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob. + // 0 = auto (POSITION_WINDOW when eligible, else M2 head split), 1 = head, 2 = window. + uint32_t kv_residency_mode; // opencoti F5 M3 neo — asymmetric GPU/CPU pipelining of attention. // 0 = off (M2 concat path runs; binary equivalent to upstream when // headinfer_gpu_heads_frac == 1.0). diff --git a/llama.cpp/src/llama-graph.cpp b/llama.cpp/src/llama-graph.cpp --- a/llama.cpp/src/llama-graph.cpp +++ b/llama.cpp/src/llama-graph.cpp @@ -2113,6 +2113,210 @@ ggml_tensor * llm_graph_context::build_attn_mha( return cur; } +// opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304), S1. GPU_STREAM tactic. +// Structurally identical to build_attn_mha_neo: the device head-group runs a +// resident ggml_flash_attn_ext, the host (spilled) head-group runs the custom +// GGML_OP_STREAMING_FLASH_ATTN op, and the two outputs are concatenated on the +// head axis. At S1 the streaming op's CUDA forward is still +// ggml_cuda_flash_attn_ext verbatim, so the output is BYTE-IDENTICAL to the NEO +// path (which builds the same graph with a plain flash_attn_ext on the host +// half; NEO's stream-pair registration only affects scheduling, not output) — +// that is the S1 gate. S2 makes the streaming op tile the host head-group's KEY +// dimension into a device slot pool with a cross-tile online-softmax (running +// max + denom in F32), so the spilled KV need not fit one resident slot; S3 +// adds the copy/compute double-buffer. The host head-group is already +// pinned-host (the M2 split allocates k_cpu/v_cpu with the pinned cpu_buft), so +// streaming rides the M2 residency — no whole-layer host reorder. +ggml_tensor * llm_graph_context::build_attn_mha_streaming( + ggml_tensor * q, + ggml_tensor * k_g, + ggml_tensor * v_g, + ggml_tensor * k_c, + ggml_tensor * v_c, + uint32_t gpu_heads, + ggml_tensor * kq_b, + ggml_tensor * kq_mask, + ggml_tensor * sinks, + ggml_tensor * v_mla, + float kq_scale, + int il) const { + GGML_ASSERT(cparams.flash_attn && "GPU_STREAM tactic requires flash-attention"); + GGML_ASSERT(kq_b == nullptr && "streaming flash attention does not support KQ bias"); + GGML_ASSERT(v_mla == nullptr && "streaming flash attention does not support MLA"); + GGML_ASSERT(gpu_heads > 0 && "GPU_STREAM requires an active head split"); + GGML_ASSERT(k_g && v_g && k_c && v_c); + + GGML_ASSERT(k_g->ne[3] == k_c->ne[3]); + GGML_ASSERT(v_g->ne[3] == v_c->ne[3]); + GGML_ASSERT(k_g->nb[1] <= k_g->nb[2] && "headinfer split requires non-transposed V"); + + // Split Q on the head axis at the gpu/cpu boundary (mirrors + // build_attn_mha_neo). GQA: Q-heads serving one KV head are contiguous in + // dim 1, so each half is a clean view. + const uint32_t n_head_q = (uint32_t) q->ne[1]; + const uint32_t n_head_kv = (uint32_t) (k_g->ne[1] + k_c->ne[1]); + GGML_ASSERT(n_head_q % n_head_kv == 0 && "GQA ratio must divide n_head_q"); + const uint32_t gqa_ratio = n_head_q / n_head_kv; + const uint32_t gpu_q_heads = gpu_heads * gqa_ratio; + const uint32_t cpu_q_heads = (n_head_kv - gpu_heads) * gqa_ratio; + GGML_ASSERT(gpu_q_heads + cpu_q_heads == n_head_q); + + const auto n_stream = k_g->ne[3]; + ggml_tensor * q_4 = ggml_view_4d(ctx0, q, + q->ne[0], q->ne[1], q->ne[2]/n_stream, n_stream, + q->nb[1], q->nb[2], q->nb[3]/n_stream, 0); + + ggml_tensor * q_g = ggml_view_4d(ctx0, q_4, + q_4->ne[0], gpu_q_heads, q_4->ne[2], q_4->ne[3], + q_4->nb[1], q_4->nb[2], q_4->nb[3], 0); + ggml_tensor * q_c = ggml_view_4d(ctx0, q_4, + q_4->ne[0], cpu_q_heads, q_4->ne[2], q_4->ne[3], + q_4->nb[1], q_4->nb[2], q_4->nb[3], + (size_t) gpu_q_heads * q_4->nb[1]); + + q_g = ggml_permute(ctx0, q_g, 0, 2, 1, 3); + q_c = ggml_permute(ctx0, q_c, 0, 2, 1, 3); + ggml_tensor * k_g_p = ggml_permute(ctx0, k_g, 0, 2, 1, 3); + ggml_tensor * v_g_p = ggml_permute(ctx0, v_g, 0, 2, 1, 3); + ggml_tensor * k_c_p = ggml_permute(ctx0, k_c, 0, 2, 1, 3); + ggml_tensor * v_c_p = ggml_permute(ctx0, v_c, 0, 2, 1, 3); + + if (k_g_p->type == GGML_TYPE_F32) k_g_p = ggml_cast(ctx0, k_g_p, GGML_TYPE_F16); + if (k_c_p->type == GGML_TYPE_F32) k_c_p = ggml_cast(ctx0, k_c_p, GGML_TYPE_F16); + if (v_g_p->type == GGML_TYPE_F32) v_g_p = ggml_cast(ctx0, v_g_p, GGML_TYPE_F16); + if (v_c_p->type == GGML_TYPE_F32) v_c_p = ggml_cast(ctx0, v_c_p, GGML_TYPE_F16); + + ggml_tensor * sinks_g = nullptr; + ggml_tensor * sinks_c = nullptr; + if (sinks) { + GGML_ASSERT((uint32_t) sinks->ne[0] == n_head_q); + sinks_g = ggml_view_1d(ctx0, sinks, gpu_q_heads, 0); + sinks_c = ggml_view_1d(ctx0, sinks, cpu_q_heads, + (size_t) gpu_q_heads * sinks->nb[0]); + } + + const float alibi_bias = hparams.f_max_alibi_bias; + const float logit_cap = hparams.attn_soft_cap ? hparams.f_attn_logit_softcapping : 0.0f; + + // Device head-group: ordinary resident flash-attention (no DMA). + ggml_tensor * cur_g = ggml_flash_attn_ext(ctx0, q_g, k_g_p, v_g_p, kq_mask, kq_scale, alibi_bias, logit_cap); + cb(cur_g, LLAMA_TENSOR_NAME_FATTN "_g", il); + ggml_flash_attn_ext_add_sinks(cur_g, sinks_g); + ggml_flash_attn_ext_set_prec (cur_g, GGML_PREC_F32); + + // Host (spilled) head-group: the streaming-FA op. At S1 its CUDA forward is + // plain FA, so cur_c == NEO's cur_c byte-for-byte; S2 tiles the key dim. + ggml_tensor * cur_c = ggml_streaming_flash_attn(ctx0, q_c, k_c_p, v_c_p, kq_mask, kq_scale, alibi_bias, logit_cap); + cb(cur_c, LLAMA_TENSOR_NAME_FATTN "_c", il); + ggml_flash_attn_ext_add_sinks(cur_c, sinks_c); + ggml_flash_attn_ext_set_prec (cur_c, GGML_PREC_F32); + + // Concatenate on the q-head axis (dim 1), same as build_attn_mha_neo. + ggml_tensor * cur = ggml_concat(ctx0, cur_g, cur_c, 1); + cb(cur, LLAMA_TENSOR_NAME_FATTN, il); + + cur = ggml_reshape_2d(ctx0, cur, cur->ne[0]*cur->ne[1], cur->ne[2]*cur->ne[3]); + + return cur; +} + +// opencoti F5 M7 Stage 3c-5 (#348) — POSITION_WINDOW tactic. No head split: ALL +// heads attend over the logical [0,n_kv) key range, decomposed by key POSITION into +// a device-resident window (k_win/v_win, the bulk) + a pinned-host overflow tail +// (k_tail/v_tail). The single two-region streaming op keeps the window FA device- +// resident and DMAs only the tail, merging both through the online-softmax combine — +// eliminating M2's O(n_kv) stream-back. k_tail==nullptr (n_kv<=window_cells) collapses +// to a single resident region, byte-identical to vanilla build_attn_mha's FA branch. +ggml_tensor * llm_graph_context::build_attn_mha_position_window( + ggml_tensor * q, + ggml_tensor * k_win, + ggml_tensor * v_win, + ggml_tensor * k_tail, + ggml_tensor * v_tail, + ggml_tensor * kq_b, + ggml_tensor * kq_mask, + ggml_tensor * sinks, + ggml_tensor * v_mla, + float kq_scale, + int il, + bool cpu_fa_tail) const { + GGML_ASSERT(cparams.flash_attn && "POSITION_WINDOW tactic requires flash-attention"); + GGML_ASSERT(kq_b == nullptr && "position window does not support KQ bias"); + GGML_ASSERT(v_mla == nullptr && "position window does not support MLA"); + GGML_ASSERT(sinks == nullptr && "position window does not support attention sinks"); + GGML_ASSERT(k_win && v_win); + + // Non-transposed V (the window-mode storage gate enforces !v_trans). + GGML_ASSERT(k_win->nb[1] <= k_win->nb[2] && "position window requires non-transposed V"); + + const auto n_stream = k_win->ne[3]; + + // Q: split streams then permute to FA layout (mirrors build_attn_mha exactly). + q = ggml_view_4d(ctx0, q, q->ne[0], q->ne[1], q->ne[2]/n_stream, n_stream, + q->nb[1], q->nb[2], q->nb[3]/n_stream, 0); + q = ggml_permute(ctx0, q, 0, 2, 1, 3); + + // Window (device bulk) → FA layout [d_head, n_win, n_head_kv, n_stream]. + ggml_tensor * k_win_p = ggml_permute(ctx0, k_win, 0, 2, 1, 3); + ggml_tensor * v_win_p = ggml_permute(ctx0, v_win, 0, 2, 1, 3); + // F32→F16 only (embedding non-causal path); quant K/V is dequantized on-lift + // inside the streaming op, so do NOT materialise an f16 copy here. + if (k_win_p->type == GGML_TYPE_F32) k_win_p = ggml_cast(ctx0, k_win_p, GGML_TYPE_F16); + if (v_win_p->type == GGML_TYPE_F32) v_win_p = ggml_cast(ctx0, v_win_p, GGML_TYPE_F16); + + // Tail (host overflow) → FA layout, or null when n_kv <= window_cells (resident). + ggml_tensor * k_tail_p = nullptr; + ggml_tensor * v_tail_p = nullptr; + if (k_tail) { + GGML_ASSERT(v_tail && "tail K present without tail V"); + k_tail_p = ggml_permute(ctx0, k_tail, 0, 2, 1, 3); + v_tail_p = ggml_permute(ctx0, v_tail, 0, 2, 1, 3); + if (k_tail_p->type == GGML_TYPE_F32) k_tail_p = ggml_cast(ctx0, k_tail_p, GGML_TYPE_F16); + if (v_tail_p->type == GGML_TYPE_F32) v_tail_p = ggml_cast(ctx0, v_tail_p, GGML_TYPE_F16); + } + + const float alibi_bias = hparams.f_max_alibi_bias; + const float logit_cap = hparams.attn_soft_cap ? hparams.f_attn_logit_softcapping : 0.0f; + + ggml_tensor * cur; + // opencoti F5 M7 Stage 3d-S3 (#353): CPU_FA_TAIL. The GPU FAs the device-resident + // window [0,wc); the CPU FAs the pinned-host tail [wc,n_kv) (no DMA of the tail + // K/V) and emits a tiny [2+DV, n_head] partial [M,S,VKQ]; the window streaming op + // (partial mode, op_params[4]=1, src[5]=partial) converts + merges it through the + // online-softmax combine. Only the partial crosses PCIe — the win on slow links. + // The tail mask is the [wc,n_kv) column slice of kq_mask, made contiguous (the + // tail-partial op asserts a contiguous mask). Gated upstream on decode (n_q==1) + + // a real tail; here k_tail_p is guaranteed non-null when cpu_fa_tail is set. + if (cpu_fa_tail && k_tail_p) { + // bug-360: window cells live on the accessor's ne[2] ([n_embd_head, n_head_kv, + // n_win, n_stream] — get_k_window:3037), NOT ne[1] (= n_head_kv). The old + // `k_win->ne[1]` sliced m_tail at offset n_head_kv (=2 on Gemma global) instead + // of n_win; decode-inert (uniform-zero causal mask) but wrong for prefill / any + // non-uniform mask. The window key count after permute is k_win_p->ne[1] == ne[2]. + const int64_t wc = k_win->ne[2]; // window cells (key/position axis, pre-permute) + const int64_t n_kv_m = kq_mask->ne[0]; // full logical key span + ggml_tensor * m_tail = ggml_view_2d(ctx0, kq_mask, n_kv_m - wc, kq_mask->ne[1], + kq_mask->nb[1], (size_t) wc * kq_mask->nb[0]); + m_tail = ggml_cont(ctx0, m_tail); // tail-partial op needs a contiguous mask + ggml_tensor * partial = ggml_flash_attn_ext_tail_partial( + ctx0, q, k_tail_p, v_tail_p, m_tail, kq_scale, alibi_bias, logit_cap); + cur = ggml_streaming_flash_attn_window( + ctx0, q, k_win_p, v_win_p, nullptr, nullptr, kq_mask, kq_scale, alibi_bias, logit_cap); + // op_params[0..2] = {scale,max_bias,softcap} (set by the ctor); [3..] zero-init. + // Slot 4 == 1 marks CPU_FA_TAIL partial mode for ggml_cuda_streaming_flash_attn. + cur->op_params[4] = 1; + cur->src[5] = partial; // the [2+DV, n_head] CPU partial + } else { + cur = ggml_streaming_flash_attn_window( + ctx0, q, k_win_p, v_win_p, k_tail_p, v_tail_p, kq_mask, kq_scale, alibi_bias, logit_cap); + } + cb(cur, LLAMA_TENSOR_NAME_FATTN, il); + ggml_flash_attn_ext_set_prec(cur, GGML_PREC_F32); + + cur = ggml_reshape_2d(ctx0, cur, cur->ne[0]*cur->ne[1], cur->ne[2]*cur->ne[3]); + return cur; +} + // opencoti F5 M3 neo — head-axis-split attention. Builds two flash_attn_ext // ops on disjoint head ranges: // GPU side: q[..., :gpu_heads*gqa, ...] × k_g × v_g (auto-routes to GPU @@ -2409,7 +2613,23 @@ ggml_tensor * llm_graph_context::build_attn( // null (build_attn for KV asserts this above at :2102), kq_b too; // the M2 gate ensures non-transposed V. See docs/features/advanced_kv.md. ggml_tensor * cur; - if (cparams.neo_pipeline_mode != 0 && mctx_cur->headinfer_split_active(il)) { + if (mctx_cur->get_layer_tactic(il) == llama_kv_cache::headinfer_tactic::GPU_STREAM) { + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304), S1. GPU_STREAM + // rides the M2 head-split: the device head-group runs a resident FA, + // the host (spilled) head-group runs the streaming-FA op, and the two + // are concatenated on the head axis — same structure as the NEO branch + // below, but the host half uses GGML_OP_STREAMING_FLASH_ATTN (which S2 + // tiles over the key dimension with op-owned async DMA). Fetches the + // SPLIT accessors, not get_k/get_v. Assigned only when the GPU_STREAM + // spill tactic is selected and the layer did not fit resident. + ggml_tensor * k_g = mctx_cur->get_k_gpu(ctx0, il); + ggml_tensor * v_g = mctx_cur->get_v_gpu(ctx0, il); + ggml_tensor * k_c = mctx_cur->get_k_cpu(ctx0, il); + ggml_tensor * v_c = mctx_cur->get_v_cpu(ctx0, il); + cur = build_attn_mha_streaming(q, k_g, v_g, k_c, v_c, + mctx_cur->headinfer_gpu_heads(il), + kq_b, kq_mask, sinks, v_mla, kq_scale, il); + } else if (cparams.neo_pipeline_mode != 0 && mctx_cur->headinfer_split_active(il)) { ggml_tensor * k_g = mctx_cur->get_k_gpu(ctx0, il); ggml_tensor * v_g = mctx_cur->get_v_gpu(ctx0, il); ggml_tensor * k_c = mctx_cur->get_k_cpu(ctx0, il); @@ -2597,10 +2817,40 @@ ggml_tensor * llm_graph_context::build_attn( const auto & kq_mask = is_swa ? inp->get_kq_mask_swa() : inp->get_kq_mask(); ggml_tensor * q = q_cur; - ggml_tensor * k = mctx_cur->get_k(ctx0, il); - ggml_tensor * v = mctx_cur->get_v(ctx0, il); - ggml_tensor * cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, kq_scale, il); + // opencoti-hook: f5-rolling-kv — M7 Stage 3c-5 (#348) POSITION_WINDOW dispatch on + // the iSWA path. Gemma-4's GLOBAL (!is_swa) overflowing layers carry a position + // window (device bulk [0,wc) + pinned-host tail [wc,n_kv)); route them through + // build_attn_mha_position_window, which keeps the window FA device-resident and + // streams only the tail via the two-region streaming op + online-softmax combine + // (the Stage-3a dst_lse keystone), eliminating M2's O(n_kv) stream-back. SWA + // layers and any MLA/KQ-bias/sinks path fall through to vanilla build_attn_mha. + const bool window_layer = + !is_swa && cparams.flash_attn && kq_b == nullptr && v_mla == nullptr && sinks == nullptr && + mctx_cur->get_layer_tactic(il) == llama_kv_cache::headinfer_tactic::POSITION_WINDOW && + mctx_cur->headinfer_window_active(il); + + ggml_tensor * cur; + if (window_layer) { + ggml_tensor * k_win = mctx_cur->get_k_window(ctx0, il); + ggml_tensor * v_win = mctx_cur->get_v_window(ctx0, il); + ggml_tensor * k_tail = mctx_cur->get_k_tail (ctx0, il); // nullptr when n_kv <= wc + ggml_tensor * v_tail = mctx_cur->get_v_tail (ctx0, il); + // opencoti F5 M7 Stage 3d-S3 (#353): CPU_FA_TAIL. When the S2 selector chose + // the CPU-FA tail tactic (slow PCIe link) AND there is a real overflow tail + // AND this is a single-token decode step (q->ne[2]==1 — the partial's per-row + // head-fastest order matches the combine only at n_q==1), FA the host tail on + // the CPU and merge a tiny partial instead of DMAing the tail. Otherwise (fast + // link, prefill, or no overflow) the unchanged two-region stream path runs. + const bool cpu_fa_tail = + (k_tail != nullptr) && (q->ne[2] == 1) && mctx_cur->headinfer_tail_is_cpu_fa(); + cur = build_attn_mha_position_window(q, k_win, v_win, k_tail, v_tail, + kq_b, kq_mask, sinks, v_mla, kq_scale, il, cpu_fa_tail); + } else { + ggml_tensor * k = mctx_cur->get_k(ctx0, il); + ggml_tensor * v = mctx_cur->get_v(ctx0, il); + cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, kq_scale, il); + } cb(cur, "kqv_out", il); if (v_rot) { diff --git a/llama.cpp/src/llama-graph.h b/llama.cpp/src/llama-graph.h --- a/llama.cpp/src/llama-graph.h +++ b/llama.cpp/src/llama-graph.h @@ -929,6 +929,57 @@ struct llm_graph_context { float kq_scale, int il) const; + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304), S1. GPU_STREAM tactic. + // Structurally identical to build_attn_mha_neo (device head-group resident + // FA + host head-group FA + head-axis concat), but the host (spilled) + // head-group runs through the custom GGML_OP_STREAMING_FLASH_ATTN op + // instead of a plain flash_attn_ext. At S1 that op's CUDA forward is still + // ggml_cuda_flash_attn_ext verbatim, so the output is byte-identical to the + // NEO path (minus NEO's stream-pair registration, which only affects + // scheduling) — the S1 gate. S2 makes the op tile the host head-group's KEY + // dimension into a device slot pool with a cross-tile online-softmax, so + // the spilled KV need not fit one resident slot; S3 adds the copy/compute + // double-buffer. The host head-group is already pinned-host (the M2 split + // allocates k_cpu/v_cpu with the pinned cpu_buft), so streaming rides the + // M2 residency — no whole-layer host reorder. Flash-attention only (asserts + // cparams.flash_attn, kq_b == nullptr, v_mla == nullptr, gpu_heads > 0). + ggml_tensor * build_attn_mha_streaming( + ggml_tensor * q, // [n_embd_head_q, n_head_q, n_tokens] + ggml_tensor * k_g, // GPU-resident K subset (gpu_heads on dim 1) + ggml_tensor * v_g, // GPU-resident V subset (gpu_heads on dim 1) + ggml_tensor * k_c, // host-resident K subset (cpu_heads on dim 1) — streamed + ggml_tensor * v_c, // host-resident V subset (cpu_heads on dim 1) — streamed + uint32_t gpu_heads, + ggml_tensor * kq_b, // must be nullptr + ggml_tensor * kq_mask, + ggml_tensor * sinks, // [n_head_q] or nullptr + ggml_tensor * v_mla, // must be nullptr + float kq_scale, + int il) const; + + // opencoti F5 M7 Stage 3c-5 — POSITION_WINDOW tactic. The KV is split by KEY + // POSITION (not head): a device-resident window [0,wc) holds the bulk, a + // pinned-host tail [wc,n_kv) holds the overflow. ALL heads attend over the + // logical [0,n_kv) range via the two-region streaming op + // (ggml_streaming_flash_attn_window): the window FA stays device-resident and + // only the small tail is DMA'd, then both merge through the online-softmax + // combine — NO M2-style O(n_kv) stream-back. k_tail==nullptr (n_kv<=wc) collapses + // to a single resident FA, byte-identical to vanilla. Flash-attention only + // (asserts cparams.flash_attn, kq_b == nullptr, v_mla == nullptr, sinks == nullptr). + ggml_tensor * build_attn_mha_position_window( + ggml_tensor * q, // [n_embd_head_q, n_head_q, n_tokens] + ggml_tensor * k_win, // device window K (all heads, [0,wc) on dim 2) + ggml_tensor * v_win, // device window V + ggml_tensor * k_tail, // host tail K (all heads, [wc,n_kv) on dim 2) or nullptr + ggml_tensor * v_tail, // host tail V or nullptr + ggml_tensor * kq_b, // must be nullptr + ggml_tensor * kq_mask, // full [n_kv] logical mask + ggml_tensor * sinks, // must be nullptr + ggml_tensor * v_mla, // must be nullptr + float kq_scale, + int il, + bool cpu_fa_tail = false) const; // S3-b2 (#353): FA the host tail on the CPU (no DMA) + llm_graph_input_attn_no_cache * build_attn_inp_no_cache() const; ggml_tensor * build_attn( diff --git a/llama.cpp/src/llama-kv-cache-iswa.cpp b/llama.cpp/src/llama-kv-cache-iswa.cpp --- a/llama.cpp/src/llama-kv-cache-iswa.cpp +++ b/llama.cpp/src/llama-kv-cache-iswa.cpp @@ -23,6 +23,9 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( uint32_t kv_size_initial, uint32_t slot_shrink_idle_ms, float headinfer_gpu_heads_frac, + uint32_t vram_target_mib, + float pcie_bw_gbps, + uint32_t kv_residency_mode, uint32_t n_seq_max, uint32_t n_ubatch, uint32_t n_pad, @@ -65,7 +68,7 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( kv_base = std::make_unique( model, type_k, type_v, v_trans, offload, unified, size_base, kv_size_initial, slot_shrink_idle_ms, - headinfer_gpu_heads_frac, + headinfer_gpu_heads_frac, vram_target_mib, pcie_bw_gbps, kv_residency_mode, n_seq_max, n_pad, 0, LLAMA_SWA_TYPE_NONE, filter_base, reuse); LLAMA_LOG_INFO("%s: creating SWA KV cache, size = %u cells\n", __func__, size_swa); @@ -73,7 +76,7 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( kv_swa = std::make_unique( model, type_k, type_v, v_trans, offload, unified, size_swa, kv_size_initial, slot_shrink_idle_ms, - headinfer_gpu_heads_frac, + headinfer_gpu_heads_frac, vram_target_mib, pcie_bw_gbps, kv_residency_mode, n_seq_max, n_pad, hparams.n_swa, hparams.swa_type, filter_swa, reuse); } diff --git a/llama.cpp/src/llama-kv-cache-iswa.h b/llama.cpp/src/llama-kv-cache-iswa.h --- a/llama.cpp/src/llama-kv-cache-iswa.h +++ b/llama.cpp/src/llama-kv-cache-iswa.h @@ -30,6 +30,17 @@ public: uint32_t slot_shrink_idle_ms, // opencoti F5 M2 headinfer — forwarded to both kv_base and kv_swa. float headinfer_gpu_heads_frac, + // opencoti F5 M7 Rolling KV — Rung 0 — forwarded to both kv_base + // and kv_swa. VRAM budget cap (MiB); 0 = maximize. Consulted only + // when headinfer_gpu_heads_frac is negative (auto). + uint32_t vram_target_mib, + // opencoti F5 M7 Rolling KV — Rung 2 (Stage 3d-S2 #352). Effective + // host->device PCIe bandwidth (GB/s); forwarded to both kv_base and + // kv_swa (only the overflowing base cache uses it). 0 = unknown. + float pcie_bw_gbps, + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob; + // forwarded to both kv_base and kv_swa. 0 = auto, 1 = head, 2 = window. + uint32_t kv_residency_mode, uint32_t n_seq_max, uint32_t n_ubatch, uint32_t n_pad, diff --git a/llama.cpp/src/llama-kv-cache.cpp b/llama.cpp/src/llama-kv-cache.cpp --- a/llama.cpp/src/llama-kv-cache.cpp +++ b/llama.cpp/src/llama-kv-cache.cpp @@ -100,6 +100,218 @@ static ggml_tensor * ggml_mul_mat_aux( // llama_kv_cache // +// opencoti F5 M7 Rolling KV — Rung 0 auto-residency — see docs/features/rolling_kv.md +// Sentinel for headinfer_gpu_heads_frac meaning "decide automatically": when the +// incoming fraction is < 0, the KV-cache constructor computes an effective frac +// from the offload device's free VRAM vs this cache's KV demand. eff_frac == 1.0 +// (KV fits the budget) allocates NO split tensors — byte-identical to vanilla +// (GPU_RESIDENT). Only a budget shortfall yields eff_frac < 1.0, which engages +// the shipped M2 CPU-spill split, sized to fit. The non-negative path is the +// manual escape hatch / upstream default and is used verbatim. +static constexpr float OPENCOTI_HEADINFER_FRAC_AUTO = -1.0f; + +// Conservative VRAM held back from the KV budget for compute scratch + cuBLAS +// workspace, which are allocated AFTER the KV cache (graph_reserve / context +// init). Model weights are already resident at KV-cache-construction time, so +// the free-VRAM read here already nets them out; this reserve only covers the +// not-yet-allocated compute buffers. --vram-target gives the operator a tighter +// knob when this default is wrong for a given workload. +static constexpr size_t OPENCOTI_HEADINFER_AUTO_RESERVE_BYTES = (size_t) 1536 * 1024 * 1024; + +// opencoti F5 M7 Rolling KV — Rung 0. Compute the effective GPU-head fraction +// for HeadInfer (M2) from a VRAM-fit check. Returns 1.0 (GPU_RESIDENT — vanilla +// fast path, no CPU split) when this cache's full KV demand fits the offload +// device's budget; otherwise the largest fraction that fits, so the M2 split is +// sized to exactly relieve the shortfall. Respects the same has_kv()/filter() +// gating as the allocation loop so base/SWA caches each size only their layers. +// Single-offload-device assumption: the budget is read from the first KV layer's +// device; if KV layers span multiple devices the result is approximate (a +// warning is logged) — later M7 rungs add a per-layer dynamic scheduler. +static float opencoti_compute_auto_gpu_heads_frac( + const llama_model & model, + const llama_hparams & hparams, + ggml_type type_k, + ggml_type type_v, + bool v_trans, + bool offload, + uint32_t kv_size, + uint32_t n_stream, + uint32_t vram_target_mib, + const std::function & filter) { + // No GPU offload → nothing is "GPU-resident" and nothing to spill; the M2 + // split is only meaningful for offloaded layers. Stay at vanilla. + if (!offload) { + return 1.0f; + } + + size_t kv_demand = 0; + ggml_backend_dev_t dev = nullptr; + ggml_backend_dev_t dev_prev = nullptr; + int n_dev_distinct = 0; + for (uint32_t il = 0; il < hparams.n_layer; ++il) { + if (!hparams.has_kv(il)) { + continue; + } + if (filter && !filter(il)) { + continue; + } + const uint32_t n_embd_k_gqa = hparams.n_embd_k_gqa(il); + const uint32_t n_embd_v_gqa = !v_trans ? hparams.n_embd_v_gqa(il) : hparams.n_embd_v_gqa_max(); + kv_demand += (size_t) ggml_row_size(type_k, n_embd_k_gqa) * kv_size; + kv_demand += (size_t) ggml_row_size(type_v, n_embd_v_gqa) * kv_size; + + ggml_backend_dev_t d = model.dev_layer(il); + if (!dev) { + dev = d; + } + if (d != dev_prev) { + dev_prev = d; + ++n_dev_distinct; + } + } + kv_demand *= (size_t) n_stream; + + if (kv_demand == 0 || !dev) { + return 1.0f; + } + if (n_dev_distinct > 1) { + LLAMA_LOG_WARN("%s: KV layers span %d devices; auto headinfer fraction is " + "approximate (budget read from the first device)\n", __func__, n_dev_distinct); + } + + size_t free_vram = 0, total_vram = 0; + ggml_backend_dev_memory(dev, &free_vram, &total_vram); + + size_t budget = free_vram; + if (vram_target_mib > 0) { + budget = std::min(budget, (size_t) vram_target_mib * 1024 * 1024); + } + budget = budget > OPENCOTI_HEADINFER_AUTO_RESERVE_BYTES + ? budget - OPENCOTI_HEADINFER_AUTO_RESERVE_BYTES + : 0; + + if (kv_demand <= budget) { + LLAMA_LOG_INFO("%s: auto residency = GPU_RESIDENT (KV %.0f MiB fits budget " + "%.0f MiB of %.0f MiB free); headinfer split OFF\n", __func__, + kv_demand / 1048576.0, budget / 1048576.0, free_vram / 1048576.0); + return 1.0f; + } + + float frac = (float) ((double) budget / (double) kv_demand); + frac = std::min(frac, 0.99f); + frac = std::max(frac, 0.0f); + LLAMA_LOG_INFO("%s: auto residency = CPU_SPILL (KV %.0f MiB exceeds budget " + "%.0f MiB of %.0f MiB free); headinfer frac = %.3f\n", __func__, + kv_demand / 1048576.0, budget / 1048576.0, free_vram / 1048576.0, (double) frac); + return frac; +} + +// opencoti F5 M7 Rolling KV — Stage 3c. Position-axis residency sizing: how many KV +// *cells* (positions) [0, C) fit resident in VRAM as the device "window"; the overflow +// [C, kv_size) becomes the pinned-host "tail" streamed device-ward per tile. Mirrors +// opencoti_compute_auto_gpu_heads_frac's budget math but on the POSITION axis (head- +// agnostic), so it sidesteps Gemma-4's 2-KV-head binary spill (the head axis can only +// express 0/50/100%). Returns kv_size — meaning "no tail", fully resident, byte-identical +// to vanilla — when the full cache fits the budget; otherwise the largest cell count C +// that fits, rounded DOWN to n_pad (covers q8_0 blck alignment, bug-225 class) and clamped +// to [n_pad, kv_size - n_pad]. Single-offload-device assumption mirrors the auto-frac sizer. +// +// BELT-AND-SUSPENDERS (bug-263, user-requested "GO + belt-and-suspenders"): the decode +// window FA processes exactly C keys, and C in the band (480, 512] maps to ntiles_KV==16 → +// gridDim.x == WARP_SIZE for head_dim > 256 — the one uniform tiling that maximizes online- +// softmax combine-order divergence (numerically valid but greedy-fragile; bug-263 is that +// sensitivity, not a defect). We snap C below that band so the resident window never sits on +// it. Universally harmless (≤ 32-cell shrink in a rare tiny-budget edge case); the symmetric +// guard for tail *tiles* lives in the streaming forward. See .wolf/buglog.json bug-263. +static uint32_t opencoti_compute_resident_window_cells( + const llama_model & model, + const llama_hparams & hparams, + ggml_type type_k, + ggml_type type_v, + bool v_trans, + bool offload, + uint32_t kv_size, + uint32_t n_stream, + uint32_t vram_target_mib, + uint32_t n_pad, + const std::function & filter) { + if (!offload || kv_size == 0) { + return kv_size; // no position-axis spill without GPU offload + } + + size_t per_cell = 0; // bytes for ONE position cell across all this cache's KV layers + ggml_backend_dev_t dev = nullptr; + for (uint32_t il = 0; il < hparams.n_layer; ++il) { + if (!hparams.has_kv(il)) { + continue; + } + if (filter && !filter(il)) { + continue; + } + const uint32_t n_embd_k_gqa = hparams.n_embd_k_gqa(il); + const uint32_t n_embd_v_gqa = !v_trans ? hparams.n_embd_v_gqa(il) : hparams.n_embd_v_gqa_max(); + per_cell += (size_t) ggml_row_size(type_k, n_embd_k_gqa); + per_cell += (size_t) ggml_row_size(type_v, n_embd_v_gqa); + if (!dev) { + dev = model.dev_layer(il); + } + } + per_cell *= (size_t) n_stream; + if (per_cell == 0 || !dev) { + return kv_size; + } + + size_t free_vram = 0, total_vram = 0; + ggml_backend_dev_memory(dev, &free_vram, &total_vram); + size_t budget = free_vram; + if (vram_target_mib > 0) { + budget = std::min(budget, (size_t) vram_target_mib * 1024 * 1024); + } + budget = budget > OPENCOTI_HEADINFER_AUTO_RESERVE_BYTES + ? budget - OPENCOTI_HEADINFER_AUTO_RESERVE_BYTES + : 0; + + if (per_cell * (size_t) kv_size <= budget) { + LLAMA_LOG_INFO("%s: position window = FULLY RESIDENT (KV %.0f MiB fits budget " + "%.0f MiB of %.0f MiB free); no host tail\n", __func__, + per_cell * (size_t) kv_size / 1048576.0, budget / 1048576.0, free_vram / 1048576.0); + return kv_size; + } + + // 3c-5 / bug-267: the window split boundary wc feeds per-tile flash-attention + // DIRECTLY — window keys [0,wc) and tail keys [wc,n_kv) each become their own + // per-tile FA sub-calls in ggml_cuda_streaming_flash_attn. For head_dim 512/576 + // (Gemma global) the CUDA FA kernel HARD-REQUIRES gqa_opt, i.e. + // K->ne[1] % FATTN_KQ_STRIDE(256) == 0 (fattn.cu get_best_fattn_kernel); a wc that + // is merely n_pad(32)-aligned makes the window's (and tail's) partial last tile + // resolve to BEST_FATTN_KERNEL_NONE → GGML_ABORT("fatal error") on the warmup run. + // n_kv is already padded to a 256-multiple by the FA graph, so aligning wc to 256 + // keeps BOTH regions 256-aligned (n_tail = n_kv - wc inherits it). 256 is a + // multiple of the q8_0 block (32, bug-225) and of every n_pad we use, so this + // strictly TIGHTENS the prior alignment. Align to lcm(n_pad, 256) (= max for the + // power-of-2 n_pad llama.cpp uses). The bug-263 (480,512] band guard is dropped: + // bug-263 is RESOLVED as FP combine-order sensitivity (acceptance gate is + // cosine+semantic, never greedy byte-equality), and 480 is not 256-aligned anyway. + const uint32_t pad = n_pad ? n_pad : 32; + const uint32_t fa_align = 256; // FATTN_KQ_STRIDE + const uint32_t walign = pad >= fa_align ? pad : fa_align; // lcm(pad,256), power-of-2 pad + uint32_t cells = (uint32_t) (budget / per_cell); + cells = (cells / walign) * walign; // round DOWN to the FA-aligned window + if (cells < walign) { + cells = walign; // never below one FA stride (~0.5 MiB — never OOMs) + } + if (kv_size > walign && cells > kv_size - walign) { + cells = ((kv_size - walign) / walign) * walign; // keep >= one FA stride of tail, stay aligned + if (cells < walign) { + cells = walign; + } + } + LLAMA_LOG_INFO("%s: position window = %u / %u cells resident (%.0f MiB budget, " + "%.0f MiB/cell-set); host tail = %u cells\n", __func__, + cells, kv_size, budget / 1048576.0, per_cell / 1048576.0, kv_size - cells); + return cells; +} + llama_kv_cache::llama_kv_cache( const llama_model & model, ggml_type type_k, @@ -111,6 +323,9 @@ llama_kv_cache::llama_kv_cache( uint32_t kv_size_initial, uint32_t slot_shrink_idle_ms, float headinfer_gpu_heads_frac, + uint32_t vram_target_mib, + float pcie_bw_gbps, + uint32_t kv_residency_mode, uint32_t n_seq_max, uint32_t n_pad, uint32_t n_swa, @@ -154,6 +369,60 @@ llama_kv_cache::llama_kv_cache( const uint32_t n_layer_kv = hparams.n_layer_kv(); + // opencoti F5 M7 Rolling KV — Rung 0 auto-residency — see docs/features/rolling_kv.md + // Resolve the effective HeadInfer fraction once for this cache. A negative + // incoming frac (OPENCOTI_HEADINFER_FRAC_AUTO) means "decide from free VRAM": + // KV fits → 1.0 (GPU_RESIDENT, byte-identical to vanilla — no split tensors); + // shortfall → the largest fit-sized fraction (engages the M2 CPU-spill split). + // A non-negative frac is the manual escape hatch / upstream default, used + // verbatim. Every split decision below reads eff_gpu_heads_frac, not the raw + // ctor argument, so the sentinel never leaks into the allocation math. + const float eff_gpu_heads_frac = headinfer_gpu_heads_frac < 0.0f + ? opencoti_compute_auto_gpu_heads_frac(model, hparams, type_k, type_v, + v_trans, offload, kv_size, n_stream, vram_target_mib, filter) + : headinfer_gpu_heads_frac; + + // opencoti F5 M7 Stage 3c — position-axis window sizing. Computed + logged now so the + // sizing math can be validated against real models (boot at various --vram-target and + // read the logged C) BEFORE the storage/forward wiring lands in the next 3c sub-stages. + // window_cells == kv_size ⇒ no host tail (fully resident, byte-identical vanilla). The + // value is intentionally unused downstream at this sub-stage. See docs/features/advanced_kv.md. + const uint32_t window_cells = opencoti_compute_resident_window_cells( + model, hparams, type_k, type_v, v_trans, offload, kv_size, n_stream, + vram_target_mib, n_pad, filter); + // opencoti F5 M7 Stage 4 (#338) — position-window mode is driven by the + // --kv-residency-mode knob (cparams.kv_residency_mode): 0=auto, 1=head, 2=window. + // The ELIGIBILITY predicate is the validated POSITION_WINDOW class: the sizer found a + // real overflow (0 < C < kv_size) on an offloaded, non-transposed-V, APPEND-ONLY + // (swa_type==NONE) cache. window mode is MUTUALLY EXCLUSIVE with the M2 head split + // (forced off below), so the shared k_per_stream/k_cpu_per_stream slots carry the + // position split, never both. + // + // bug-268: window mode is valid ONLY on an APPEND-ONLY cache whose cell index equals + // the key position (the keystone — see plan). That holds for the iSWA base / non-SWA + // cache (swa_type == NONE, n_swa == 0, find_slot append-only) but NOT for the + // sliding-window cache, whose cells rotate/reuse (idx != position) — a tiny + // --vram-target made the sizer return C 0 && window_cells < kv_size + && swa_type == LLAMA_SWA_TYPE_NONE; + // auto (0): window when eligible. head (1): never window (force the M2 split). + // window (2): force window when eligible; warn + fall back when ineligible. + const bool window_mode = window_eligible && kv_residency_mode != 1; + if (kv_residency_mode == 2 && !window_eligible) { + LLAMA_LOG_WARN("%s: --kv-residency-mode window requested but this cache is ineligible " + "(offload=%d v_trans=%d window_cells=%u/%u swa_type=%d) — falling back to " + "head/resident\n", __func__, offload, v_trans, window_cells, kv_size, (int) swa_type); + } + if (window_mode) { + LLAMA_LOG_INFO("%s: rolling-kv POSITION_WINDOW mode ON (--kv-residency-mode %s) — " + "window %u / %u cells, host tail %u (M2 head split forced off)\n", __func__, + kv_residency_mode == 2 ? "window" : "auto", + window_cells, kv_size, kv_size - window_cells); + } + // define a comparator for the buft -> ctx map to ensure that the order is well-defined: struct ggml_backend_buft_comparator { bool operator()(const ggml_backend_buffer_type_t & lhs, const ggml_backend_buffer_type_t & rhs) const { @@ -273,6 +542,7 @@ llama_kv_cache::llama_kv_cache( // pageable CPU buft when no CUDA dev / no dev_host_buffer_type. Pinned // host buffers DMA at full PCIe bandwidth; pageable hits ~50%. ggml_backend_buffer_type_t cpu_buft = ggml_backend_cpu_buffer_type(); + if (offload) { auto * dev = model.dev_layer(il); buft = ggml_backend_dev_buffer_type(dev); @@ -311,12 +581,15 @@ llama_kv_cache::llama_kv_cache( uint32_t n_embd_v_gpu = n_embd_v_gqa; uint32_t n_embd_k_cpu = 0; uint32_t n_embd_v_cpu = 0; - if (headinfer_gpu_heads_frac < 1.0f && offload && !v_trans) { + // opencoti F5 M7 Stage 3c — POSITION_WINDOW and the M2 head split are + // mutually exclusive: when window mode is on, force the head split off + // so k_per_stream carries the position window (ALL heads), never both. + if (!window_mode && eff_gpu_heads_frac < 1.0f && offload && !v_trans) { const uint32_t n_head_kv = hparams.n_head_kv(il); const uint32_t n_embd_head_k = hparams.n_embd_head_k(il); const uint32_t n_embd_head_v = hparams.n_embd_head_v(il); if (n_head_kv > 1) { - uint32_t g = (uint32_t) std::lround((double) headinfer_gpu_heads_frac * (double) n_head_kv); + uint32_t g = (uint32_t) std::lround((double) eff_gpu_heads_frac * (double) n_head_kv); g = std::max(1u, std::min(g, n_head_kv - 1)); gpu_heads = g; n_embd_k_gpu = g * n_embd_head_k; @@ -328,6 +601,14 @@ llama_kv_cache::llama_kv_cache( } } + // opencoti F5 M7 Stage 3c — position-window cell counts. In window + // mode the device tensor spans only the resident window [0, wc) and the + // host tail tensor spans [wc, kv_size) (tail_c cells); both carry ALL + // heads. Outside window mode wc == kv_size and tail_c == 0, so the + // tensor shapes below are byte-identical to the pre-3c layout. + const uint32_t wc = window_mode ? window_cells : kv_size; + const uint32_t tail_c = window_mode ? (kv_size - window_cells) : 0; + // opencoti F4 M3 Phase 4 — per-stream tensor split. // Each stream gets its own OWNING 3-D tensor of shape // [n_embd, kv_size, 1]. Dim 3 == 1 keeps stride math @@ -349,11 +630,13 @@ llama_kv_cache::llama_kv_cache( throw std::runtime_error("failed to create ggml context for kv cache"); } + // opencoti F5 M7 Stage 3c — `wc` is kv_size outside window mode + // (byte-identical), or the resident window length in window mode. ggml_tensor * k_s = has_k - ? ggml_new_tensor_3d(ctx_s, type_k, n_embd_k_gpu, kv_size, 1) + ? ggml_new_tensor_3d(ctx_s, type_k, n_embd_k_gpu, wc, 1) : nullptr; ggml_tensor * v_s = has_v - ? ggml_new_tensor_3d(ctx_s, type_v, n_embd_v_gpu, kv_size, 1) + ? ggml_new_tensor_3d(ctx_s, type_v, n_embd_v_gpu, wc, 1) : nullptr; // Tensor naming preserves the pre-Phase-4 name in unified @@ -382,34 +665,44 @@ llama_kv_cache::llama_kv_cache( // in its own host buffer (Phase 4-B). Vectors stay empty // when gpu_heads == 0 so headinfer_split_active() can test // by vector size. - if (gpu_heads > 0) { - // opencoti F5 M7-A (0034) — was ggml_backend_cpu_buffer_type() - // (pageable). Now uses the layer-scope cpu_buft resolved above: - // pinned host buffer when CUDA is the layer dev, plain CPU buft - // otherwise. Pinned unlocks full PCIe DMA for M2's CPU-half - // stream-back AND is the foundation Rolling KV (M7) builds on. + // opencoti F5 M7-A (0034) — host-resident KV subset. M2 head split + // (gpu_heads > 0) stores the CPU head-half here; F5 M7 Stage 3c + // window mode (window_mode) instead stores the ALL-heads position + // tail [wc, kv_size) here. The two are mutually exclusive (window + // mode forces gpu_heads == 0), so exactly one branch sizes the + // tensor. The host buffer is the pinned cpu_buft when CUDA is the + // layer dev — full PCIe DMA for the stream-back. + if (gpu_heads > 0 || window_mode) { ggml_context * ctx_cpu_s = ctx_for_buft(cpu_buft, s); if (!ctx_cpu_s) { throw std::runtime_error("failed to create CPU ggml context for headinfer kv cache"); } + // window mode: ALL heads, tail_c cells; M2: CPU head subset, kv_size cells. + const uint32_t host_k_embd = window_mode ? n_embd_k_gqa : n_embd_k_cpu; + const uint32_t host_v_embd = window_mode ? n_embd_v_gqa : n_embd_v_cpu; + const uint32_t host_cells = window_mode ? tail_c : kv_size; ggml_tensor * k_cpu_s = has_k - ? ggml_new_tensor_3d(ctx_cpu_s, type_k, n_embd_k_cpu, kv_size, 1) + ? ggml_new_tensor_3d(ctx_cpu_s, type_k, host_k_embd, host_cells, 1) : nullptr; ggml_tensor * v_cpu_s = has_v - ? ggml_new_tensor_3d(ctx_cpu_s, type_v, n_embd_v_cpu, kv_size, 1) + ? ggml_new_tensor_3d(ctx_cpu_s, type_v, host_v_embd, host_cells, 1) : nullptr; + // Name distinguishes the two layouts for state-IO grep: + // "cache_k_tail" (position tail) vs "cache_k_cpu" (head half). + const char * k_base = window_mode ? "cache_k_tail" : "cache_k_cpu"; + const char * v_base = window_mode ? "cache_v_tail" : "cache_v_cpu"; if (k_cpu_s) { if (n_stream == 1) { - ggml_format_name(k_cpu_s, "cache_k_cpu_l%d", il); + ggml_format_name(k_cpu_s, "%s_l%d", k_base, il); } else { - ggml_format_name(k_cpu_s, "cache_k_cpu_l%d_s%u", il, s); + ggml_format_name(k_cpu_s, "%s_l%d_s%u", k_base, il, s); } } if (v_cpu_s) { if (n_stream == 1) { - ggml_format_name(v_cpu_s, "cache_v_cpu_l%d", il); + ggml_format_name(v_cpu_s, "%s_l%d", v_base, il); } else { - ggml_format_name(v_cpu_s, "cache_v_cpu_l%d_s%u", il, s); + ggml_format_name(v_cpu_s, "%s_l%d_s%u", v_base, il, s); } } k_cpu_per_stream.push_back(k_cpu_s); @@ -429,11 +722,13 @@ llama_kv_cache::llama_kv_cache( // are byte-identical to pre-Phase-4 because the view has // identical ne/nb and (in unified mode) the same underlying // buffer layout. + // opencoti F5 M7 Stage 3c — the view spans the device tensor's + // actual ne[1] (== wc), which is kv_size outside window mode. k_stream.push_back(k_s - ? ggml_view_2d(ctx_s, k_s, n_embd_k_gpu, kv_size, k_s->nb[1], 0) + ? ggml_view_2d(ctx_s, k_s, n_embd_k_gpu, wc, k_s->nb[1], 0) : nullptr); v_stream.push_back(v_s - ? ggml_view_2d(ctx_s, v_s, n_embd_v_gpu, kv_size, v_s->nb[1], 0) + ? ggml_view_2d(ctx_s, v_s, n_embd_v_gpu, wc, v_s->nb[1], 0) : nullptr); } @@ -445,6 +740,7 @@ llama_kv_cache::llama_kv_cache( k_stream, v_stream, k_cpu_per_stream, v_cpu_per_stream, gpu_heads, + window_mode ? window_cells : 0u, }); } @@ -472,6 +768,12 @@ llama_kv_cache::llama_kv_cache( } } + // opencoti F5 M7 Rolling KV — Rung 1. Derive the per-layer tactic table + + // streaming plan from the split decisions just baked into `layers`, and log + // the histogram. eff_bw (Stage 3d-S2 #352) feeds the overflow tail-tactic + // crossover selector. Compute-inert at Rung 1 (nothing reads the table yet). + build_rolling_kv_plan(pcie_bw_gbps); + // opencoti F4 M3 Phase 1 — see docs/decisions/0001-lazy-slot-context.md // Eager malloc, deferred zero-fill. The address-space reservation done // by ggml_backend_alloc_ctx_tensors_from_buft is essentially free on @@ -616,35 +918,41 @@ void llama_kv_cache::ensure_cleared(uint32_t up_to_cells) { // tensor. M2's CPU heads follow the same pattern via // k_cpu_per_stream / v_cpu_per_stream (empty when no split). for (const auto & layer : layers) { + // opencoti F5 M7 Stage 3c — in window mode the device tensor holds only + // the resident window cells [0, wc) and the host tensor holds the tail + // cells [wc, kv_size) at local index (cell - wc). Clamp every memset to + // the destination tensor's own cell capacity (ne[1]) so a partial-window + // tensor is never written past its end (the warmup zero-fill bug). wc == 0 + // ⇒ M2 / vanilla layout (host tensor shares the device cell range), and + // the clamp is a no-op (b <= kv_size == ne[1]) → byte-identical to before. + const uint32_t wc = layer.window_cells; + // Clear logical cells [a, b) of a tensor whose local cell 0 maps to + // logical cell `base`, clamped to the tensor's capacity (ne[1]). + auto clear_range = [&](ggml_tensor * t, uint32_t base) { + if (!t) { + return; + } + const uint32_t cap = (uint32_t) t->ne[1]; + const uint32_t lo0 = a > base ? a - base : 0; + const uint32_t hi0 = b > base ? b - base : 0; + const uint32_t lo = lo0 < cap ? lo0 : cap; + const uint32_t hi = hi0 < cap ? hi0 : cap; + if (hi > lo) { + ggml_backend_tensor_memset(t, 0, (size_t) lo * t->nb[1], (size_t) (hi - lo) * t->nb[1]); + } + }; for (uint32_t st = 0; st < n_stream; ++st) { - auto * k = layer.k_per_stream[st]; - if (k) { - const size_t off = (size_t) a * (size_t) k->nb[1]; - const size_t sz = (size_t) (b - a) * (size_t) k->nb[1]; - ggml_backend_tensor_memset(k, 0, off, sz); - } - auto * v = layer.v_per_stream[st]; - if (v) { - const size_t off = (size_t) a * (size_t) v->nb[1]; - const size_t sz = (size_t) (b - a) * (size_t) v->nb[1]; - ggml_backend_tensor_memset(v, 0, off, sz); - } - // M2 headinfer host-resident head subset (same [a, b) cells). + // Device window (or full) tensor: local cell 0 == logical cell 0. + clear_range(layer.k_per_stream[st], 0); + clear_range(layer.v_per_stream[st], 0); + // Host tensor: window tail's local cell 0 == logical cell wc; the M2 + // head subset (wc == 0) shares the device cell range (base 0). + const uint32_t host_base = wc; if (st < layer.k_cpu_per_stream.size()) { - auto * k_cpu = layer.k_cpu_per_stream[st]; - if (k_cpu) { - const size_t off = (size_t) a * (size_t) k_cpu->nb[1]; - const size_t sz = (size_t) (b - a) * (size_t) k_cpu->nb[1]; - ggml_backend_tensor_memset(k_cpu, 0, off, sz); - } + clear_range(layer.k_cpu_per_stream[st], host_base); } if (st < layer.v_cpu_per_stream.size()) { - auto * v_cpu = layer.v_cpu_per_stream[st]; - if (v_cpu) { - const size_t off = (size_t) a * (size_t) v_cpu->nb[1]; - const size_t sz = (size_t) (b - a) * (size_t) v_cpu->nb[1]; - ggml_backend_tensor_memset(v_cpu, 0, off, sz); - } + clear_range(layer.v_cpu_per_stream[st], host_base); } } } @@ -1910,6 +2218,51 @@ ggml_tensor * llama_kv_cache::get_k(ggml_context * ctx, int32_t il, uint32_t n_k return result; } + // opencoti F5 M7 Stage 3c — position-window reassembly. The device tensor + // holds the resident window [0, wc); the pinned host tail holds [wc, kv_size). + // Build a window view over the device tensor (dim-2 length n_win) + a tail + // view over the host tensor (dim-2 length n_tail) and concat on the POSITION + // axis (dim 2). The backend scheduler streams the host tail to the compute + // backend — same mechanism as the M2 head concat above. When n_kv <= wc the + // tail is empty → a single window view, byte-identical to vanilla resident. + // This is the correctness fallback (3c-4); the perf path (separate window/tail + // FA + online-softmax combine, no streamback) is build_attn_mha_position_window. + if (layer.window_cells > 0 && !layer.k_cpu_per_stream.empty() && layer.k_cpu_per_stream[0]) { + const uint32_t n_embd_head_k = hparams.n_embd_head_k(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_k_gqa2 = (uint64_t) n_head_kv * n_embd_head_k; + const uint32_t wc = layer.window_cells; + const uint32_t n_win = n_kv < wc ? n_kv : wc; + const uint32_t n_tail = n_kv > wc ? n_kv - wc : 0; + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * k_s = layer.k_per_stream[s_cache]; + ggml_tensor * kc_s = layer.k_cpu_per_stream[s_cache]; + ggml_tensor * win_v = ggml_view_4d(ctx, k_s, + n_embd_head_k, n_head_kv, n_win, 1, + ggml_row_size(k_s->type, n_embd_head_k), + ggml_row_size(k_s->type, n_embd_k_gqa2), + ggml_row_size(k_s->type, n_embd_k_gqa2 * (uint64_t) k_s->ne[1]), + 0); + if (n_tail == 0) { + return win_v; + } + ggml_tensor * tail_v = ggml_view_4d(ctx, kc_s, + n_embd_head_k, n_head_kv, n_tail, 1, + ggml_row_size(kc_s->type, n_embd_head_k), + ggml_row_size(kc_s->type, n_embd_k_gqa2), + ggml_row_size(kc_s->type, n_embd_k_gqa2 * (uint64_t) kc_s->ne[1]), + 0); + return ggml_concat(ctx, win_v, tail_v, 2); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; + } + // opencoti F4 M3 Phase 4-C: per-stream view + concat along stream axis // (dim 3). In unified mode (sinfo.s1 == sinfo.s0 == 0) the loop does // not fire, yielding a single 4-D view of layer.k_per_stream[0] — @@ -1983,6 +2336,46 @@ ggml_tensor * llama_kv_cache::get_v(ggml_context * ctx, int32_t il, uint32_t n_k return result; } + // opencoti F5 M7 Stage 3c — position-window reassembly (mirror of get_k). + // window mode is gated to !v_trans at construction, so V is always the + // non-transposed (heads in dim 1) layout here. Concat device window ⊕ host + // tail on the position axis (dim 2); empty tail (n_kv <= wc) → window only. + if (layer.window_cells > 0 && !layer.v_cpu_per_stream.empty() && layer.v_cpu_per_stream[0]) { + const uint32_t n_embd_head_v = hparams.n_embd_head_v(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_v_gqa2 = (uint64_t) n_head_kv * n_embd_head_v; + const uint32_t wc = layer.window_cells; + const uint32_t n_win = n_kv < wc ? n_kv : wc; + const uint32_t n_tail = n_kv > wc ? n_kv - wc : 0; + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * v_s = layer.v_per_stream[s_cache]; + ggml_tensor * vc_s = layer.v_cpu_per_stream[s_cache]; + ggml_tensor * win_v = ggml_view_4d(ctx, v_s, + n_embd_head_v, n_head_kv, n_win, 1, + ggml_row_size(v_s->type, n_embd_head_v), + ggml_row_size(v_s->type, n_embd_v_gqa2), + ggml_row_size(v_s->type, n_embd_v_gqa2 * (uint64_t) v_s->ne[1]), + 0); + if (n_tail == 0) { + return win_v; + } + ggml_tensor * tail_v = ggml_view_4d(ctx, vc_s, + n_embd_head_v, n_head_kv, n_tail, 1, + ggml_row_size(vc_s->type, n_embd_head_v), + ggml_row_size(vc_s->type, n_embd_v_gqa2), + ggml_row_size(vc_s->type, n_embd_v_gqa2 * (uint64_t) vc_s->ne[1]), + 0); + return ggml_concat(ctx, win_v, tail_v, 2); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; + } + // opencoti F4 M3 Phase 4-C: per-stream view + concat along stream // axis. Same shape as get_k, but with v_trans branching. Unified mode // (sinfo.s1 == sinfo.s0 == 0) keeps a single 4-D view → byte-identical. @@ -2035,6 +2428,307 @@ ggml_tensor * llama_kv_cache::get_v(ggml_context * ctx, int32_t il, uint32_t n_k // per-step O(n_kv) CPU→GPU stream-back. Dimensions and strides are identical // to the corresponding halves inside M2's get_k/get_v. +// opencoti F5 M7 Stage 3d-S1 (#351) — compute microbench (auto-adapt loop step 2, +// docs/features/rolling_kv.md). Measure THIS host's GPU flash-attention compute time +// per KV cell so the Stage-3d-S2 tactic selector can size tile_bytes_max = +// compute_ms * eff_bw_gbps * 0.8 and choose stream-vs-CPU-FA from MEASURED hardware +// characteristics — never hardcoded (PCIe 3.0 x1 … 5.0 x16 all auto-adapt). +// +// Host-side: builds a synthetic decode-shape ggml_flash_attn_ext over `bench_cells` +// keys and runs it on a transient backend over the model's device — the already-loaded +// CUDA DSO dispatches the kernel, so NO DSO rebuild is needed. Mirrors the real +// build_attn_mha FA contract (post-permute q[hd,n_q,n_head], k/v[hd,n_kv,n_head_kv], +// F16 mask). Never aborts boot: every shortfall (no GPU, unsupported op, non-GQA +// geometry, alloc miss) returns 0.0f and the selector keeps the stream default. +static float opencoti_fa_compute_probe_ms( + const llama_model & model, const llama_hparams & hparams, + int32_t il, ggml_type type_k, ggml_type type_v, int64_t bench_cells, + bool on_cpu) { + // GPU probe (on_cpu=false): the model's compute device for layer il — the already-loaded + // CUDA DSO dispatches the FA kernel. CPU probe (on_cpu=true, Stage 3d-S2 #352 companion): + // the host CPU backend (ggml-cpu / iqk FA) so the selector can MODEL the no-DMA CPU-FA tail + // and choose stream-vs-CPU from measured hardware — never hardcoded. + ggml_backend_dev_t dev = on_cpu + ? ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU) + : model.dev_layer(il); + if (!dev) { + return 0.0f; + } + // NB: `ggml_backend_dev_type` names both the enum and the accessor function; in C++ + // `ggml_backend_dev_type(dev)` would parse as a functional cast, so read the type via props. + ggml_backend_dev_props dprops = {}; + ggml_backend_dev_get_props(dev, &dprops); + // Compare enumerator values directly — do NOT declare a variable of type + // `ggml_backend_dev_type` (it names both the enum and a function; the function + // hides the type in ordinary lookup, so such a declaration mis-parses). + const bool type_ok = on_cpu + ? (dprops.type == GGML_BACKEND_DEVICE_TYPE_CPU) + : (dprops.type == GGML_BACKEND_DEVICE_TYPE_GPU); + if (!type_ok) { + return 0.0f; // device-type mismatch; selector falls back (stream default / no CPU model) + } + + const int64_t hd_k = (int64_t) hparams.n_embd_head_k(il); + const int64_t hd_v = (int64_t) hparams.n_embd_head_v(il); + const int64_t n_head = (int64_t) hparams.n_head(il); + const int64_t n_head_kv = (int64_t) hparams.n_head_kv(il); + // GQA contract: ggml_flash_attn_ext build asserts can_mul_mat(k,q) ⇒ n_head % n_head_kv == 0. + if (hd_k <= 0 || hd_v <= 0 || n_head <= 0 || n_head_kv <= 0 || (n_head % n_head_kv) != 0 || bench_cells < 256) { + return 0.0f; + } + + ggml_backend_t backend = ggml_backend_dev_init(dev, nullptr); + if (!backend) { + return 0.0f; + } + + const size_t meta_size = ggml_tensor_overhead() * 16 + ggml_graph_overhead(); + ggml_init_params ip = { meta_size, nullptr, /*.no_alloc =*/ true }; + ggml_context * ctx = ggml_init(ip); + if (!ctx) { + ggml_backend_free(backend); + return 0.0f; + } + + // Post-permute decode FA shapes (mirror build_attn_mha): head count lives in ne[2] + // (GQA broadcast q->ne[2] % k->ne[2] == 0); the kq_mask is F16 (the op asserts F16). + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, hd_k, 1, n_head, 1); + ggml_tensor * k = ggml_new_tensor_4d(ctx, type_k, hd_k, bench_cells, n_head_kv, 1); + ggml_tensor * v = ggml_new_tensor_4d(ctx, type_v, hd_v, bench_cells, n_head_kv, 1); + ggml_tensor * mask = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, bench_cells, 1, 1, 1); + const float scale = 1.0f / sqrtf((float) hd_k); + ggml_tensor * fa = ggml_flash_attn_ext(ctx, q, k, v, mask, scale, 0.0f, 0.0f); + ggml_flash_attn_ext_set_prec(fa, GGML_PREC_F32); + + if (!ggml_backend_dev_supports_op(dev, fa)) { + ggml_free(ctx); + ggml_backend_free(backend); + return 0.0f; + } + + ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend); + if (!buf) { + ggml_free(ctx); + ggml_backend_free(backend); + return 0.0f; + } + ggml_backend_buffer_clear(buf, 0); // zero inputs: mask 0 ⇒ attend-all, K/V 0 ⇒ finite output, no NaN + + ggml_cgraph * gf = ggml_new_graph(ctx); + ggml_build_forward_expand(gf, fa); + + // warmup (kernel selection / JIT / first-touch), then time a few iterations + ggml_backend_graph_compute(backend, gf); + ggml_backend_synchronize(backend); + + const int iters = 5; + const auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < iters; ++i) { + ggml_backend_graph_compute(backend, gf); + } + ggml_backend_synchronize(backend); + const auto t1 = std::chrono::steady_clock::now(); + const double total_ms = std::chrono::duration(t1 - t0).count(); + + ggml_backend_buffer_free(buf); + ggml_free(ctx); + ggml_backend_free(backend); + + return (float) (total_ms / (double) iters); // ms for one FA over bench_cells keys +} + +// opencoti F5 M7 Rolling KV — Rung 1 (M7-B). Derive the per-layer tactic table +// from the split decisions already baked into `layers` during construction, and +// log the histogram. At Rung 1 the only tactics that appear are GPU_RESIDENT +// (gpu_heads == 0, full KV on device, vanilla fast path) and CPU_SPILL +// (gpu_heads > 0, the shipped M2 head-split). The streaming tactics +// (GPU_STREAM / GPU_STREAM_SUBTILE) + the tile-sizing fields (eff_bw_gbps from +// W1 pcie_profile threaded via cparams, compute_ms from a live FA microbench, +// tile_bytes / n_slots, and the slot pool) are assigned at Rung 2 alongside the +// double-buffer kernel that consumes them. This keeps the table compute-inert: +// nothing on the attention path reads layer_tactic until Rung 2, so the graph +// at small ctx stays byte-identical to Rung 0 / vanilla. +void llama_kv_cache::build_rolling_kv_plan(float eff_bw_gbps) { + layer_tactic.assign(layers.size(), headinfer_tactic::GPU_RESIDENT); + rkv_plan = rolling_kv_plan{}; + rkv_plan.eff_bw_gbps = eff_bw_gbps; // Stage 3d-S2 (#352): MEASURED host->device PCIe bandwidth (0 = unknown) + + // opencoti F5 M7 Rolling KV — M7-E E1 (#305): stream-by-default. An + // over-budget layer (gpu_heads > 0 ⇒ KV did not fit resident) is served by + // the GPU_STREAM tactic — the custom streaming-FA op that tiles attention + // over the key dimension with a cross-tile online-softmax — instead of + // CPU_SPILL (the W2 iqk CPU-FA engine). + // + // The S-ladder (S1→S3d) proved GPU_STREAM byte-identical on f16 and coherent + // on q8_0 (dequant-on-lift) incl. the n_head_cpu>1 Gemma split, so GPU_STREAM + // is the DEFAULT relief tactic (Decision-4 graduated relief). CPU_SPILL (the W2 + // iqk CPU-FA path) remains in-tree as a tactic but is no longer auto-selected. + // The decision is made ONCE at construction here; the runtime promote/demote- + // under-load scheduler is the deferred E2. + const headinfer_tactic spill_tactic = headinfer_tactic::GPU_STREAM; + + // opencoti F5 M7 Stage 3c-5 — POSITION_WINDOW takes precedence: a window layer + // (window_cells > 0) keeps the bulk device-resident [0,wc) and streams only the + // overflow tail [wc,kv) via the two-region streaming op. It is mutually exclusive + // with the M2/E1 head split (gpu_heads == 0 whenever window_cells > 0). + for (size_t j = 0; j < layers.size(); ++j) { + const headinfer_tactic t = + layers[j].window_cells > 0 ? headinfer_tactic::POSITION_WINDOW : + layers[j].gpu_heads > 0 ? spill_tactic : + headinfer_tactic::GPU_RESIDENT; + layer_tactic[j] = t; + rkv_plan.hist[(int) t]++; + } + + // "active" tracks whether any *streaming* tactic is engaged. + rkv_plan.active = + rkv_plan.hist[(int) headinfer_tactic::GPU_STREAM] > 0 || + rkv_plan.hist[(int) headinfer_tactic::GPU_STREAM_SUBTILE] > 0 || + rkv_plan.hist[(int) headinfer_tactic::POSITION_WINDOW] > 0; + + LLAMA_LOG_INFO("%s: rolling-kv tactic table: %zu KV layers — " + "GPU_RESIDENT=%u GPU_STREAM=%u SUBTILE=%u CPU_SPILL=%u POSITION_WINDOW=%u CPU_FA_TAIL=%u " + "(relief=%s; M7-E E1 / Stage 3c-5)\n", + __func__, layers.size(), + rkv_plan.hist[0], rkv_plan.hist[1], rkv_plan.hist[2], rkv_plan.hist[3], rkv_plan.hist[4], rkv_plan.hist[5], + spill_tactic == headinfer_tactic::GPU_STREAM ? "GPU_STREAM" : "CPU_SPILL"); + + // opencoti F5 M7 Stage 3d-S1 (#351) — run the compute microbench ONLY when a + // POSITION_WINDOW layer exists (overflow is actually happening). At small ctx where + // every layer is GPU_RESIDENT the boot stays byte-identical (no probe, no extra log). + // The measured fa_ms_per_cell is consumed by the Stage-3d-S2 selector; here we just + // populate + log it (Rung-1 compute-inert: the attention path still reads only the + // tactic table, unchanged). + if (rkv_plan.hist[(int) headinfer_tactic::POSITION_WINDOW] > 0) { + int32_t bench_il = -1; + uint32_t bench_wc = 0; + ggml_type tk = GGML_TYPE_F16; + ggml_type tv = GGML_TYPE_F16; + for (size_t j = 0; j < layers.size(); ++j) { + if (layers[j].window_cells > 0 && !layers[j].k_per_stream.empty() && !layers[j].v_per_stream.empty()) { + bench_il = (int32_t) layers[j].il; + bench_wc = layers[j].window_cells; + tk = layers[j].k_per_stream[0]->type; + tv = layers[j].v_per_stream[0]->type; + break; + } + } + if (bench_il >= 0) { + const int64_t bench_cells = 8192; // 256-aligned, ~32 MiB transient K+V; amortizes launch slack + const float probe_ms = opencoti_fa_compute_probe_ms(model, hparams, bench_il, tk, tv, bench_cells, /*on_cpu=*/false); + if (probe_ms > 0.0f) { + rkv_plan.fa_ms_per_cell = probe_ms / (float) bench_cells; + rkv_plan.compute_ms = rkv_plan.fa_ms_per_cell * (float) kv_size_max; // per-layer FA ms @ full ctx (reference) + LLAMA_LOG_INFO("%s: compute microbench — FA %.3f ms / %lld cells = %.6f ms/cell " + "(global layer %d, %s K / %s V); est %.2f ms/layer @ %u cells\n", __func__, + (double) probe_ms, (long long) bench_cells, (double) rkv_plan.fa_ms_per_cell, + bench_il, ggml_type_name(tk), ggml_type_name(tv), + (double) rkv_plan.compute_ms, kv_size_max); + } else { + LLAMA_LOG_INFO("%s: compute microbench — unavailable (no GPU probe); Stage-3d-S2 selector " + "will use the analytic/stream default\n", __func__); + } + + // opencoti F5 M7 Stage 3d-S2 (#352) — overflow tail-tactic crossover selector. + // Runs only when the GPU microbench succeeded (fa_ms_per_cell measured). Decides, from + // MEASURED hardware (GPU fa_ms_per_cell, CPU cpu_fa_ms_per_cell, eff_bw_gbps), whether + // the host overflow tail is best STREAMED to the device (POSITION_WINDOW — hidden under + // FA compute on fast links) or FLASH-ATTENDED on the CPU (CPU_FA_TAIL — no DMA, wins on + // slow links). Nothing hardcoded — the only constant is the documented 0.8 overlap + // safety (pcie-profile.h:8). INFORMATIONAL at S2: layer_tactic stays POSITION_WINDOW so + // the forward is byte-unchanged; rkv_plan.tail_tactic carries the choice for the S2 gate + // and the S3 CPU-FA-tail execution branch. + if (rkv_plan.fa_ms_per_cell > 0.0f) { + // CPU companion probe — same synthetic FA on the host CPU backend (ggml-cpu / iqk). + const float cpu_probe_ms = opencoti_fa_compute_probe_ms(model, hparams, bench_il, tk, tv, bench_cells, /*on_cpu=*/true); + if (cpu_probe_ms > 0.0f) { + rkv_plan.cpu_fa_ms_per_cell = cpu_probe_ms / (float) bench_cells; + } + + const uint32_t wc = bench_wc; + const uint32_t n_tail = (kv_size_max > wc) ? (kv_size_max - wc) : 0u; + const size_t per_cell_bytes = + (size_t) ggml_row_size(tk, hparams.n_embd_k_gqa(bench_il)) + + (size_t) ggml_row_size(tv, hparams.n_embd_v_gqa(bench_il)); + const double tail_bytes = (double) n_tail * (double) per_cell_bytes; + const double fa = (double) rkv_plan.fa_ms_per_cell; + const double cfa = (double) rkv_plan.cpu_fa_ms_per_cell; + + // Per-layer, per-decode-step models (GB/s = 1e6 bytes/ms): + // stream : the GPU must FA over ALL kv_size keys anyway; the tail DMA overlaps + // under that compute → stream_ms = max(fa*kv_size, tail_transfer_ms) + // cpu-tail: GPU does only the window FA; the CPU does the tail FA concurrently, + // no DMA → cpu_fa_tail_ms = max(fa*wc, cfa*n_tail) + const double resident_compute_ms = fa * (double) kv_size_max; + const double tail_transfer_ms = (eff_bw_gbps > 0.0f) + ? tail_bytes / ((double) eff_bw_gbps * 1.0e6) : 0.0; + // tile_bytes_max = bytes hideable under the FA compute (compute_ms*eff_bw*0.8). + rkv_plan.tile_bytes = (eff_bw_gbps > 0.0f) + ? (size_t) (resident_compute_ms * (double) eff_bw_gbps * 1.0e6 * 0.8) : 0; + const double stream_ms = std::max(resident_compute_ms, tail_transfer_ms); + + headinfer_tactic chosen = headinfer_tactic::POSITION_WINDOW; + const char * why = "stream (default)"; + double cpu_fa_tail_ms = 0.0; + if (eff_bw_gbps <= 0.0f) { + why = "stream (bandwidth unknown)"; + } else if (tail_bytes <= (double) rkv_plan.tile_bytes) { + why = "stream (tail <= tile_bytes_max -> hidden under FA compute)"; + } else if (cfa > 0.0f) { + // opencoti #354 — CPU_FA_TAIL is GATED OUT of auto-select. We still MODEL its + // cost (for the audit-trail log + the physics crossover record) but never CHOOSE + // it: routing the host tail through CPU flash-attention fails the single-needle + // RULER retrieval bar. The scalar _f16_one_chunk ceilings at niah~80 (q8=60), and + // the iqk-wired tail partial measured niah=0.0 — while the GPU-stream + // POSITION_WINDOW path holds niah=100. On a slow link this keeps streaming the + // tail (some decode tps for correctness — the right trade). Re-enable only when a + // CPU tail FA is proven >= vanilla-2pp on niah (the #354 gate). + cpu_fa_tail_ms = std::max(fa * (double) wc, cfa * (double) n_tail); + // CPU_FA_TAIL stays MODELED here (audit log) but is gated OUT of auto-select + // (#354): the doctrine-clean GPU-stream POSITION_WINDOW path holds niah=100, so + // even on a slow link we keep streaming the tail — trading some decode tps for + // correctness. The tactic + op remain in-tree, dormant, for a future CPU tail FA + // that is proven >= vanilla-2pp on niah (the #354 re-enable gate). + why = (cpu_fa_tail_ms < stream_ms) + ? "stream (CPU_FA_TAIL gated out #354 — faster but loses the needle)" + : "stream (<= cpu-fa tail)"; + } else { + why = "stream (no CPU probe -> CPU_FA_TAIL unmodeled)"; + } + rkv_plan.tail_tactic = chosen; + + LLAMA_LOG_INFO("%s: Stage-3d-S2 selector — eff_bw=%.2f GB/s; window=%u tail=%u cells (%.1f MiB); " + "fa=%.6f cpu_fa=%.6f ms/cell; resident_compute=%.3f tail_transfer=%.3f stream=%.3f cpu_fa_tail=%.3f ms; " + "tile_bytes_max=%.1f MiB -> tail_tactic=%s [%s]\n", __func__, + (double) eff_bw_gbps, wc, n_tail, tail_bytes / (1024.0*1024.0), + fa, cfa, resident_compute_ms, tail_transfer_ms, stream_ms, cpu_fa_tail_ms, + (double) rkv_plan.tile_bytes / (1024.0*1024.0), + chosen == headinfer_tactic::CPU_FA_TAIL ? "CPU_FA_TAIL" : "POSITION_WINDOW", why); + } + } + } +} + +// opencoti F5 M7 Rolling KV — Rung 1. Tactic for model layer `il`. Empty table +// (Rung 0 path / no auto decision) ⇒ GPU_RESIDENT. +// opencoti F5 M7 Stage 3d-S3 (#353): the S2 selector stored its overflow tail-tactic +// choice (stream vs CPU-FA) in rkv_plan.tail_tactic (plan-wide; the global windowed +// layers are homogeneous). The graph reads this to route the CPU-FA-tail branch. +bool llama_kv_cache::headinfer_tail_is_cpu_fa() const { + return rkv_plan.tail_tactic == headinfer_tactic::CPU_FA_TAIL; +} + +llama_kv_cache::headinfer_tactic llama_kv_cache::get_layer_tactic(int32_t il) const { + if (layer_tactic.empty()) { + return headinfer_tactic::GPU_RESIDENT; + } + const auto it = map_layer_ids.find(il); + if (it == map_layer_ids.end() || (size_t) it->second >= layer_tactic.size()) { + return headinfer_tactic::GPU_RESIDENT; + } + return layer_tactic[it->second]; +} + bool llama_kv_cache::headinfer_split_active(int32_t il) const { const auto it = map_layer_ids.find(il); if (it == map_layer_ids.end()) { @@ -2184,6 +2878,159 @@ ggml_tensor * llama_kv_cache::get_v_cpu(ggml_context * ctx, int32_t il, uint32_t return result; } +// opencoti F5 M7 Stage 3c-5 — POSITION_WINDOW perf-path accessors. Unlike the +// get_k/get_v dim-2 concat (correctness fallback), these return the device window +// and pinned-host tail regions SEPARATELY so build_attn_mha_position_window can run +// a window FA over the resident bulk and stream only the overflow tail. Layout is +// identical to get_k_gpu ([n_embd_head_k, n_head_kv, cells, n_stream]); the caller +// applies the same permute(0,2,1,3) the FA path uses. The device window tensor has +// ne[1] == window_cells, so its dim-3 (stream/position) stride is computed off +// k_s->ne[1] (the audit point — NOT kv_size); the host tail likewise off kc_s->ne[1]. +ggml_tensor * llama_kv_cache::get_k_window(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const { + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + GGML_ASSERT(layer.window_cells > 0 && "get_k_window called without an active position window"); + + const uint32_t n_embd_head_k = hparams.n_embd_head_k(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_k_gqa2 = (uint64_t) n_head_kv * n_embd_head_k; + const uint32_t wc = layer.window_cells; + const uint32_t n_win = n_kv < wc ? n_kv : wc; + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * k_s = layer.k_per_stream[s_cache]; + return ggml_view_4d(ctx, k_s, + n_embd_head_k, n_head_kv, n_win, 1, + ggml_row_size(k_s->type, n_embd_head_k), + ggml_row_size(k_s->type, n_embd_k_gqa2), + ggml_row_size(k_s->type, n_embd_k_gqa2 * (uint64_t) k_s->ne[1]), + 0); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; +} + +ggml_tensor * llama_kv_cache::get_k_tail(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const { + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + GGML_ASSERT(layer.window_cells > 0 + && !layer.k_cpu_per_stream.empty() + && layer.k_cpu_per_stream[0] + && "get_k_tail called without an active position window"); + + const uint32_t n_embd_head_k = hparams.n_embd_head_k(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_k_gqa2 = (uint64_t) n_head_kv * n_embd_head_k; + const uint32_t wc = layer.window_cells; + const uint32_t n_tail = n_kv > wc ? n_kv - wc : 0; + if (n_tail == 0) { + return nullptr; // no overflow → fully resident window, single-region FA + } + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * kc_s = layer.k_cpu_per_stream[s_cache]; + return ggml_view_4d(ctx, kc_s, + n_embd_head_k, n_head_kv, n_tail, 1, + ggml_row_size(kc_s->type, n_embd_head_k), + ggml_row_size(kc_s->type, n_embd_k_gqa2), + ggml_row_size(kc_s->type, n_embd_k_gqa2 * (uint64_t) kc_s->ne[1]), + 0); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; +} + +ggml_tensor * llama_kv_cache::get_v_window(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const { + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + GGML_ASSERT(layer.window_cells > 0 && "get_v_window called without an active position window"); + GGML_ASSERT(!v_trans && "position window requires non-transposed V"); + + const uint32_t n_embd_head_v = hparams.n_embd_head_v(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_v_gqa2 = (uint64_t) n_head_kv * n_embd_head_v; + const uint32_t wc = layer.window_cells; + const uint32_t n_win = n_kv < wc ? n_kv : wc; + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * v_s = layer.v_per_stream[s_cache]; + return ggml_view_4d(ctx, v_s, + n_embd_head_v, n_head_kv, n_win, 1, + ggml_row_size(v_s->type, n_embd_head_v), + ggml_row_size(v_s->type, n_embd_v_gqa2), + ggml_row_size(v_s->type, n_embd_v_gqa2 * (uint64_t) v_s->ne[1]), + 0); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; +} + +ggml_tensor * llama_kv_cache::get_v_tail(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const { + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + GGML_ASSERT(layer.window_cells > 0 + && !layer.v_cpu_per_stream.empty() + && layer.v_cpu_per_stream[0] + && "get_v_tail called without an active position window"); + GGML_ASSERT(!v_trans && "position window requires non-transposed V"); + + const uint32_t n_embd_head_v = hparams.n_embd_head_v(il); + const uint32_t n_head_kv = hparams.n_head_kv(il); + const uint64_t n_embd_v_gqa2 = (uint64_t) n_head_kv * n_embd_head_v; + const uint32_t wc = layer.window_cells; + const uint32_t n_tail = n_kv > wc ? n_kv - wc : 0; + if (n_tail == 0) { + return nullptr; + } + + auto build_stream_view = [&](uint32_t s_cache) -> ggml_tensor * { + ggml_tensor * vc_s = layer.v_cpu_per_stream[s_cache]; + return ggml_view_4d(ctx, vc_s, + n_embd_head_v, n_head_kv, n_tail, 1, + ggml_row_size(vc_s->type, n_embd_head_v), + ggml_row_size(vc_s->type, n_embd_v_gqa2), + ggml_row_size(vc_s->type, n_embd_v_gqa2 * (uint64_t) vc_s->ne[1]), + 0); + }; + + ggml_tensor * result = build_stream_view(sinfo.s0); + for (uint32_t s = sinfo.s0 + 1; s <= sinfo.s1; ++s) { + result = ggml_concat(ctx, result, build_stream_view(s), 3); + } + return result; +} + +bool llama_kv_cache::headinfer_window_active(int32_t il) const { + const auto it = map_layer_ids.find(il); + if (it == map_layer_ids.end()) { + return false; + } + const auto & layer = layers[it->second]; + return layer.window_cells > 0 + && !layer.k_cpu_per_stream.empty() + && layer.k_cpu_per_stream[0] != nullptr; +} + +uint32_t llama_kv_cache::headinfer_window_cells(int32_t il) const { + const auto it = map_layer_ids.find(il); + if (it == map_layer_ids.end()) { + return 0; + } + return layers[it->second].window_cells; +} + ggml_tensor * llama_kv_cache::cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il, const slot_info & sinfo) const { const int32_t ikv = map_layer_ids.at(il); @@ -2258,6 +3105,65 @@ ggml_tensor * llama_kv_cache::cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggm return tie; } + // opencoti F5 M7 Stage 3c — position-window scatter. Route each row to the + // device window (dest cell idx < wc) or the pinned host tail (idx >= wc, at + // local index idx-wc). set_input_k_idxs has already written window-local + // (idx) / tail-local (idx-wc) values into k_idxs. For the append-only base + // cache the destination cells are position-sorted within a stream, so the + // window rows form a contiguous prefix [0,kw) and the tail rows the suffix + // [kw,tps). Two set_rows (window + tail) per stream, tied into one dep node. + // + // The split point kw is read from sinfo.idxs at graph-build. The full-cache + // reserve context installs a DUMMY sinfo whose idxs[s] has a single element + // while the worst-case ubatch has many tokens — so when idxs[s].size() != + // tokens_per_stream we are at reserve and build a conservative both-full + // graph for memory planning (no kernel runs at reserve, so index values are + // irrelevant; only the topology must cover the real graph, which uses <= the + // same two set_rows per stream). + if (layer.window_cells > 0 && !layer.k_cpu_per_stream.empty() && layer.k_cpu_per_stream[0]) { + const uint32_t wc = layer.window_cells; + ggml_tensor * k_cur2 = ggml_view_2d(ctx, k_cur, n_embd_gqa, n_tokens, k_cur->nb[2], 0); + const uint32_t batch_ns = sinfo.n_stream(); + GGML_ASSERT(batch_ns >= 1); + const int64_t tokens_per_stream = (batch_ns == 1) ? n_tokens : (n_tokens / (int64_t) batch_ns); + GGML_ASSERT((int64_t) batch_ns * tokens_per_stream == n_tokens); + + ggml_tensor * tie = nullptr; + auto accumulate = [&](ggml_tensor * store) { + if (!tie) { tie = store; return; } + tie = ggml_add(ctx, + ggml_cast(ctx, ggml_view_1d(ctx, tie, ggml_blck_size(tie->type), 0), GGML_TYPE_F32), + ggml_cast(ctx, ggml_view_1d(ctx, store, ggml_blck_size(store->type), 0), GGML_TYPE_F32)); + }; + auto emit = [&](ggml_tensor * dst, int64_t row0, int64_t nrows, size_t base) { + if (nrows <= 0) return; + ggml_tensor * cur = ggml_view_2d(ctx, k_cur2, n_embd_gqa, nrows, + k_cur2->nb[1], (base + (size_t) row0) * k_cur2->nb[1]); + ggml_tensor * idx = ggml_view_1d(ctx, k_idxs, nrows, (base + (size_t) row0) * k_idxs->nb[0]); + accumulate(ggml_set_rows(ctx, dst, cur, idx)); + }; + for (uint32_t s = 0; s < batch_ns; ++s) { + ggml_tensor * k_s = layer.k_per_stream[sinfo.strm[s]]; + ggml_tensor * kc_s = layer.k_cpu_per_stream[sinfo.strm[s]]; + const size_t base = (size_t) s * tokens_per_stream; + if (sinfo.idxs[s].size() != (size_t) tokens_per_stream) { + // reserve dummy sinfo → conservative both-full graph for planning + emit(k_s, 0, tokens_per_stream, base); + emit(kc_s, 0, tokens_per_stream, base); + continue; + } + int64_t kw = 0; + while (kw < tokens_per_stream && (uint32_t) sinfo.idxs[s][kw] < wc) ++kw; + for (int64_t r = kw; r < tokens_per_stream; ++r) { + GGML_ASSERT((uint32_t) sinfo.idxs[s][r] >= wc + && "position-window scatter requires position-sorted destination idxs"); + } + emit(k_s, 0, kw, base); // window prefix [0,kw) + emit(kc_s, kw, tokens_per_stream - kw, base); // tail suffix [kw,tps) + } + return tie; + } + k_cur = ggml_view_2d(ctx, k_cur, n_embd_gqa, n_tokens, k_cur->nb[2], 0); // opencoti F4 M3 Phase 4-C: per-stream scatter via N separate @@ -2389,6 +3295,46 @@ ggml_tensor * llama_kv_cache::cpy_v(ggml_context * ctx, ggml_tensor * v_cur, ggm GGML_ASSERT(batch_ns >= 1); if (!v_trans) { + // opencoti F5 M7 Stage 3c — position-window scatter (mirror of cpy_k). + // window mode is gated to !v_trans, so V is always reached here. Same + // window prefix / tail suffix split + reserve-dummy guard as cpy_k. + if (layer.window_cells > 0 && !layer.v_cpu_per_stream.empty() && layer.v_cpu_per_stream[0]) { + const uint32_t wc = layer.window_cells; + ggml_tensor * v_cur2 = ggml_view_2d(ctx, v_cur, n_embd_gqa, n_tokens, v_cur->nb[2], 0); + const int64_t tokens_per_stream = (batch_ns == 1) ? n_tokens : (n_tokens / (int64_t) batch_ns); + GGML_ASSERT((int64_t) batch_ns * tokens_per_stream == n_tokens); + + ggml_tensor * tie = nullptr; + auto accumulate = [&](ggml_tensor * store) { + if (!tie) { tie = store; return; } + tie = ggml_add(ctx, + ggml_cast(ctx, ggml_view_1d(ctx, tie, ggml_blck_size(tie->type), 0), GGML_TYPE_F32), + ggml_cast(ctx, ggml_view_1d(ctx, store, ggml_blck_size(store->type), 0), GGML_TYPE_F32)); + }; + auto emit = [&](ggml_tensor * dst, int64_t row0, int64_t nrows, size_t base) { + if (nrows <= 0) return; + ggml_tensor * cur = ggml_view_2d(ctx, v_cur2, n_embd_gqa, nrows, + v_cur2->nb[1], (base + (size_t) row0) * v_cur2->nb[1]); + ggml_tensor * idx = ggml_view_1d(ctx, v_idxs, nrows, (base + (size_t) row0) * v_idxs->nb[0]); + accumulate(ggml_set_rows(ctx, dst, cur, idx)); + }; + for (uint32_t s = 0; s < batch_ns; ++s) { + ggml_tensor * v_s = layer.v_per_stream[sinfo.strm[s]]; + ggml_tensor * vc_s = layer.v_cpu_per_stream[sinfo.strm[s]]; + const size_t base = (size_t) s * tokens_per_stream; + if (sinfo.idxs[s].size() != (size_t) tokens_per_stream) { + emit(v_s, 0, tokens_per_stream, base); + emit(vc_s, 0, tokens_per_stream, base); + continue; + } + int64_t kw = 0; + while (kw < tokens_per_stream && (uint32_t) sinfo.idxs[s][kw] < wc) ++kw; + emit(v_s, 0, kw, base); + emit(vc_s, kw, tokens_per_stream - kw, base); + } + return tie; + } + v_cur = ggml_view_2d(ctx, v_cur, n_embd_gqa, n_tokens, v_cur->nb[2], 0); if (batch_ns == 1) { @@ -2535,9 +3481,15 @@ void llama_kv_cache::set_input_k_idxs(ggml_tensor * dst, const llama_ubatch * ub // global indices wrote out-of-bounds against the per-stream // tensor and crashed CUDA decode. + // opencoti F5 M7 Stage 3c — position-window local indexing. When the cache is + // windowed, a destination cell idx >= wc lives in the host tail at local index + // idx-wc; idx < wc stays window-local. wc is uniform across the cache's layers. + // wc == 0 (no window) ⇒ byte-identical to the vanilla local indexing below. + const uint32_t wc_k = layers.empty() ? 0u : layers.front().window_cells; for (uint32_t s = 0; s < sinfo.n_stream(); ++s) { for (uint32_t i = 0; i < sinfo.size(); ++i) { - data[s*sinfo.size() + i] = sinfo.idxs[s][i]; + const uint32_t idx = (uint32_t) sinfo.idxs[s][i]; + data[s*sinfo.size() + i] = (wc_k > 0 && idx >= wc_k) ? (int64_t) (idx - wc_k) : (int64_t) idx; } } } @@ -2555,9 +3507,14 @@ void llama_kv_cache::set_input_v_idxs(ggml_tensor * dst, const llama_ubatch * ub // byte-identical regardless (sinfo.strm[0] == 0 makes offs == 0). if (!v_trans) { + // opencoti F5 M7 Stage 3c — position-window local indexing (mirror of + // set_input_k_idxs). window mode is gated to !v_trans, so the tail-local + // adjustment only applies here. wc == 0 ⇒ byte-identical to vanilla. + const uint32_t wc_v = layers.empty() ? 0u : layers.front().window_cells; for (uint32_t s = 0; s < sinfo.n_stream(); ++s) { for (uint32_t i = 0; i < sinfo.size(); ++i) { - data[s*sinfo.size() + i] = sinfo.idxs[s][i]; + const uint32_t idx = (uint32_t) sinfo.idxs[s][i]; + data[s*sinfo.size() + i] = (wc_v > 0 && idx >= wc_v) ? (int64_t) (idx - wc_v) : (int64_t) idx; } } } else { @@ -3234,6 +4191,34 @@ void llama_kv_cache::state_write_data(llama_io_write_i & io, const cell_ranges_t io.write_tensor(k, range.first * k_row_gpu, range_size * k_row_gpu); io.write_tensor(k_cpu, range.first * k_row_cpu, range_size * k_row_cpu); } + } else if (layer.window_cells > 0 + && !layer.k_cpu_per_stream.empty() + && layer.k_cpu_per_stream[cr.strm]) { + // opencoti F5 M7 Stage 3c-6 — position-window split-aware state write (K). + // In window mode k (k_stream) holds the resident cells [0,wc) and k_cpu the + // overflow tail [wc,kv_size) at local idx-wc; both carry ALL heads, so each + // cell's row is k_size_row. Emit each range's cells in LOGICAL position order + // (window prefix [a,min(b,wc)) then tail suffix [max(a,wc),b)), so the on-disk + // bytes are byte-identical to a vanilla / M2 save → a window save loads into a + // resident cache and vice-versa (cross-residency-mode compatible). + auto * k_cpu = layer.k_cpu_per_stream[cr.strm]; + const uint32_t wc = layer.window_cells; + GGML_ASSERT((uint64_t) ggml_row_size(k_cpu->type, n_embd_k_gqa) == k_size_row); + for (const auto & range : cr.data) { + const uint32_t a = range.first, b = range.second; + const uint32_t w_hi = b < wc ? b : wc; // window cells [a, w_hi) + const uint32_t t_lo = a > wc ? a : wc; // tail cells [t_lo, b) + size_t wrote = 0; + if (w_hi > a) { + io.write_tensor(k, (size_t) a * k_size_row, (size_t) (w_hi - a) * k_size_row); + wrote += (size_t) (w_hi - a) * k_size_row; + } + if (b > t_lo) { + io.write_tensor(k_cpu, (size_t) (t_lo - wc) * k_size_row, (size_t) (b - t_lo) * k_size_row); + wrote += (size_t) (b - t_lo) * k_size_row; + } + GGML_ASSERT(wrote == (size_t) (b - a) * k_size_row && "3c-6 K window/tail parity"); + } } else { // Read each range of cells of k_size length and write out for (const auto & range : cr.data) { @@ -3283,6 +4268,30 @@ void llama_kv_cache::state_write_data(llama_io_write_i & io, const cell_ranges_t io.write_tensor(v, range.first * v_row_gpu, range_size * v_row_gpu); io.write_tensor(v_cpu, range.first * v_row_cpu, range_size * v_row_cpu); } + } else if (layer.window_cells > 0 + && !layer.v_cpu_per_stream.empty() + && layer.v_cpu_per_stream[cr.strm]) { + // opencoti F5 M7 Stage 3c-6 — position-window split-aware state write (V). + // Mirror of the K branch: window prefix from v (k_stream) then tail suffix + // from v_cpu, in logical cell order → byte-identical on-disk layout. + auto * v_cpu = layer.v_cpu_per_stream[cr.strm]; + const uint32_t wc = layer.window_cells; + GGML_ASSERT((uint64_t) ggml_row_size(v_cpu->type, n_embd_v_gqa) == v_size_row); + for (const auto & range : cr.data) { + const uint32_t a = range.first, b = range.second; + const uint32_t w_hi = b < wc ? b : wc; + const uint32_t t_lo = a > wc ? a : wc; + size_t wrote = 0; + if (w_hi > a) { + io.write_tensor(v, (size_t) a * v_size_row, (size_t) (w_hi - a) * v_size_row); + wrote += (size_t) (w_hi - a) * v_size_row; + } + if (b > t_lo) { + io.write_tensor(v_cpu, (size_t) (t_lo - wc) * v_size_row, (size_t) (b - t_lo) * v_size_row); + wrote += (size_t) (b - t_lo) * v_size_row; + } + GGML_ASSERT(wrote == (size_t) (b - a) * v_size_row && "3c-6 V window/tail parity"); + } } else { // Read each range of cells of v_size length and write out for (const auto & range : cr.data) { @@ -3530,6 +4539,28 @@ bool llama_kv_cache::state_read_data(llama_io_read_i & io, uint32_t strm, uint32 io.read_tensor(k_cpu, sinfo.idxs[0][i] * k_row_cpu, k_row_cpu); } } + } else if (layer.window_cells > 0 + && !layer.k_cpu_per_stream.empty() + && layer.k_cpu_per_stream[strm]) { + // opencoti F5 M7 Stage 3c-6 — symmetric position-window read (K). The + // on-disk stream is cell_count rows in logical position order (vanilla + // layout, written above); route each row to the device window (dest + // position < wc) or the host tail (>= wc, local pos-wc) by its + // destination cell position. Handles contiguous and scattered sinfo + // uniformly (a contiguous run can still straddle wc). New io.read_tensor + // idiom (no ptr-returning io.read in this base): each read_tensor + // consumes one row from the stream in logical order. + auto * k_cpu = layer.k_cpu_per_stream[strm]; + const uint32_t wc = layer.window_cells; + for (uint32_t i = 0; i < cell_count; ++i) { + const uint32_t p = sinfo.is_contiguous() + ? (sinfo.head() + i) : (uint32_t) sinfo.idxs[0][i]; + if (p < wc) { + io.read_tensor(k, (size_t) p * k_size_row, k_size_row); + } else { + io.read_tensor(k_cpu, (size_t) (p - wc) * k_size_row, k_size_row); + } + } } else if (sinfo.is_contiguous()) { // Fast path: contiguous cells, single memcpy io.read_tensor(k, sinfo.head() * k_size_row, cell_count * k_size_row); @@ -3599,6 +4630,24 @@ bool llama_kv_cache::state_read_data(llama_io_read_i & io, uint32_t strm, uint32 io.read_tensor(v_cpu, sinfo.idxs[0][i] * v_row_cpu, v_row_cpu); } } + } else if (layer.window_cells > 0 + && !layer.v_cpu_per_stream.empty() + && layer.v_cpu_per_stream[strm]) { + // opencoti F5 M7 Stage 3c-6 — symmetric position-window read (V). + // Mirror of the K branch: route each row to v (window) or v_cpu + // (tail, local pos-wc) by its destination cell position. New + // io.read_tensor idiom — one row per call, sequential stream order. + auto * v_cpu = layer.v_cpu_per_stream[strm]; + const uint32_t wc = layer.window_cells; + for (uint32_t i = 0; i < cell_count; ++i) { + const uint32_t p = sinfo.is_contiguous() + ? (sinfo.head() + i) : (uint32_t) sinfo.idxs[0][i]; + if (p < wc) { + io.read_tensor(v, (size_t) p * v_size_row, v_size_row); + } else { + io.read_tensor(v_cpu, (size_t) (p - wc) * v_size_row, v_size_row); + } + } } else if (sinfo.is_contiguous()) { // Fast path: contiguous cells, single memcpy io.read_tensor(v, sinfo.head() * v_size_row, cell_count * v_size_row); @@ -3787,6 +4836,27 @@ ggml_tensor * llama_kv_cache_context::get_v_cpu(ggml_context * ctx, int32_t il) return kv->get_v_cpu(ctx, il, n_kv, sinfos[i_cur]); } +// opencoti F5 M7 Stage 3c-5 — per-batch wrappers around the position-window accessors. +ggml_tensor * llama_kv_cache_context::get_k_window(ggml_context * ctx, int32_t il) const { + return kv->get_k_window(ctx, il, n_kv, sinfos[i_cur]); +} + +ggml_tensor * llama_kv_cache_context::get_v_window(ggml_context * ctx, int32_t il) const { + return kv->get_v_window(ctx, il, n_kv, sinfos[i_cur]); +} + +ggml_tensor * llama_kv_cache_context::get_k_tail(ggml_context * ctx, int32_t il) const { + return kv->get_k_tail(ctx, il, n_kv, sinfos[i_cur]); +} + +ggml_tensor * llama_kv_cache_context::get_v_tail(ggml_context * ctx, int32_t il) const { + return kv->get_v_tail(ctx, il, n_kv, sinfos[i_cur]); +} + +bool llama_kv_cache_context::headinfer_window_active(int32_t il) const { + return kv->headinfer_window_active(il); +} + bool llama_kv_cache_context::headinfer_split_active(int32_t il) const { return kv->headinfer_split_active(il); } @@ -3795,6 +4865,16 @@ uint32_t llama_kv_cache_context::headinfer_gpu_heads(int32_t il) const { return kv->headinfer_gpu_heads(il); } +// opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304). +llama_kv_cache::headinfer_tactic llama_kv_cache_context::get_layer_tactic(int32_t il) const { + return kv->get_layer_tactic(il); +} + +// opencoti F5 M7 Stage 3d-S3 (#353). +bool llama_kv_cache_context::headinfer_tail_is_cpu_fa() const { + return kv->headinfer_tail_is_cpu_fa(); +} + ggml_tensor * llama_kv_cache_context::cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il) const { return kv->cpy_k(ctx, k_cur, k_idxs, il, sinfos[i_cur]); } diff --git a/llama.cpp/src/llama-kv-cache.h b/llama.cpp/src/llama-kv-cache.h --- a/llama.cpp/src/llama-kv-cache.h +++ b/llama.cpp/src/llama-kv-cache.h @@ -116,6 +116,21 @@ public: // to host, streamed back per step). 1.0 = off. Effective // only for offloaded layers with !v_trans (flash-attn). float headinfer_gpu_heads_frac, + // opencoti F5 M7 Rolling KV — Rung 0 — see docs/features/rolling_kv.md + // VRAM budget cap (MiB) for auto residency. 0 = use all free + // VRAM (maximize). Only consulted when headinfer_gpu_heads_frac + // is negative (auto); ignored for an explicit fraction. + uint32_t vram_target_mib, + // opencoti F5 M7 Rolling KV — Rung 2 (Stage 3d-S2 #352). Effective + // host->device PCIe bandwidth (GB/s), resolved at boot from the W1 + // pcie_profile (ReBAR probe / sysfs / --pcie-bw-gbps override). Fed to + // build_rolling_kv_plan so the tail-tactic selector sizes the stream + // vs CPU-FA crossover from MEASURED bandwidth. 0 = unknown (stream default). + float pcie_bw_gbps, + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency-tactic knob. + // 0 = auto (POSITION_WINDOW when eligible, else M2 head split), + // 1 = head (force M2 split), 2 = window (force POSITION_WINDOW). + uint32_t kv_residency_mode, uint32_t n_seq_max, uint32_t n_pad, uint32_t n_swa, @@ -234,6 +249,28 @@ public: bool headinfer_split_active(int32_t il) const; uint32_t headinfer_gpu_heads (int32_t il) const; + // opencoti F5 M7 Stage 3c-5 — position-window accessors for the POSITION_WINDOW + // perf path (build_attn_mha_position_window). Unlike get_k/get_v (which dim-2 + // concat the window + tail back into one [n_embd, n_kv] view — the correctness + // fallback), these return the RAW resident-window and pinned-host-tail region + // views separately so the streaming op can keep the bulk device-resident and DMA + // only the overflow. ALL heads, sliced on the POSITION axis: + // get_k/v_window → device view, n_win = min(n_kv, window_cells) cells [0,n_win) + // get_k/v_tail → host view, n_tail = n_kv > wc ? n_kv - wc : 0 cells, or + // nullptr when n_kv <= window_cells (no overflow → resident). + // Callers MUST gate on headinfer_window_active(il); for non-window layers the + // window accessors fall back to the equivalent of get_k/get_v and tail is nullptr. + ggml_tensor * get_k_window(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; + ggml_tensor * get_v_window(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; + ggml_tensor * get_k_tail (ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; + ggml_tensor * get_v_tail (ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; + bool headinfer_window_active(int32_t il) const; + uint32_t headinfer_window_cells (int32_t il) const; + // opencoti F5 M7 Stage 3d-S3 (#353): true when the S2 selector chose CPU_FA_TAIL + // (FA the host tail on the CPU, no DMA) for the windowed layers — read by the + // graph to route build_attn_mha_position_window down the CPU-partial branch. + bool headinfer_tail_is_cpu_fa() const; + // store k_cur and v_cur in the cache based on the provided head location ggml_tensor * cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il, const slot_info & sinfo) const; ggml_tensor * cpy_v(ggml_context * ctx, ggml_tensor * v_cur, ggml_tensor * v_idxs, int32_t il, const slot_info & sinfo) const; @@ -277,6 +314,46 @@ public: void set_input_k_rot(ggml_tensor * dst) const; void set_input_v_rot(ggml_tensor * dst) const; + // opencoti F5 M7 Rolling KV — Rung 1 (M7-B/C) runtime tactic table. + // Per-KV-layer residency tactic, decided from VRAM budget + PCIe bandwidth + // at construction and re-decidable at runtime by later rungs. Rung 1 + // POPULATES + LOGS this table and the streaming plan but does NOT consume + // them in the compute path (the build_attn_mha dispatch + slot pool land at + // Rung 2 / M7-D/E). At small ctx where KV fits, every layer is + // GPU_RESIDENT and the graph is byte-identical to Rung 0 / vanilla. + enum class headinfer_tactic : uint8_t { + GPU_RESIDENT = 0, // full KV on device, no DMA (vanilla fast path) + GPU_STREAM = 1, // ping-pong tile streamed per step (Rung 2) + GPU_STREAM_SUBTILE = 2, // tile > slot: online-softmax subtile (Rung 2) + CPU_SPILL = 3, // M2 head-split CPU half (shipped tactic) + POSITION_WINDOW = 4, // Stage 3c: device window [0,wc) + host tail [wc,kv) streamed + CPU_FA_TAIL = 5, // Stage 3d-S2: device window [0,wc) + CPU-FA host tail (no DMA) + }; + + struct rolling_kv_plan { + bool active = false; // any layer not GPU_RESIDENT-or-spill-only + float eff_bw_gbps = 0.0f; // host->device bandwidth (W1 pcie_profile) + float compute_ms = 0.0f; // est. FA ms / layer at target ctx + float fa_ms_per_cell = 0.0f; // opencoti Stage 3d-S1 (#351): MEASURED GPU flash-attn ms per KV cell (compute microbench). 0 ⇒ unmeasured (selector falls back to the stream default). + float cpu_fa_ms_per_cell = 0.0f; // opencoti Stage 3d-S2 (#352): MEASURED CPU flash-attn ms per KV cell (companion probe). 0 ⇒ unmeasured (CPU_FA_TAIL not modeled → stream). + size_t tile_bytes = 0; // tile_bytes_max = compute_ms·eff_bw·0.8 (bytes hideable under FA compute) + uint32_t n_slots = 0; // ping-pong slots [2,4]; 0 until Rung 2 allocates + // opencoti Stage 3d-S2 (#352): the bandwidth-vs-CPU crossover choice for the + // OVERFLOW tail, decided from MEASURED hardware (fa_ms_per_cell, cpu_fa_ms_per_cell, + // eff_bw_gbps) — nothing hardcoded (auto-adapts PCIe 3.0 x1 … 5.0 x16). Either + // POSITION_WINDOW (stream the tail, hidden under FA compute on fast links) or + // CPU_FA_TAIL (FA the host tail on the CPU, no DMA, on slow links). Informational at + // S2 (layer_tactic stays POSITION_WINDOW so the forward is unchanged); S3 consumes it. + headinfer_tactic tail_tactic = headinfer_tactic::POSITION_WINDOW; + uint32_t hist[6] = { 0, 0, 0, 0, 0, 0 }; // tactic histogram, indexed by headinfer_tactic (CPU_FA_TAIL==5) + }; + + const rolling_kv_plan & get_rolling_kv_plan() const { return rkv_plan; } + + // Tactic for KV-cache layer `il` (model layer id). GPU_RESIDENT when the + // table is empty (Rung 0 path / no auto decision). + headinfer_tactic get_layer_tactic(int32_t il) const; + private: const llama_model & model; const llama_hparams & hparams; @@ -320,6 +397,15 @@ private: std::vector v_cpu_per_stream; uint32_t gpu_heads = 0; + // opencoti F5 M7 Stage 3c — position-window residency (mutually exclusive + // with the head split above: gpu_heads == 0 whenever this is set). When + // window_cells > 0 the device tensors k_per_stream[s]/v_per_stream[s] hold + // ALL heads but only the resident window [0, window_cells) — so ne[1] == + // window_cells, NOT kv_size (the audit point) — and k_cpu_per_stream[s]/ + // v_cpu_per_stream[s] hold ALL heads for the pinned-host tail + // [window_cells, kv_size). 0 ⇒ no position window (vanilla / M2 layout). + uint32_t window_cells = 0; + // Convenience accessors for unified-mode (n_stream == 1) call // sites. Returns nullptr in non-unified mode — that's the // signal that the caller must use per-stream access instead. @@ -415,6 +501,21 @@ private: // model layer id -> KV cache layer id std::unordered_map map_layer_ids; + // opencoti F5 M7 Rolling KV — Rung 1. Per-KV-cache-layer tactic (same index + // space as `layers`), and the derived streaming plan. Populated + logged at + // construction by build_rolling_kv_plan(); inert until Rung 2. Empty vector + // ⇒ every layer GPU_RESIDENT (Rung 0 path). + std::vector layer_tactic; + rolling_kv_plan rkv_plan; + + // opencoti F5 M7 Rolling KV — Rung 1. Derive layer_tactic[] + rkv_plan from + // the per-layer split decision already baked into `layers` (gpu_heads > 0 ⇒ + // CPU_SPILL, else GPU_RESIDENT), then log the histogram. Called once at the + // end of construction. Compute-inert: nothing reads the table until Rung 2. + // eff_bw_gbps (Stage 3d-S2 #352) feeds the overflow tail-tactic crossover + // selector (stream vs CPU-FA); 0 = unknown ⇒ stream default. + void build_rolling_kv_plan(float eff_bw_gbps); + size_t total_size() const; size_t size_k_bytes() const; @@ -509,6 +610,22 @@ public: bool headinfer_split_active(int32_t il) const; uint32_t headinfer_gpu_heads (int32_t il) const; + // opencoti F5 M7 Stage 3c-5 — per-batch POSITION_WINDOW accessors. Delegate to + // the underlying cache's get_k/v_window / get_k/v_tail with the current n_kv + + // slot info. get_k/v_tail return nullptr when n_kv <= window_cells (no overflow). + ggml_tensor * get_k_window(ggml_context * ctx, int32_t il) const; + ggml_tensor * get_v_window(ggml_context * ctx, int32_t il) const; + ggml_tensor * get_k_tail (ggml_context * ctx, int32_t il) const; + ggml_tensor * get_v_tail (ggml_context * ctx, int32_t il) const; + bool headinfer_window_active(int32_t il) const; + + // opencoti F5 M7 Rolling KV — Rung 2 (W4 M7-D #304). Per-layer tactic for + // the streaming dispatch at llama-graph.cpp build_attn. Delegates to the + // underlying cache's tactic table (built at ctor by build_rolling_kv_plan). + llama_kv_cache::headinfer_tactic get_layer_tactic(int32_t il) const; + // opencoti F5 M7 Stage 3d-S3 (#353): delegate the CPU_FA_TAIL tail-tactic query. + bool headinfer_tail_is_cpu_fa() const; + // store k_cur and v_cur in the cache based on the provided head location // note: the heads in k_cur and v_cur should be laid out contiguously in memory // - k_cur [n_embd_head_k, n_head_k, n_tokens] diff --git a/llama.cpp/src/llama-memory-hybrid-iswa.cpp b/llama.cpp/src/llama-memory-hybrid-iswa.cpp --- a/llama.cpp/src/llama-memory-hybrid-iswa.cpp +++ b/llama.cpp/src/llama-memory-hybrid-iswa.cpp @@ -50,6 +50,12 @@ llama_memory_hybrid_iswa::llama_memory_hybrid_iswa( 0, // opencoti F5 M2 headinfer — hybrid-iswa does not thread headinfer; 1.0 = off. 1.0f, + // opencoti F5 M7 Rolling KV — hybrid-iswa does not thread auto residency; 0 = unused. + 0, + // opencoti F5 M7 Stage 3d-S2 — hybrid-iswa does not thread PCIe bandwidth; 0 = unknown (stream default). + 0.0f, + // opencoti F5 M7 Stage 4 — hybrid-iswa does not thread the residency-mode knob; 0 = auto. + 0, n_seq_max, n_ubatch, n_pad, diff --git a/llama.cpp/src/llama-memory-hybrid.cpp b/llama.cpp/src/llama-memory-hybrid.cpp --- a/llama.cpp/src/llama-memory-hybrid.cpp +++ b/llama.cpp/src/llama-memory-hybrid.cpp @@ -49,6 +49,12 @@ llama_memory_hybrid::llama_memory_hybrid( 0, // opencoti F5 M2 headinfer — hybrid memory does not thread headinfer; 1.0 = off. 1.0f, + // opencoti F5 M7 Rolling KV — hybrid memory does not thread auto residency; 0 = unused. + 0, + // opencoti F5 M7 Stage 3d-S2 — hybrid memory does not thread PCIe bandwidth; 0 = unknown (stream default). + 0.0f, + // opencoti F5 M7 Stage 4 — hybrid memory does not thread the residency-mode knob; 0 = auto. + 0, n_seq_max, n_pad, n_swa, diff --git a/llama.cpp/src/llama-model.cpp b/llama.cpp/src/llama-model.cpp --- a/llama.cpp/src/llama-model.cpp +++ b/llama.cpp/src/llama-model.cpp @@ -2079,6 +2079,12 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, cparams.slot_shrink_idle_ms, // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md cparams.headinfer_gpu_heads_frac, + // opencoti F5 M7 Rolling KV — Rung 0 — see docs/features/rolling_kv.md + cparams.vram_target_mib, + // opencoti F5 M7 Rolling KV — Rung 2 (Stage 3d-S2 #352) + cparams.pcie_bw_gbps, + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency knob + cparams.kv_residency_mode, cparams.n_seq_max, cparams.n_ubatch, 1, @@ -2101,6 +2107,12 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, cparams.slot_shrink_idle_ms, // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md cparams.headinfer_gpu_heads_frac, + // opencoti F5 M7 Rolling KV — Rung 0 — see docs/features/rolling_kv.md + cparams.vram_target_mib, + // opencoti F5 M7 Rolling KV — Rung 2 (Stage 3d-S2 #352) + cparams.pcie_bw_gbps, + // opencoti F5 M7 Rolling KV — Stage 4 (#338) residency knob + cparams.kv_residency_mode, cparams.n_seq_max, 1, hparams.n_swa, diff --git a/llama.cpp/tools/server/server.cpp b/llama.cpp/tools/server/server.cpp --- a/llama.cpp/tools/server/server.cpp +++ b/llama.cpp/tools/server/server.cpp @@ -145,6 +145,11 @@ int server_main(int argc, char ** argv, LOG_INF("%s: pcie profile: %.1f GB/s effective (x%d @ %.1f GT/s, source=%s) — %s\n", __func__, (double) pp.effective_bw_gbps, pp.link_width, (double) pp.link_gtps, pp.source.c_str(), pp.detail.c_str()); + // opencoti F5 M7 Stage 3d-S2 (#352): write the RESOLVED bandwidth back into params so it + // flows params → cparams.pcie_bw_gbps (common.cpp) → the KV-cache ctor → build_rolling_kv_plan's + // overflow tail-tactic selector. Under autodetect params.pcie_bw_gbps starts 0; this captures + // the detected/override value. ctx_server.load_model(params) (below) reads the updated params. + params.pcie_bw_gbps = pp.effective_bw_gbps; } server_http_context ctx_http;