diff --git "a/patches/0078-dca.patch" "b/patches/0078-dca.patch" new file mode 100644--- /dev/null +++ "b/patches/0078-dca.patch" @@ -0,0 +1,3515 @@ +From: opencoti +Subject: [PATCH 0078] DCA — Dual Chunk Attention for Gemma 4 A4B + Qwen2/3/3.5 (#593) + +Dual Chunk Attention (DCA) for long-context prefill on a single 3090, enabling the +Gemma 4 26B-A4B (iSWA) 1M-context path and the qwen2 / qwen3 / qwen3moe / qwen35 / +qwen35moe families. Chunked attention with per-regime RoPE (INTRA same-chunk causal, +SUCC prior chunk, INTER older chunks) lets a 262144-trained model serve far beyond +its native ctx-cap (the server slot honors >n_ctx_train when --dca on). + +The fused multi-chunk kernel dispatches on head_dim (#623): the DCA switch selects +dispatch_gqa for DKQ in {128 (qwen3), 256 (qwen35moe), 512 (gemma4)} instead of +the old hardcoded <512,512>, and the dca_fused path forces nstages<=1 + Q_in_reg=false +at all three kernel sites (iter, process_tile, host) so the narrower-head Qwen configs +(nstages_target=2, Q_in_reg=true natively) don't skip the K-load or read unpopulated Q +registers -> NaN. These clamps are no-ops on Gemma's DKQ=512 (native nstages_target=1, +Q_in_reg=false): the compiled kernel for the 512 path is byte-identical pre/post-fix. + +Flags (common/arg.cpp): --dca on|off, --dca-chunk-size N (default 8192), +--dca-yarn-factor F. A new LSE-emitting flash-attn op (GGML_OP_FLASH_ATTN_EXT_LSE) +streams the three regimes and combines them via log-sum-exp; the multi-chunk path +fuses all three key-tile phases into ONE kernel pass (no host LSE combine, no +streaming-LSE recompute). Two host-side optimisations keep the dense-mask cost +independent of n_kv: O6 derives the SUCC/INTER band bounds analytically (killing 2 +of 3 dense masks), and O9 narrows the INTRA mask to the single kernel-read window +[cq*c, (cq+1)*c) under analytical_bands (chunk_size % n_ubatch == 0). + +MEASURED on RTX 5000-class (bf16 KV, v7-coder Q4_K_M, NIAH needle): DCA-on retrieves +the needle at 2/4/8/16/32 chunks (16k..256k) and, after O9, prefill throughput at +256k is at parity with native (~1.0x, down from the pre-opt 1.59x). Scales to +512k/768k/1M (DCA-only — past the native cap). INERT when --dca off (the default): +the build is byte-identical to stock on the non-DCA path. + +qwen35 MTP loader (#622): the qwen35 arch is a Qwen3-Next hybrid that may carry an MTP +(nextn) head as the last block (e.g. Qwen3.6-27B-Omnimerge-v4: block_count=65, +nextn_predict_layers=1, blk.64 = MTP). The loader now reads nextn_predict_layers, +excludes the MTP block from the recurrent/transformer loops and the n_layer switch +(llama-model.cpp), loads blk..nextn.* instead of ssm/attn for the last +nextn_predict_layers blocks, and skips the MTP block in llm_build_qwen35 forward +(qwen35.cpp). llama-arch.cpp re-classifies the six NEXTN_* tensors +LLM_TENSOR_LAYER_OUTPUT -> LLM_TENSOR_LAYER_REPEATING (matching upstream): they are +stored per-block, so the loader must not fault on the block index. This makes the +qwen35 MTP model loadable and, with DCA on, the fused multi-chunk path over its 16 +hd256 full-attn layers is coherence-gate verified. + +Context-shift fragmentation fallback (#617B): #617's analytical INDEX bands +(chunk = key_index / chunk_size) assume cache cell-index == absolute position. A classic +context shift (seq_rm + seq_add past n_ctx) relabels pos[] in place WITHOUT compaction, +so index != pos permanently and the index-based kernel visits the wrong physical cells +(model collapses to a token-loop). A sticky is_fragmented latch (llama-kv-cells.h, set on +pos_add / cell-rm, cleared only on clear()) gates a dense POSITION-based SUCC/INTER+INTRA +mask fallback (set_input_dca fill + two extra op mask srcs, fattn-mma full-range-when- +masked): nullptr on the contiguous fast path (zero extra VRAM/compute, byte-identical to +the non-fragmented binary), dense n_kv-wide masks only once fragmented. Reachable only on +full-attention DCA archs (qwen2/3/3.5) generating past n_ctx; iSWA archs (gemma4) disable +KV-shift entirely so never hit it. + +Decode chunk-boundary fix (#635): the INTER-regime query position is 2*c, NOT c. With +#617's K pre-roped once at j%c, an INTER-Q roped at the constant c gives relative distance +c-(j%c) in [1,c], which OVERLAPS the recent/SUCC window (SUCC true rel runs to 2c-1) — a +chunk-end far key then masquerades as the most-recent token and floods ~50% of attention +mass at the first cq>=2 decode query, collapsing greedy decode ~2x before the model's +intrinsic limit. INTER-Q = 2*c puts every far key at rel [c+1,2c], strictly beyond the +window edge c: retrievable but unable to out-compete genuine recent keys. One host line +(set_input_dca in llama-kv-cache.cpp). Verified: greedy count-up matches the DCA-off +intrinsic break at chunk 256/512/1024 in BOTH decode modes, and a NIAH needle planted in +a far INTER chunk is still retrieved. Decode keeps the upstream stream-k path (an earlier +single-block-per-tile decode workaround was built on a falsified stream-k-FTZ-combine +hypothesis and has been reverted). + +diff --git a/llama.cpp/BUILD.mk b/llama.cpp/BUILD.mk +--- a/llama.cpp/BUILD.mk ++++ b/llama.cpp/BUILD.mk +@@ -199,6 +199,7 @@ LLAMA_SRCS_CPP := \ + llama.cpp/src/models/talkie.cpp \ + llama.cpp/src/models/wavtokenizer-dec.cpp \ + llama.cpp/src/models/xverse.cpp \ ++ llama.cpp/src/dca.cpp \ + llama.cpp/src/llama-adapter.cpp \ + llama.cpp/src/llama-arch.cpp \ + llama.cpp/src/llama-batch.cpp \ +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 +@@ -1505,6 +1505,34 @@ common_params_context common_params_parser_init(common_params & params, llama_ex + 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 dca — Dual Chunk Attention training-free long-context — see docs/features/gemma4_dca.md ++ add_opt(common_arg( ++ {"--dca"}, "MODE", ++ "Dual Chunk Attention (training-free long-context extension): off (default), on. " ++ "Routes full-attention layers (Gemma-4 global / Qwen all) through 3-regime chunked-position " ++ "attention so no query-key distance exceeds the pretrain window. Compose with " ++ "--rope-scaling yarn + --override-kv .context_length for context beyond native.", ++ [](common_params & params, const std::string & value) { ++ if (value == "off" || value == "false" || value == "0") params.dca_enabled = false; ++ else if (value == "on" || value == "true" || value == "1") params.dca_enabled = true; ++ else throw std::invalid_argument("--dca must be on|off"); ++ } ++ ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DCA")); ++ add_opt(common_arg( ++ {"--dca-chunk-size"}, "N", ++ "DCA chunk length in tokens. 0 = auto (the model's original/pretrain context window). " ++ "Intra-chunk positions are native; inter-chunk distances are clamped within the window.", ++ [](common_params & params, const std::string & value) { ++ params.dca_chunk_size = (uint32_t) std::stoul(value); ++ } ++ ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DCA_CHUNK_SIZE")); ++ add_opt(common_arg( ++ {"--dca-yarn-factor"}, "F", ++ "DCA YaRN mscale stretch composed with chunked attention. 1.0 = none (default).", ++ [](common_params & params, const std::string & value) { ++ params.dca_yarn_factor = std::stof(value); ++ } ++ ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DCA_YARN_FACTOR")); + // 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 +@@ -1645,6 +1645,10 @@ struct llama_context_params common_context_params_to_llama(const common_params & + cparams.neo_pipeline_mode = params.neo_pipeline_mode; + // opencoti F5 W3 #292 fused-MoE — see docs/features/advanced_kv.md + cparams.fused_moe_up_gate = params.fused_moe_up_gate; ++ // opencoti F5 dca — see docs/features/gemma4_dca.md ++ cparams.dca_enabled = params.dca_enabled; ++ cparams.dca_chunk_size = params.dca_chunk_size; ++ cparams.dca_yarn_factor = params.dca_yarn_factor; + cparams.n_rs_seq = params.speculative.need_n_rs_seq(); + cparams.n_batch = params.n_batch; + cparams.n_ubatch = params.n_ubatch; +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 +@@ -459,6 +459,11 @@ struct common_params { + // false = off (default; unfused 3-node path). true = fuse decode MoE + // up+gate+GLU into one GGML_OP_MOE_FUSED_UP_GATE op for the CUDA backend. + bool fused_moe_up_gate = false; ++ // opencoti F5 dca — Dual Chunk Attention training-free long-context (docs/features/gemma4_dca.md) ++ // off by default; chunk 0 = auto (<- n_ctx_orig_yarn); yarn_factor 1.0 = no YaRN stretch. ++ bool dca_enabled = false; ++ uint32_t dca_chunk_size = 0; ++ float dca_yarn_factor = 1.0f; + // 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 / +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 +@@ -599,6 +599,8 @@ extern "C" { + + GGML_OP_TURBO_WHT, // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391 + ++ GGML_OP_FLASH_ATTN_EXT_LSE, // opencoti F5 dca — #593 (O+lse packed) ++ + GGML_OP_COUNT, + }; + +@@ -2487,6 +2489,64 @@ extern "C" { + float max_bias, + float logit_softcap); + ++ // opencoti F5 dca (#593) — flash-attention emitting NORMALIZED O packed with per-row ++ // LSE for the DCA 3-regime graph-math combine. dst = [DV+1, n_head, n_q, n_batch]; row ++ // [0,DV) = O, row [DV] = lse (= m + logf(sum exp), softcap+sinks folded in). Same ++ // (q,k,v,mask,scale,max_bias,logit_softcap) contract as ggml_flash_attn_ext. CUDA-only. ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): pos_q/chunk_size/regime select the analytical ++ // band-bounds path for the SUCC/INTER regimes — pass mask=NULL + the device I32 query-position ++ // vector pos_q (len q->ne[1]) + the DCA chunk_size + regime (1=SUCC, 2=INTER); the CUDA op derives ++ // the per-tile KV band on-device (no dense [n_kv x n_tokens] mask, no host O(n_kv*n_tokens) fill + ++ // H2D). INTRA (regime 0) keeps its native-causal mask: mask=, pos_q=NULL, chunk_size=0, ++ // regime=0 — the legacy form (src[4]=NULL, op_params[3..4]=0, byte-identical to pre-O6). ++ GGML_API struct ggml_tensor * ggml_flash_attn_ext_lse( ++ 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, ++ struct ggml_tensor * pos_q, ++ int chunk_size, ++ int regime); ++ ++ // opencoti F5 dca perf (T87.pD Branch B) — FUSED single-pass DCA flash-attention. Replaces the ++ // 3 per-regime ggml_flash_attn_ext_lse ops + the host LSE combine with ONE FA pass that walks the ++ // full causal range [0,i] in 3 contiguous key-phases (INTER→SUCC→INTRA), each using its OWN ++ // pre-roped Q variant, with a single continuous online-softmax. The 3 regime bands PARTITION the ++ // causal range so this is bit-faithful to the combine (no double counting). Encoded as a ++ // GGML_OP_FLASH_ATTN_EXT_LSE node with regime=3 (FUSED); the CUDA forward branches on it. ++ // q_intra/q_succ/q_inter : the 3 pre-roped+permuted Q variants (same shape). ++ // k, v : roped-once K, V (as for ggml_flash_attn_ext). ++ // mask_intra : the native-causal INTRA mask (its KV_min/KV_max already bound the ++ // INTRA phase to [cq*c, causal_limit] via the O4 band-trim). ++ // mask_succ/mask_inter : opencoti F5 dca (#617B) FRAGMENTED-fallback dense masks. NULL on the ++ // contiguous fast path (the always-case ⇒ kernel uses analytical index ++ // bands, byte-identical to #617). Non-NULL only when the KV cache is ++ // non-contiguous after a context-shift (cell index != absolute pos): ++ // the kernel then runs each phase over the FULL key range and selects ++ // keys via these position-based masks (mutually exclusive across the 3). ++ // pos_q_real : device I32 [n_q] REAL abs query positions; kernel derives cq=pos/chunk. ++ // Output dst = [DV+1, n_head, n_q, n_batch] (same as the LSE op; the [DV] lse row is UNUSED — the ++ // single fused pass needs no combine). max_bias must be 0 (no ALiBi on the analytical bands). ++ GGML_API struct ggml_tensor * ggml_flash_attn_ext_dca_fused( ++ struct ggml_context * ctx, ++ struct ggml_tensor * q_intra, ++ struct ggml_tensor * q_succ, ++ struct ggml_tensor * q_inter, ++ struct ggml_tensor * k, ++ struct ggml_tensor * v, ++ struct ggml_tensor * mask_intra, ++ struct ggml_tensor * mask_succ, ++ struct ggml_tensor * mask_inter, ++ struct ggml_tensor * pos_q_real, ++ float scale, ++ float max_bias, ++ float logit_softcap, ++ int chunk_size); ++ + 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 +@@ -1273,7 +1273,8 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra + // 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)) { ++ (node->op == GGML_OP_STREAMING_FLASH_ATTN || node->op == GGML_OP_MOE_FUSED_UP_GATE || ++ node->op == GGML_OP_FLASH_ATTN_EXT_LSE /* opencoti F5 dca #593 */)) { + *cur_backend_id = 0; + SET_CAUSE(node, "4.cuda-only"); + } +@@ -1638,6 +1639,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s + // when offloading MoE weights, we can reduce the amount of data copied by copying only the experts that are used + ggml_tensor * node = split->graph.nodes[0]; + if (split->graph.n_nodes > 0 && ++ input->buffer != NULL && // opencoti F5 dca: graph inputs may lack a buffer here; guard the usage query + ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(input->buffer) && ( + (node->src[0] == input_cpy && node->op == GGML_OP_MUL_MAT_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 +@@ -2110,6 +2110,7 @@ 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_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 (CUDA-only, shares abort) + case GGML_OP_STREAMING_FLASH_ATTN: + { + // opencoti-hook: f5-rolling-kv — W4 M7-D #304. CUDA-only op. The +@@ -2402,6 +2403,7 @@ 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_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 (CUDA-only) + case GGML_OP_STREAMING_FLASH_ATTN: + { + // opencoti-hook: f5-rolling-kv — W4 M7-D #304. CUDA-only; never +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 +@@ -476,6 +476,7 @@ static bool GGML_CALL ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev + // 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_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 (CUDA-only) + case GGML_OP_STREAMING_FLASH_ATTN: + case GGML_OP_MOE_FUSED_UP_GATE: + return false; +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 +@@ -25,6 +25,7 @@ typedef void (* fattn_kernel_t)( + const char * __restrict__ mask, + const char * __restrict__ sinks, + const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + float * __restrict__ dst, + float2 * __restrict__ dst_meta, + const float scale, +@@ -47,6 +48,16 @@ typedef void (* fattn_kernel_t)( + // 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; ++// opencoti F5 dca perf (T87.pD-opt O2): launch_fattn sets this true iff a finalize kernel ran (so ++// dst_lse was populated), false on the in-kernel single-block path. Lets the DCA _lse op fall back to ++// streaming_lse_kernel only when the free emit didn't happen. Defined in fattn.cu. ++extern thread_local bool opencoti_fattn_dst_lse_written; ++// opencoti F5 dca perf (T87.pD-opt O6 Tier B): analytical band-bounds side-channel for DCA SUCC/INTER. ++// When pos_q != nullptr, launch_fattn fills KV_min/KV_max from pos_q[0]/chunk for the regime (no dense ++// mask scan; mask passed as nullptr) and drains these back. Defined in fattn.cu. See UPSTREAM_SYNC.md. ++extern thread_local const int * opencoti_fattn_dca_pos_q; ++extern thread_local int opencoti_fattn_dca_chunk; ++extern thread_local int opencoti_fattn_dca_regime; + + 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); +@@ -645,10 +656,17 @@ constexpr __device__ dequantize_V_t get_dequantize_V() { + } + } + ++// opencoti F5 dca perf (T87.pD-opt O4): this kernel emits BOTH the high edge (KV_max, the suffix-trim ++// upstream uses to skip fully-masked HIGH key-bands) AND a low edge (KV_min) so the FA kernels can also ++// skip fully-masked LOW key-bands. For ordinary causal masks KV_min == 0 (no-op, byte-identical); for ++// Dual Chunk Attention INTRA/SUCC regimes (valid band [lo,hi) starting above 0) and for sliding-window ++// masks, KV_min lets each query-tile attend only its contiguous band. KV_min may be nullptr (callers ++// that don't want the low trim), in which case only KV_max is written. + template + __launch_bounds__(FATTN_KQ_STRIDE/2, 1) + static __global__ void flash_attn_mask_to_KV_max( +- const half2 * __restrict__ mask, int * __restrict__ KV_max, const int ne30, const int s31, const int s33) { ++ const half2 * __restrict__ mask, int * __restrict__ KV_max, int * __restrict__ KV_min, ++ const int ne30, const int s31, const int s33) { + const int ne31 = gridDim.x; + const int tid = threadIdx.x; + const int sequence = blockIdx.y; +@@ -692,11 +710,79 @@ static __global__ void flash_attn_mask_to_KV_max( + // In either case, walk back the decrementation by FATTN_KQ_STRIDE. + KV_max_sj += FATTN_KQ_STRIDE; + ++ // opencoti F5 dca perf (T87.pD-opt O4): symmetric LOW-edge scan. Walk UP from the bottom tile and ++ // stop at the first tile that is not entirely -inf; its lower edge is the lowest KV position any ++ // query in this tile attends to. All threads must participate (warp_reduce + buf_iw) so this runs ++ // before the thread-0-only return below. Clamp to KV_max so a fully-masked row yields min==max (an ++ // empty [min,max) band → the FA kernel does no work, matching the existing KV_max==0 behaviour). ++ int KV_min_sj = 0; ++ if (KV_min != nullptr) { ++ for (; KV_min_sj < ne30 * FATTN_KQ_STRIDE; KV_min_sj += FATTN_KQ_STRIDE) { ++ int all_inf = 1; ++ ++#pragma unroll ++ for (int j = 0; j < ncols1; ++j) { ++ const float2 tmp = __half22float2(mask[j*s31 + KV_min_sj/2 + tid]); ++ all_inf = all_inf && int(isinf(tmp.x)) && int(isinf(tmp.y)); ++ } ++ ++ all_inf = warp_reduce_all(all_inf); ++ if (tid % WARP_SIZE == 0) { ++ buf_iw[tid / WARP_SIZE] = all_inf; ++ } ++ __syncthreads(); ++ all_inf = buf_iw[tid % WARP_SIZE]; ++ __syncthreads(); ++ all_inf = warp_reduce_all(all_inf); ++ ++ if (!all_inf) { ++ break; ++ } ++ } ++ if (KV_min_sj > KV_max_sj) { ++ KV_min_sj = KV_max_sj; ++ } ++ } ++ + if (threadIdx.x != 0) { + return; + } + + KV_max[sequence*ne31 + jt] = KV_max_sj; ++ if (KV_min != nullptr) { ++ KV_min[sequence*ne31 + jt] = KV_min_sj; ++ } ++} ++ ++// opencoti F5 dca perf (T87.pD-opt O6 Tier B): ANALYTICAL per-tile KV band producer for the DCA ++// SUCC/INTER regimes. Replaces the dense-mask scan (flash_attn_mask_to_KV_max) — there is no mask. The ++// SUCC/INTER band for query position p1 is a CONTIGUOUS causal key range fixed by cq = p1/chunk, and ++// (guard: chunk % n_ubatch == 0) cq is uniform across the whole ubatch, so a single [lo,hi) covers every ++// query tile. Key index == position (single-stream contiguous prefill), and chunk is a multiple of ++// FATTN_KQ_STRIDE so the edges are tile-aligned (exact). One thread per query tile jt writes that band; ++// the FA kernels clamp their key-tile loop to [KV_min, KV_max) (the O4 kb0_min / kb0_stop machinery). ++// SUCC (regime 1): [ (cq-1)*chunk , cq*chunk ) INTER (regime 2): [ 0 , (cq-1)*chunk ) ++// An empty band (e.g. INTER at cq==1 → hi==0 → lo==hi) ⇒ the FA does zero key-tiles ⇒ LSE = -inf, which ++// the host 3-regime combine weights as 0 (identical to the all-(-inf)-mask-row behaviour). n_stream==1 ++// for DCA so the flat thread index is exactly the (sequence==0) jt the kernels read as KV_*[0*iter_j+jt]. ++static __global__ void flash_attn_dca_bounds_to_KV( ++ const int * __restrict__ pos_q, int * __restrict__ KV_max, int * __restrict__ KV_min, ++ const int n_tiles, const int chunk, const int regime, const int n_kv_keys) { ++ const int jt = blockIdx.x * blockDim.x + threadIdx.x; ++ if (jt >= n_tiles) { ++ return; ++ } ++ const int p1 = pos_q[0]; // cq uniform across the ubatch (guard: chunk % n_ubatch == 0) ++ const int cq = p1 / chunk; ++ int lo, hi; ++ if (regime == 1) { lo = (cq - 1) * chunk; hi = cq * chunk; } // SUCC ++ else { lo = 0; hi = (cq - 1) * chunk; } // INTER (regime == 2) ++ if (lo < 0) { lo = 0; } ++ if (hi < 0) { hi = 0; } ++ if (hi > n_kv_keys) { hi = n_kv_keys; } ++ if (lo > hi) { lo = hi; } ++ KV_max[jt] = hi; ++ KV_min[jt] = lo; + } + + template // D == head size +@@ -1006,6 +1092,7 @@ void launch_fattn( + ggml_cuda_pool_alloc K_f16(pool); + ggml_cuda_pool_alloc V_f16(pool); + ggml_cuda_pool_alloc KV_max(pool); ++ ggml_cuda_pool_alloc KV_min(pool); // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + ggml_cuda_pool_alloc dst_tmp(pool); + ggml_cuda_pool_alloc dst_tmp_meta(pool); + +@@ -1100,8 +1187,24 @@ void launch_fattn( + const int iter_k = K->ne[1] / FATTN_KQ_STRIDE; + + KV_max.alloc(ne_KV_max); ++ KV_min.alloc(ne_KV_max); // opencoti F5 dca perf (T87.pD-opt O4): low-edge band, same shape as KV_max + flash_attn_mask_to_KV_max<<>> +- ((const half2 *) mask->data, KV_max.ptr, iter_k, s31, s33); ++ ((const half2 *) mask->data, KV_max.ptr, KV_min.ptr, iter_k, s31, s33); ++ CUDA_CHECK(cudaGetLastError()); ++ } else if (opencoti_fattn_dca_pos_q != nullptr && K->ne[1] % FATTN_KQ_STRIDE == 0) { ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): analytical band for a DCA SUCC/INTER regime — there ++ // is no mask to scan (mask==nullptr), so the mask-scan branch above never fires for these. Fill ++ // KV_min/KV_max from the query positions for ANY n_q (incl. decode n_q==1), so the band is always ++ // enforced. n_stream==1 for DCA ⇒ ne_KV_max == ntiles_x and the flat jt index is correct. ++ const dim3 blocks_num_KV_max(ntiles_x, Q->ne[3], 1); ++ const int ne_KV_max = blocks_num_KV_max.x * blocks_num_KV_max.y; ++ KV_max.alloc(ne_KV_max); ++ KV_min.alloc(ne_KV_max); ++ const int threads = 64; ++ const int blocks = (ne_KV_max + threads - 1) / threads; ++ flash_attn_dca_bounds_to_KV<<>>( ++ opencoti_fattn_dca_pos_q, KV_max.ptr, KV_min.ptr, ++ ne_KV_max, opencoti_fattn_dca_chunk, opencoti_fattn_dca_regime, (int) K->ne[1]); + CUDA_CHECK(cudaGetLastError()); + } + +@@ -1212,6 +1315,7 @@ void launch_fattn( + mask ? ((const char *) mask->data) : nullptr, + sinks ? ((const char *) sinks->data) : nullptr, + KV_max.ptr, ++ KV_min.ptr, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band (nullptr when prepass didn't run) + !stream_k && parallel_blocks > 1 ? dst_tmp.ptr : (float *) KQV->data, dst_tmp_meta.ptr, + scale, max_bias, m0, m1, n_head_log2, logit_softcap, + Q->ne[0], ne01, Q->ne[2], Q->ne[3], Q->nb[1], Q->nb[2], Q->nb[3], +@@ -1228,10 +1332,14 @@ void launch_fattn( + // 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. ++ // opencoti F5 dca perf (T87.pD-opt O2): report (don't assert) whether a finalize kernel will run and ++ // thus whether dst_lse was populated. Callers that have a fallback (the DCA _lse op -> streaming_lse_ ++ // kernel) read opencoti_fattn_dst_lse_written and recompute only when false. The rolling-kv decode ++ // path always has stream_k engaged (fixup_ran true), so its behaviour is unchanged. + 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"); ++ opencoti_fattn_dst_lse_written = (fixup_ran || combine_ran); + } + + if (stream_k) { +diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh +--- a/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh ++++ b/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh +@@ -527,9 +527,11 @@ static __device__ __forceinline__ void flash_attn_ext_f16_load_mask( + } + } + ++ + template ++ typename T_A_KQ, typename T_B_KQ, typename T_C_KQ, typename T_A_VKQ, typename T_B_VKQ, typename T_C_VKQ, ++ bool dca_fused = false> + static __device__ __forceinline__ void flash_attn_ext_f16_iter( + const float2 * const __restrict__ Q_f2, + const half2 * const __restrict__ K_h2, +@@ -565,8 +567,20 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + constexpr int nbatch_fa = ggml_cuda_fattn_mma_get_nbatch_fa(DKQ, DV, ncols); + constexpr int nbatch_K2 = ggml_cuda_fattn_mma_get_nbatch_K2(DKQ, DV, ncols); + constexpr int nbatch_V2 = ggml_cuda_fattn_mma_get_nbatch_V2(DKQ, DV, ncols); +- constexpr bool Q_in_reg = ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols); +- constexpr int nstages = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2); ++ // opencoti F5 dca (#623): MUST mirror flash_attn_ext_f16_process_tile's forced Q_in_reg=false on the ++ // fused path. The config returns true for DKQ 128/256; if iter trusted it, iter would read Q from the ++ // Q_B register array (the Q_in_reg branch) — but Q_B is only populated under `if constexpr (!dca_fused)`, ++ // so on the DCA path it is uninitialised garbage -> NaN KQ -> '!' loop. Forcing false makes iter load Q ++ // from tile_Q (filled by dca_load_tile_Q), matching process_tile's layout. DKQ=512 already returns false. ++ constexpr bool Q_in_reg = dca_fused ? false : ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols); ++ // opencoti F5 dca (#623): the dca_fused 3-phase path calls iter with last_iter=false and NO caller-side ++ // cp.async prologue, so iter must be self-contained per kb0 (load its own tile_K/mask/V). That only holds ++ // at nstages<=1. For narrow head dims (DKQ 128/256) the Ampere config returns nstages_target=2 (double- ++ // buffered: the caller is expected to prime tile_K and iter SKIPS the K-load at the `nstages<=1` guard ++ // below) -> tile_K read uninitialised -> NaN KQ -> garbage ('!' loop). DKQ=512 already gets nstages=1 ++ // (why Gemma's fused path works). Clamp the fused path to single-stage so 128/256 mirror the validated 512. ++ constexpr int nstages_cfg = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2); ++ constexpr int nstages = dca_fused && nstages_cfg > 1 ? 1 : nstages_cfg; + + constexpr int stride_tile_Q = DKQ/2 + 4; + constexpr int stride_tile_K = nbatch_K2 + 4; +@@ -593,7 +607,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + (V_h2 + int64_t(k_VKQ_0)*stride_V, tile_V, nbatch_V2, stride_V, k_VKQ_sup); + } else { + constexpr bool use_cp_async = nstages == 1; +- if (ncols2 > 1 || mask_h) { ++ if (mask_h) { // opencoti O6-B: null mask_h ⇒ band-only (KV_min/KV_max); stock ncols2>1 always has a mask + flash_attn_ext_f16_load_mask + (mask_h + k_VKQ_0, tile_mask, stride_mask, k_VKQ_sup, jt*ncols1, ne01); + } +@@ -692,7 +706,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + float KQ_rowsum_add[cols_per_thread] = {0.0f}; + + if constexpr (cols_per_warp == 8) { +- if (ncols2 > 1 || mask_h) { ++ if (mask_h) { // opencoti O6-B: null mask_h ⇒ band-only (KV_min/KV_max); stock ncols2>1 always has a mask + #pragma unroll + for (int i00 = 0; i00 < nbatch_fa; i00 += np*T_C_KQ::I) { + const int i0 = i00 + (threadIdx.y % np)*T_C_KQ::I; +@@ -754,7 +768,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + } + } + } else { // not Turing mma or T_B_KQ::I > 8 +- if (ncols2 > 1 || mask_h) { ++ if (mask_h) { // opencoti O6-B: null mask_h ⇒ band-only (KV_min/KV_max); stock ncols2>1 always has a mask + #pragma unroll + for (int i00 = 0; i00 < nbatch_fa; i00 += np*T_C_KQ::J) { + const int i0 = i00 + (threadIdx.y % np)*T_C_KQ::J; +@@ -916,6 +930,8 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + } + } + #endif // defined(TURING_MMA_AVAILABLE) ++ // opencoti DCA #635 probe B REMOVED — see the note where the debug globals used to live: decode is ++ // stream-k single-phase-per-block, so this in-kernel rescale point is not the SUCC->INTRA seam. + } + + // Convert KQ C tiles into B tiles for VKQ calculation: +@@ -939,7 +955,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( + cp_async_wait_all(); + __syncthreads(); + if (!last_iter) { +- if (ncols2 > 1 || mask_h) { ++ if (mask_h) { // opencoti O6-B: null mask_h ⇒ band-only (KV_min/KV_max); stock ncols2>1 always has a mask + flash_attn_ext_f16_load_mask + (mask_h + k_VKQ_0 + nbatch_fa, tile_mask, stride_mask, k_VKQ_sup, jt*ncols1, ne01); + } +@@ -1113,7 +1129,54 @@ template struct mma_tile_sizes { + }; + #endif // defined(TURING_MMA_AVAILABLE) + +-template ++// opencoti F5 dca perf (T87.pD Branch B): fill tile_Q SRAM from a Q source (one regime's pre-roped Q). ++// Extracted verbatim from process_tile's Q-load (the !Q_in_reg SRAM fill). The fused kernel calls this ++// once per phase to swap the live Q variant (Q_in_reg==false for D=512 ⇒ Q lives in tile_Q, not regs). ++// Caller controls __syncthreads (before: prior phase's tile_Q reads done; after: fill visible). ++template ++static __device__ __forceinline__ void dca_load_tile_Q( ++ half2 * const __restrict__ tile_Q, const float2 * const __restrict__ Q_f2, const float scale, ++ const int jt, const int zt_gqa, const uint3 ne01, const int gqa_ratio, ++ const int stride_Q1, const int stride_Q2) { ++ constexpr int warp_size = ggml_cuda_get_physical_warp_size(); ++ constexpr int ncols = ncols1 * ncols2; ++ constexpr int stride_tile_Q = DKQ/2 + 4; ++ const half2 scale_h2 = make_half2(scale, scale); ++#pragma unroll ++ for (int stride_k : {warp_size, warp_size/2, warp_size/4, warp_size/8}) { ++ const int k0_start = stride_k == warp_size ? 0 : DKQ/2 - (DKQ/2) % (2*stride_k); ++ const int k0_stop = DKQ/2 - (DKQ/2) % (1*stride_k); ++ const int stride_jc = warp_size / stride_k; ++ if (k0_start == k0_stop) { ++ continue; ++ } ++#pragma unroll ++ for (int jc0 = 0; jc0 < ncols; jc0 += nwarps*stride_jc) { ++ const int jc = jc0 + threadIdx.y*stride_jc + (stride_k == warp_size ? 0 : threadIdx.x / stride_k); ++ if (jc0 + nwarps*stride_jc > ncols && jc >= ncols) { ++ break; ++ } ++ const int j = jc / ncols2; ++ const int c = jc % ncols2; ++ if ((ncols1 == 1 || jt*ncols1 + j < int(ne01.z)) && (ncols2 == 1 || zt_gqa*ncols2 + c < gqa_ratio)) { ++#pragma unroll ++ for (int k0 = k0_start; k0 < k0_stop; k0 += stride_k) { ++ const int k = k0 + (stride_k == warp_size ? threadIdx.x : threadIdx.x % stride_k); ++ const float2 tmp = Q_f2[(jt*ncols1 + j)*stride_Q1 + c*stride_Q2 + k]; ++ tile_Q[jc*stride_tile_Q + k] = scale_h2 * make_half2(tmp.x, tmp.y); ++ } ++ } else { ++#pragma unroll ++ for (int k0 = k0_start; k0 < k0_stop; k0 += stride_k) { ++ const int k = k0 + (stride_k == warp_size ? threadIdx.x : threadIdx.x % stride_k); ++ tile_Q[jc*stride_tile_Q + k] = make_half2(0.0f, 0.0f); ++ } ++ } ++ } ++ } ++} ++ ++template + static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + const float2 * const __restrict__ Q_f2, + const half2 * const __restrict__ K_h2, +@@ -1137,7 +1200,18 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + const int jt, + const int zt_gqa, + const int kb0_start, +- const int kb0_stop) { ++ const int kb0_stop, ++ const int kb0_min, // opencoti F5 dca perf (T87.pD-opt O4): lowest non-masked k-block; loop starts here ++ const float2 * const __restrict__ Q_succ_f2 = nullptr, // opencoti F5 dca perf (Branch B fused): SUCC-phase pre-roped Q ++ const float2 * const __restrict__ Q_inter_f2 = nullptr, // INTER-phase pre-roped Q ++ const int dca_cq = 0, // chunk index cq (uniform per launch under analytical bands) ++ const int dca_c_blk = 0, // chunk_size / nbatch_fa (chunk width in k-blocks) ++ // opencoti F5 dca (#617B): FRAGMENTED-fallback dense masks. NULL ⇒ contiguous fast path (the ++ // analytical b1/b2 band partition below). Non-NULL (post context-shift, index!=pos) ⇒ each of ++ // the 3 phases runs the FULL owned [kb0_start,kb0_stop) and selects keys via its position-based ++ // mask (mutually exclusive across the 3 ⇒ each causal key contributes in exactly one phase). ++ const half * const __restrict__ mask_succ_h = nullptr, ++ const half * const __restrict__ mask_inter_h = nullptr) { + #if defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE) + //In this kernel Q, K, V are matrices while i, j, k are matrix indices. + +@@ -1157,8 +1231,19 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + constexpr int nbatch_K2 = ggml_cuda_fattn_mma_get_nbatch_K2 (DKQ, DV, ncols); + constexpr int nbatch_V2 = ggml_cuda_fattn_mma_get_nbatch_V2 (DKQ, DV, ncols); + constexpr int nbatch_combine = ggml_cuda_fattn_mma_get_nbatch_combine(DKQ, DV, ncols); +- constexpr bool Q_in_reg = ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols); +- constexpr int nstages = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2); ++ // opencoti F5 dca (#623): the fused-DCA path swaps the 3 pre-roped Q variants by RELOADING tile_Q at ++ // each INTER/SUCC/INTRA phase boundary (see the dca_fused block below + dca_load_tile_Q), so the per- ++ // key iter MUST read Q from tile_Q — i.e. Q_in_reg==false. The config only returns false for DKQ=512 ++ // (Gemma 4); for DKQ=128/256 (qwen3 / qwen35moe) it returns true, which (a) routes the iter to the ++ // never-filled Q_B register array and (b) aliases tile_K onto tile_Q (line below), so the K-load ++ // clobbers the Q the phase loop just wrote — both corrupt the attention (observed: multi-chunk ++ // '!!!!'/garbage at hd128). Force false on the fused path at every head dim; no-op at DKQ=512. ++ constexpr bool Q_in_reg = dca_fused ? false : ggml_cuda_fattn_mma_get_Q_in_reg(DKQ, DV, ncols); ++ // opencoti F5 dca (#623): clamp the fused path to single-stage (see flash_attn_ext_f16_iter for the full ++ // rationale). nstages drives the tile_V/tile_mask shared-mem offsets just below (~:1126); the kernel, ++ // iter, and the host ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case smem sizing must all use this clamp. ++ constexpr int nstages_cfg = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2); ++ constexpr int nstages = dca_fused && nstages_cfg > 1 ? 1 : nstages_cfg; + + if (cols_per_warp > ncols) { + NO_DEVICE_CODE; +@@ -1196,6 +1281,135 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + KQ_max[col] = -FLT_MAX/2.0f; + } + ++ // opencoti F5 dca perf (T87.pD Branch B): single fused pass over the partitioned causal range. ++ // Sorted-key-axis phases: INTER [0,b1) Q_inter no-mask | SUCC [b1,b2) Q_succ no-mask | INTRA ++ // [b2,kb0_stop) Q_intra(=Q_f2)+causal mask. One continuous online-softmax (VKQ_C/KQ_max/KQ_rowsum ++ // persist across phases) ⇒ bit-faithful to the 3-op + host LSE combine (the bands partition [0,i]). ++ // Q swapped by reloading tile_Q at each boundary (Q_in_reg==false ⇒ iter reads Q from tile_Q). ++ // last_iter is dead at nstages<=1 (used only under nstages>1) ⇒ pass false everywhere; per-block ++ // k_VKQ_sup: full nbatch_fa for the whole-block INTER/SUCC bands, real ne11-bound for INTRA's tail. ++ if constexpr (dca_fused) { ++ if (mask_inter_h == nullptr) { ++ // ===== CONTIGUOUS FAST PATH (cell index == absolute position) — UNCHANGED from #617/#618 ===== ++ // Stream-k: this CUDA block owns key-blocks [kb0_start, kb0_stop) of the output tile. The three DCA ++ // phases partition the causal range — INTER [0,b1) | SUCC [b1,b2) | INTRA [b2,...] — so clip each to ++ // the owned sub-range and run only the non-empty intersection. kb0_start/kb0_stop/dca_cq are ++ // block-uniform (one tile per block, KV_max per-tile) ⇒ every thread takes the same phase guards, ++ // keeping the __syncthreads balanced. Partial (O, KQ_max, KQ_rowsum) over any contiguous k-range ++ // combine via the shared stream-k fixup regardless of which Q-phase produced them — the per-key ++ // pre-roped Q is baked into each partial, so the fixup is phase-agnostic. ++ int b1 = (dca_cq - 1) * dca_c_blk; if (b1 < 0) { b1 = 0; } ++ int b2 = dca_cq * dca_c_blk; if (b2 < 0) { b2 = 0; } ++ const int inter_lo = kb0_start; const int inter_hi = min(b1, kb0_stop); ++ const int succ_lo = max(kb0_start, b1); const int succ_hi = min(b2, kb0_stop); ++ const int intra_lo = max(kb0_start, b2); const int intra_hi = kb0_stop; ++ // Phase order is INTRA -> SUCC -> INTER (near keys FIRST). Online-softmax is order-invariant in ++ // exact math, so within ONE block this ordering is purely cosmetic; running the highest-logit ++ // current-chunk band first just keeps the running KQ_max high so later bands rescale cleanly. (The ++ // earlier "INTRA-first avoids an in-iter FTZ flush" rationale was FALSIFIED in #635: the real ++ // SOFTMAX_FTZ wipeout was in the CROSS-BLOCK stream-k fixup, not flash_attn_ext_f16_iter. The fix ++ // is to run DCA fused DECODE single-block-per-tile so there is no cross-block fixup at all — see the ++ // nblocks override in ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case. Each block still owns one ++ // contiguous [kb0_start,kb0_stop) band; the 3 sub-bands are just visited high-to-low.) ++ if (intra_hi > intra_lo) { // PHASE INTRA (same-chunk causal; mask supplies the within-last-block tail) ++ // Mirror the proven stock ncols2>1 path EXACTLY: oob_check=false, k_VKQ_sup=nbatch_fa. The ++ // INTRA mask (full n_kv width, -inf beyond the same-chunk causal diagonal) zeroes every key ++ // past the causal bound i — including the partial tail of the last block — just as it does ++ // for every non-DCA Gemma global layer. kb0_stop (from KV_max) already trims to the bound; ++ // the INTRA band starts at the exact block boundary b2 (chunk % nbatch_fa == 0). ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = intra_lo; kb0 < intra_hi; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_f2, K_h2, V_h2, mask_h, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ } ++ if (succ_hi > succ_lo) { // PHASE SUCC (immediately-prior chunk, Q_succ, band-only — no mask) ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_succ_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = succ_lo; kb0 < succ_hi; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_succ_f2, K_h2, V_h2, nullptr, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ } ++ if (inter_hi > inter_lo) { // PHASE INTER (older chunks, Q_inter, band-only — no mask) ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_inter_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = inter_lo; kb0 < inter_hi; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_inter_f2, K_h2, V_h2, nullptr, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ } ++ } else { ++ // ===== FRAGMENTED FALLBACK (opencoti F5 dca #617B) — cell index != absolute position ===== ++ // After a context-shift the cache is position-relabeled in place, so the analytical index bands ++ // above pick the WRONG physical cells. Instead run ALL THREE phases over the FULL owned range ++ // [kb0_start,kb0_stop); each phase applies its dense position-based mask (mask_inter/succ/intra) ++ // which the host filled mutually-exclusively (INTRA ck==cq, SUCC ck==cq-1, INTER ck SUCC -> INTER (near-keys-first) order as the contiguous fast path above. The dense ++ // masks are mutually exclusive so each causal key contributes in exactly one phase; within ONE block ++ // the online-softmax is order-invariant, so the ordering is cosmetic (see the falsified-FTZ note on ++ // the fast path). Decode also runs single-block-per-tile here (#635 nblocks override) so there is no ++ // cross-block fixup on the shifted-cache path either. ++ // PHASE INTRA (Q_intra=Q_f2, dense mask_intra=mask_h) ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = kb0_start; kb0 < kb0_stop; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_f2, K_h2, V_h2, mask_h, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ // PHASE SUCC (Q_succ, dense mask_succ) ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_succ_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = kb0_start; kb0 < kb0_stop; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_succ_f2, K_h2, V_h2, mask_succ_h, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ // PHASE INTER (Q_inter, dense mask_inter) ++ __syncthreads(); ++ dca_load_tile_Q(tile_Q, Q_inter_f2, scale, jt, zt_gqa, ne01, gqa_ratio, stride_Q1, stride_Q2); ++ __syncthreads(); ++ for (int kb0 = kb0_start; kb0 < kb0_stop; ++kb0) { ++ flash_attn_ext_f16_iter ++ ++ (Q_inter_f2, K_h2, V_h2, mask_inter_h, dstk, dstk_fixup, scale, slope, logit_softcap, ++ ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, ++ KQ_max, KQ_rowsum, jt, kb0, nbatch_fa); ++ } ++ } ++ } ++ ++ if constexpr (!dca_fused) { + // Load Q data into tile_Q, either temporarily or permanently. + // Q in registers is faster, but register pressure is the biggest bottleneck. + // The loading is done with decreasing granularity for D for better memory bandwidth. +@@ -1253,7 +1467,18 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + + __syncthreads(); + ++ // opencoti F5 dca perf (T87.pD-opt O4): start the K/V loop at the lowest non-masked k-block so ++ // fully-masked LOW key-bands (DCA INTRA/SUCC regimes, SWA windows) are skipped entirely. Clamp to ++ // kb0_stop-1 so the unconditional last iteration below stays in range; processing one extra masked ++ // block is the online-softmax identity (contributes nothing), and the cp_async preload below keys ++ // off this same kb0 so the multi-stage pipeline stays consistent. kb0_min==0 for causal masks (no-op). + int kb0 = kb0_start; ++ if (kb0_min > kb0) { ++ kb0 = kb0_min; ++ } ++ if (kb0_stop > 0 && kb0 > kb0_stop - 1) { ++ kb0 = kb0_stop - 1; ++ } + + // Preload mask and K data for first iteration when using cp_async with multiple stages: + if constexpr (nstages > 1) { +@@ -1261,7 +1486,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + constexpr bool use_cp_async = true; + constexpr bool oob_check = false; + constexpr int k_VKQ_sup = nbatch_fa; +- if (ncols2 > 1 || mask_h) { ++ if (mask_h) { // opencoti O6-B: null mask_h ⇒ band-only (KV_min/KV_max); stock ncols2>1 always has a mask + flash_attn_ext_f16_load_mask + (mask_h + kb0*nbatch_fa, tile_mask, stride_mask, k_VKQ_sup, jt*ncols1, ne01); + } +@@ -1311,6 +1536,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + ne01, ne02, stride_K, stride_V, stride_mask, tile_Q, tile_K, tile_V, tile_mask, Q_B, VKQ_C, + KQ_max, KQ_rowsum, jt, kb0, k_VKQ_sup); + } ++ } // end if constexpr (!dca_fused) + + // With multi-stage loading there is no __syncthreads at the end of the iter, + // there can be a race condition on shared memory access for combining/writing back results. +@@ -1695,7 +1921,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( + GGML_UNUSED_VARS(Q_f2, K_h2, V_h2, mask_h, sinks_f, dstk, dstk_fixup, + scale, slope, logit_softcap, ne01, ne02, gqa_ratio, + stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, +- jt, kb0_start, kb0_stop); ++ jt, kb0_start, kb0_stop, kb0_min); + NO_DEVICE_CODE; + #endif // defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE) + } +@@ -1709,6 +1935,7 @@ static __global__ void flash_attn_ext_f16( + const char * __restrict__ mask, + const char * __restrict__ sinks, + const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + float * __restrict__ dst, + float2 * __restrict__ dst_meta, + const float scale, +@@ -1806,7 +2033,7 @@ static __global__ void flash_attn_ext_f16( + + const float2 * Q_f2 = (const float2 *) (Q + nb03*sequence + nb02*zt_Q); + const half2 * K_h2 = (const half2 *) (K + nb13*sequence + nb12*z_KV); +- const half * mask_h = ncols2 == 1 && !mask ? nullptr : ++ const half * mask_h = !mask ? nullptr : + (const half *) (mask + nb33*(sequence % ne33)); + float2 * dstk = ((float2 *) dst) + (sequence*ne01.z*ne02 + zt_Q) * (DV/2); + +@@ -1818,17 +2045,20 @@ static __global__ void flash_attn_ext_f16( + if (KV_max) { + kb0_stop = min(kb0_stop, KV_max[sequence*iter_j + jt] / nbatch_fa); + } ++ // opencoti F5 dca perf (T87.pD-opt O4): lowest non-masked k-block for this query-tile (0 when no ++ // low trim, i.e. causal masks → unchanged). process_tile starts its K/V loop here. ++ const int kb0_min = KV_min ? (KV_min[sequence*iter_j + jt] / nbatch_fa) : 0; + constexpr bool is_fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer. + if (kb0_start == 0) { + constexpr bool needs_fixup = false; // CUDA block is working on an entire tile. + flash_attn_ext_f16_process_tile + (Q_f2, K_h2, V_h2, mask_h, sinks_f, dstk, dst_meta, scale, slope, logit_softcap, +- ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop); ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop, kb0_min); + } else { + constexpr bool needs_fixup = true; // CUDA block is missing the beginning of a tile. + flash_attn_ext_f16_process_tile + (Q_f2, K_h2, V_h2, mask_h, sinks_f, dstk, dst_meta, scale, slope, logit_softcap, +- ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop); ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop, kb0_min); + } + + kbc += iter_k; +@@ -1852,7 +2082,7 @@ static __global__ void flash_attn_ext_f16( + + const float2 * Q_f2 = (const float2 *) (Q + nb03*sequence + nb02*zt_Q); + const half2 * K_h2 = (const half2 *) (K + nb13*sequence + nb12*z_KV); +- const half * mask_h = ncols2 == 1 && !mask ? nullptr : ++ const half * mask_h = !mask ? nullptr : + (const half *) (mask + nb33*(sequence % ne33)); + float2 * dstk = ((float2 *) dst) + (sequence*ne01.z*ne02 + zt_Q) * (DV/2); + +@@ -1864,14 +2094,16 @@ static __global__ void flash_attn_ext_f16( + if (KV_max) { + kb0_stop = min(kb0_stop, KV_max[sequence*iter_j + jt] / nbatch_fa); + } ++ // opencoti F5 dca perf (T87.pD-opt O4): lowest non-masked k-block for this query-tile (see above). ++ const int kb0_min = KV_min ? (KV_min[sequence*iter_j + jt] / nbatch_fa) : 0; + + constexpr bool is_fixup = true; // Last index writes its data to fixup buffer to avoid data races with other blocks. + constexpr bool needs_fixup = false; + flash_attn_ext_f16_process_tile + (Q_f2, K_h2, V_h2, mask_h, sinks_f, dstk, dst_meta, scale, slope, logit_softcap, +- ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop); ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop, kb0_min); + #else +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +@@ -1956,6 +2188,355 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml + } + + ++// opencoti F5 dca perf (T87.pD Branch B + stream-k): FUSED single-pass DCA kernel. Uses the stock kbc ++// stream-k decomposition (load-balances the causal triangle ⇒ dst_meta partials + host fixup), walking ++// the partitioned causal range in INTER→SUCC→INTRA phases via process_tile<...,dca_fused=true> which ++// clips each phase to the block's owned [kb0_start,kb0_stop). Extra params vs the stock kernel: Q_succ/ ++// Q_inter (pre-roped Q for the older-chunk phases), pos_q (I32 real abs query pos), chunk (tokens). ++template ++__launch_bounds__(ggml_cuda_fattn_mma_get_nthreads(DKQ, DV, ncols1*ncols2), ggml_cuda_fattn_mma_get_occupancy(DKQ, DV, ncols1*ncols2)) ++static __global__ void flash_attn_ext_f16_dca_fused( ++ const char * __restrict__ Q, ++ const char * __restrict__ K, ++ const char * __restrict__ V, ++ const char * __restrict__ mask, ++ const char * __restrict__ mask_succ, // opencoti F5 dca (#617B): dense SUCC mask (NULL ⇒ analytical fast path) ++ const char * __restrict__ mask_inter, // opencoti F5 dca (#617B): dense INTER mask (NULL ⇒ analytical fast path) ++ const char * __restrict__ sinks, // unused (DCA has no attention sinks) ++ const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // unused in fused (per-phase bounds derived analytically) ++ const char * __restrict__ Q_succ, ++ const char * __restrict__ Q_inter, ++ const int * __restrict__ pos_q, ++ const int chunk, ++ float * __restrict__ dst, ++ float2 * __restrict__ dst_meta, // unused (single block per tile) ++ const float scale, ++ const float max_bias, ++ const float m0, ++ const float m1, ++ const uint32_t n_head_log2, ++ const float logit_softcap, ++ const int32_t ne00, const uint3 ne01, const int32_t ne02, const int32_t ne03, ++ const int32_t nb01, const int32_t nb02, const int32_t nb03, ++ const int32_t ne10, const int32_t ne11, const int32_t ne12, const int32_t ne13, ++ const int32_t nb11, const int32_t nb12, const int64_t nb13, ++ const int32_t nb21, const int32_t nb22, const int64_t nb23, ++ const int32_t ne31, const int32_t ne32, const int32_t ne33, ++ const int32_t nb31, const int32_t nb32, const int64_t nb33) { ++#if defined(FLASH_ATTN_AVAILABLE) && (defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)) || defined(AMD_MFMA_AVAILABLE)) ++ if (use_logit_softcap && !(DKQ == 128 || DKQ == 256 || DKQ == 512)) { NO_DEVICE_CODE; return; } ++ ++ constexpr int warp_size = ggml_cuda_get_physical_warp_size(); ++ constexpr int ncols = ncols1 * ncols2; ++ constexpr int nbatch_fa = ggml_cuda_fattn_mma_get_nbatch_fa(DKQ, DV, ncols); ++ constexpr int nthreads = ggml_cuda_fattn_mma_get_nthreads(DKQ, DV, ncols); ++ constexpr int nwarps = nthreads / warp_size; ++ ++ const int gqa_ratio = ne02 / ne12; ++ const int stride_Q1 = nb01 / sizeof(float2); ++ const int stride_Q2 = nb02 / sizeof(float2); ++ const int stride_K = nb11 / sizeof(half2); ++ const int stride_mask = nb31 / sizeof(half); ++ const int stride_V = V_is_K_view ? stride_K : nb21 / sizeof(half2); ++ ++ const int iter_k = (ne11 + (nbatch_fa - 1)) / nbatch_fa; ++ const int iter_j = (int(ne01.z) + (ncols1 - 1)) / ncols1; ++ const int iter_z_gqa = (gqa_ratio + (ncols2 - 1)) / ncols2; ++ ++ // Stream-k decomposition (verbatim from the stock flash_attn_ext_f16): the fixed grid of CUDA blocks ++ // is mapped onto (ntiles_dst × iter_k) work via the flat kbc index, so each block owns a contiguous ++ // [kb0_start, kb0_stop) slice of one or more output tiles. This load-balances the causal triangle ++ // (late-query tiles do ~iter_k k-iters, early ones ~1) that one-block-per-tile could not. The only ++ // DCA delta vs stock: per tile we also resolve Q_succ_f2/Q_inter_f2 + the chunk index dca_cq, and call ++ // process_tile<...,dca_fused=true> which clips its 3 phases to the owned k-range. dst_meta carries the ++ // partial (O,max,rowsum); the host launch runs the shared stream-k fixup to combine seams. ++ const int dca_c_blk = chunk / nbatch_fa; // chunk width in k-blocks (loop-invariant) ++ ++ int kbc = int64_t(blockIdx.x + 0)*(iter_k*iter_j*iter_z_gqa*ne12*ne03) / gridDim.x; ++ const int kbc_stop = int64_t(blockIdx.x + 1)*(iter_k*iter_j*iter_z_gqa*ne12*ne03) / gridDim.x; ++ ++ int kb0_start = kbc % iter_k; ++ int kb0_stop = min(iter_k, kb0_start + kbc_stop - kbc); ++ ++ while (kbc < kbc_stop && kb0_stop == iter_k) { ++ const int sequence = kbc /(iter_k*iter_j*iter_z_gqa*ne12); ++ const int z_KV = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence)/(iter_k*iter_j*iter_z_gqa); ++ const int zt_gqa = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence - iter_k*iter_j*iter_z_gqa * z_KV)/(iter_k*iter_j); ++ const int jt = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence - iter_k*iter_j*iter_z_gqa * z_KV - iter_k*iter_j * zt_gqa) / iter_k; ++ const int zt_Q = z_KV*gqa_ratio + zt_gqa*ncols2; ++ ++ const float2 * Q_f2 = (const float2 *) (Q + nb03*sequence + nb02*zt_Q); ++ const float2 * Q_succ_f2 = (const float2 *) (Q_succ + nb03*sequence + nb02*zt_Q); ++ const float2 * Q_inter_f2 = (const float2 *) (Q_inter + nb03*sequence + nb02*zt_Q); ++ const half2 * K_h2 = (const half2 *) (K + nb13*sequence + nb12*z_KV); ++ const half * mask_h = !mask ? nullptr : (const half *) (mask + nb33*(sequence % ne33)); ++ // opencoti F5 dca (#617B): per-sequence bases for the FRAGMENTED-fallback dense masks (same ++ // shape/stride as mask ⇒ same nb33/ne33 offset). NULL on the contiguous fast path. ++ const half * mask_succ_h = !mask_succ ? nullptr : (const half *) (mask_succ + nb33*(sequence % ne33)); ++ const half * mask_inter_h = !mask_inter ? nullptr : (const half *) (mask_inter + nb33*(sequence % ne33)); ++ float2 * dstk = ((float2 *) dst) + (sequence*int(ne01.z)*ne02 + zt_Q) * (DV/2); ++ const half2 * V_h2 = V_is_K_view ? K_h2 : (const half2 *) (V + nb23*sequence + nb22*z_KV); ++ ++ // opencoti F5 dca (#617B): KV_max clamps to the INTRA causal index bound — VALID only when ++ // index==pos. On the fragmented path the host passes KV_max==null (positions aren't index- ++ // ordered) so this is skipped and every phase spans the full owned KV, masks doing all selection. ++ if (KV_max) { kb0_stop = min(kb0_stop, KV_max[sequence*iter_j + jt] / nbatch_fa); } ++ const int dca_cq = pos_q[jt*ncols1] / chunk; // cq uniform per tile (analytical bands) ++ const float slope = 1.0f; // DCA: max_bias==0 ⇒ no ALiBi ++ constexpr bool is_fixup = false; ++ if (kb0_start == 0) { ++ constexpr bool needs_fixup = false; // block works on an entire tile ++ flash_attn_ext_f16_process_tile ++ (Q_f2, K_h2, V_h2, mask_h, nullptr, dstk, dst_meta, scale, slope, logit_softcap, ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, ++ kb0_start, kb0_stop, /*kb0_min=*/0, Q_succ_f2, Q_inter_f2, dca_cq, dca_c_blk, mask_succ_h, mask_inter_h); ++ } else { ++ constexpr bool needs_fixup = true; // block is missing the beginning of a tile ++ flash_attn_ext_f16_process_tile ++ (Q_f2, K_h2, V_h2, mask_h, nullptr, dstk, dst_meta, scale, slope, logit_softcap, ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, ++ kb0_start, kb0_stop, /*kb0_min=*/0, Q_succ_f2, Q_inter_f2, dca_cq, dca_c_blk, mask_succ_h, mask_inter_h); ++ } ++ ++ kbc += iter_k; ++ kbc -= kbc % iter_k; ++ kb0_start = 0; ++ kb0_stop = min(iter_k, kbc_stop - kbc); ++ } ++ ++ if (kbc >= kbc_stop) { ++ return; ++ } ++ ++ const int sequence = kbc /(iter_k*iter_j*iter_z_gqa*ne12); ++ const int z_KV = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence)/(iter_k*iter_j*iter_z_gqa); ++ const int zt_gqa = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence - iter_k*iter_j*iter_z_gqa * z_KV)/(iter_k*iter_j); ++ const int jt = (kbc - iter_k*iter_j*iter_z_gqa*ne12 * sequence - iter_k*iter_j*iter_z_gqa * z_KV - iter_k*iter_j * zt_gqa) / iter_k; ++ const int zt_Q = z_KV*gqa_ratio + zt_gqa*ncols2; ++ ++ const float2 * Q_f2 = (const float2 *) (Q + nb03*sequence + nb02*zt_Q); ++ const float2 * Q_succ_f2 = (const float2 *) (Q_succ + nb03*sequence + nb02*zt_Q); ++ const float2 * Q_inter_f2 = (const float2 *) (Q_inter + nb03*sequence + nb02*zt_Q); ++ const half2 * K_h2 = (const half2 *) (K + nb13*sequence + nb12*z_KV); ++ const half * mask_h = !mask ? nullptr : (const half *) (mask + nb33*(sequence % ne33)); ++ // opencoti F5 dca (#617B): FRAGMENTED-fallback dense mask bases (tail tile). NULL on the fast path. ++ const half * mask_succ_h = !mask_succ ? nullptr : (const half *) (mask_succ + nb33*(sequence % ne33)); ++ const half * mask_inter_h = !mask_inter ? nullptr : (const half *) (mask_inter + nb33*(sequence % ne33)); ++ float2 * dstk = ((float2 *) dst) + (sequence*int(ne01.z)*ne02 + zt_Q) * (DV/2); ++ const half2 * V_h2 = V_is_K_view ? K_h2 : (const half2 *) (V + nb23*sequence + nb22*z_KV); ++ ++ if (KV_max) { kb0_stop = min(kb0_stop, KV_max[sequence*iter_j + jt] / nbatch_fa); } // null when fragmented ++ const int dca_cq = pos_q[jt*ncols1] / chunk; ++ const float slope = 1.0f; ++ constexpr bool is_fixup = true; // last index writes to the fixup buffer to avoid races ++ constexpr bool needs_fixup = false; ++ flash_attn_ext_f16_process_tile ++ (Q_f2, K_h2, V_h2, mask_h, nullptr, dstk, dst_meta, scale, slope, logit_softcap, ++ ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, ++ kb0_start, kb0_stop, /*kb0_min=*/0, Q_succ_f2, Q_inter_f2, dca_cq, dca_c_blk, mask_succ_h, mask_inter_h); ++#else ++ GGML_UNUSED_VARS(Q, K, V, mask, mask_succ, mask_inter, sinks, KV_max, KV_min, Q_succ, Q_inter, pos_q, chunk, dst, dst_meta, ++ scale, max_bias, m0, m1, n_head_log2, logit_softcap, ne00, ne01, ne02, ne03, nb01, nb02, nb03, ++ ne10, ne11, ne12, ne13, nb11, nb12, nb13, nb21, nb22, nb23, ne31, ne32, ne33, nb31, nb32, nb33); ++ NO_DEVICE_CODE; ++#endif ++} ++ ++// opencoti F5 dca perf (T87.pD Branch B + stream-k): host launch for the fused kernel. dst is the O-only ++// temp built by ggml_cuda_flash_attn_ext_lse's regime==3 branch (op_params[3]=chunk; src[4]=pos_q, ++// src[5]=Q_succ, src[6]=Q_inter). Mirrors launch_fattn: K/V already f16 (DCA), stream-k decomposition + ++// dst_tmp_meta partials + fixup (dst_lse=nullptr — fused needs no LSE out), KV_max from the INTRA mask ++// gives the per-tile causal bound. No parallel-blocks path (stream-k always covers the triangle). ++template ++void ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ++ const ggml_tensor * KQV = dst; ++ 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]; ++ const ggml_tensor * PQ = dst->src[4]; // pos_q (I32) ++ const ggml_tensor * Qs = dst->src[5]; // Q_succ ++ const ggml_tensor * Qi = dst->src[6]; // Q_inter ++ const ggml_tensor * Ms = dst->src[7]; // opencoti F5 dca (#617B): dense SUCC mask (NULL ⇒ fast path) ++ const ggml_tensor * Mi = dst->src[8]; // opencoti F5 dca (#617B): dense INTER mask (NULL ⇒ fast path) ++ GGML_ASSERT(K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16 && "dca fused: f16 K/V required"); ++ GGML_ASSERT(PQ && Qs && Qi && mask && "dca fused: missing src tensors"); ++ // opencoti F5 dca (#617B): FRAGMENTED fallback engaged iff the host attached the dense SUCC/INTER ++ // masks (only on a non-contiguous, post-context-shift cache). Both present or both absent. ++ const bool fragmented = (Ms != nullptr); ++ GGML_ASSERT((Ms != nullptr) == (Mi != nullptr) && "dca fused: succ/inter masks are all-or-nothing"); ++ ++ const int id = ggml_cuda_get_device(); ++ const int cc = ggml_cuda_info().devices[id].cc; ++ constexpr int ncols = ncols1 * ncols2; ++ ++ const int nthreads = ggml_cuda_fattn_mma_get_nthreads (DKQ, DV, ncols, cc); ++ const int nbatch_fa = ggml_cuda_fattn_mma_get_nbatch_fa (DKQ, DV, ncols, cc); ++ const int nbatch_K2 = ggml_cuda_fattn_mma_get_nbatch_K2 (DKQ, DV, ncols, cc); ++ const int nbatch_V2 = ggml_cuda_fattn_mma_get_nbatch_V2 (DKQ, DV, ncols, cc); ++ const int nbatch_combine = ggml_cuda_fattn_mma_get_nbatch_combine(DKQ, DV, ncols, cc); ++ // opencoti F5 dca (#623): MUST match the kernel's forced Q_in_reg=false on the fused path (see ++ // flash_attn_ext_f16_process_tile) so this host shared-mem sizing matches the tile_Q + separate ++ // tile_K layout the kernel uses. This is the DCA fused case function ⇒ always false. ++ const bool Q_in_reg = false; ++ // opencoti F5 dca (#623): match the kernel's single-stage clamp on the fused path (see ++ // flash_attn_ext_f16_process_tile / _iter). For DKQ 128/256 the config returns nstages_target=2 but the ++ // fused kernel forces single-stage, so size the KV smem (nbytes_shared_KV below) for 1 stage to agree. ++ const int nstages = std::min(1, ggml_cuda_fattn_mma_get_nstages(DKQ, DV, ncols1, ncols2, cc)); ++ ++ const int cols_per_warp = std::min(ncols, get_cols_per_warp(cc)); ++ const int warp_size_host = ggml_cuda_info().devices[ctx.device].warp_size; ++ const int nwarps = nthreads / warp_size_host; ++ constexpr bool V_is_K_view = DKQ == 576; ++ ++ const size_t nbytes_shared_KV_1stage = nbatch_fa * std::max(nbatch_K2 + 4, nbatch_V2 + 4) * sizeof(half2); ++ const size_t nbytes_shared_KV_2stage = nbatch_fa * (nbatch_K2 + 4 + nbatch_V2 + 4) * sizeof(half2); ++ const size_t nbytes_shared_Q = ncols * (DKQ/2 + 4) * sizeof(half2); ++ const size_t nbytes_shared_mask = ncols1 * (nbatch_fa/2 + 4) * sizeof(half2); ++ const size_t nbytes_shared_combine = nwarps*cols_per_warp * (nbatch_combine + 4) * sizeof(half2); ++ const size_t nbytes_shared_KV = nstages <= 1 ? nbytes_shared_KV_1stage : nbytes_shared_KV_2stage; ++ const size_t nbytes_shared_total = std::max(nbytes_shared_combine, Q_in_reg ? ++ std::max(nbytes_shared_Q, nbytes_shared_KV + nbytes_shared_mask) : ++ nbytes_shared_Q + nbytes_shared_KV + nbytes_shared_mask); ++ ++ float scale = 1.0f, max_bias = 0.0f, logit_softcap = 0.0f; ++ memcpy(&scale, (const float *) KQV->op_params + 0, sizeof(float)); ++ memcpy(&max_bias, (const float *) KQV->op_params + 1, sizeof(float)); ++ memcpy(&logit_softcap, (const float *) KQV->op_params + 2, sizeof(float)); ++ const int chunk = ggml_get_op_params_i32(dst, 3); ++ if (logit_softcap != 0.0f) { scale /= logit_softcap; } ++ ++ ggml_cuda_pool & pool = ctx.pool(); ++ cudaStream_t main_stream = ctx.stream(); ++ ++ const int ntiles_x = (Q->ne[1] + ncols1 - 1) / ncols1; ++ const int gqa_ratio = Q->ne[2] / K->ne[2]; ++ const int ntiles_z_gqa = (gqa_ratio + ncols2 - 1) / ncols2; ++ const int ntiles_dst = ntiles_x * ntiles_z_gqa * K->ne[2] * Q->ne[3]; ++ ++ ggml_cuda_pool_alloc KV_max(pool); ++ ggml_cuda_pool_alloc KV_min(pool); ++ int * KV_max_ptr = nullptr; ++ // opencoti F5 dca (#617B): KV_max is the INTRA causal INDEX bound — valid only when index==pos. On ++ // the fragmented (shifted) path positions aren't index-ordered, so leave KV_max null: the kernel ++ // then spans the full KV per phase and the dense position-based masks do all selection. ++ if (!fragmented && K->ne[1] % FATTN_KQ_STRIDE == 0) { // KV_max trims INTRA to the causal bound (mask still enforces it) ++ const int s31 = mask->nb[1] / sizeof(half2); ++ const int s33 = mask->nb[3] / sizeof(half2); ++ const dim3 blk(ntiles_x, Q->ne[3], 1); ++ const dim3 bd(FATTN_KQ_STRIDE/2, 1, 1); ++ const int ne_KV_max = blk.x*blk.y; ++ const int iter_k = K->ne[1] / FATTN_KQ_STRIDE; ++ KV_max.alloc(ne_KV_max); ++ KV_min.alloc(ne_KV_max); ++ flash_attn_mask_to_KV_max<<>> ++ ((const half2 *) mask->data, KV_max.ptr, KV_min.ptr, iter_k, s31, s33); ++ CUDA_CHECK(cudaGetLastError()); ++ KV_max_ptr = KV_max.ptr; ++ } ++ ++ const uint32_t n_head = Q->ne[2]; ++ const uint32_t n_head_log2 = 1u << uint32_t(floorf(log2f(float(n_head)))); ++ const uint3 ne01 = init_fastdiv_values(Q->ne[1]); ++ ++ // opencoti F5 dca perf (T87.pD Branch B + stream-k): load-balance the causal triangle exactly like the ++ // stock D=512 path. Replicate launch_fattn's stream-k decomposition + fixup so the fused kernel inherits ++ // it: blocks own [kb0_start,kb0_stop) of the (ntiles_dst × iter_k) work, partial (O,max,rowsum) land in ++ // dst_tmp_meta, a fixup combines seams per output tile. dst_lse=nullptr — the fused op needs no LSE out. ++ const int nsm = ggml_cuda_info().devices[id].nsm; ++ const int ntiles_KV = (K->ne[1] + nbatch_fa - 1) / nbatch_fa; ++ const dim3 block_dim(warp_size_host, nwarps, 1); ++ ggml_cuda_pool_alloc dst_tmp_meta(pool); ++ ++ auto launch = [&](auto kernel) { ++#if !defined(GGML_USE_MUSA) ++ static bool raised[GGML_CUDA_MAX_DEVICES] = {false}; ++ if (!raised[id]) { ++ CUDA_CHECK(cudaFuncSetAttribute((const void *) kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total)); ++ raised[id] = true; ++ } ++#endif ++ int max_blocks_per_sm = 1; ++ CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, (const void *) kernel, block_dim.x*block_dim.y*block_dim.z, nbytes_shared_total)); ++ GGML_ASSERT(max_blocks_per_sm > 0); ++ ++ // stream-k block count (verbatim from launch_fattn): round to a multiple of ntiles_dst when the ++ // occupancy loss is <=5% so each tile gets equal blocks (skips the general fixup). ++ const int max_blocks = max_blocks_per_sm*nsm; ++ const int tiles_nwaves = (ntiles_dst + max_blocks - 1) / max_blocks; ++ const int tiles_efficiency_percent = 100 * ntiles_dst / (max_blocks*tiles_nwaves); ++ const bool use_stream_k = cc >= GGML_CUDA_CC_ADA_LOVELACE || amd_wmma_available(cc) || tiles_efficiency_percent < 75; ++ int nblocks = ntiles_dst; ++ if (use_stream_k) { ++ const int nblocks_raw = std::min(max_blocks, ntiles_KV*ntiles_dst); ++ const int nblocks_rounded = (nblocks_raw / ntiles_dst) * ntiles_dst; ++ const int eff_loss = nblocks_rounded > 0 ? 100*(nblocks_raw - nblocks_rounded)/nblocks_raw : 100; ++ nblocks = eff_loss <= 5 ? nblocks_rounded : nblocks_raw; ++ } ++ ++ // opencoti DCA #635: decode keeps the upstream stream-k path unchanged. The chunk-boundary ++ // corruption was NOT a stream-k cross-block combine artifact — that hypothesis was falsified ++ // (decode is single-block-deterministic, and forcing stream-k decode WITH the real fix below is ++ // clean). It was a host-side INTER-Q position bug, fixed in set_input_dca (llama-kv-cache.cpp: ++ // INTER-Q = 2c, not c, so far keys land at rel-distance [c+1,2c] instead of overlapping the ++ // recent/SUCC window). Verified on the count-up reproducer in BOTH decode modes + NIAH retrieval. ++ ++ float2 * dst_meta_ptr = nullptr; ++ if (ntiles_dst % nblocks != 0) { // fixup needed only when SMs work on fractional tiles ++ dst_tmp_meta.alloc(size_t(nblocks) * ncols * (2 + DV/2)); ++ dst_meta_ptr = dst_tmp_meta.ptr; ++ } ++ ++ const dim3 grid(nblocks, 1, 1); ++ kernel<<>>( ++ (const char *) Q->data, (const char *) K->data, (const char *) V->data, ++ (const char *) mask->data, ++ Ms ? (const char *) Ms->data : nullptr, // opencoti F5 dca (#617B): dense SUCC mask (null ⇒ fast path) ++ Mi ? (const char *) Mi->data : nullptr, // opencoti F5 dca (#617B): dense INTER mask (null ⇒ fast path) ++ nullptr, KV_max_ptr, nullptr, ++ (const char *) Qs->data, (const char *) Qi->data, (const int *) PQ->data, chunk, ++ (float *) KQV->data, dst_meta_ptr, ++ scale, 0.0f /*max_bias*/, 1.0f /*m0*/, 1.0f /*m1*/, n_head_log2, logit_softcap, ++ Q->ne[0], ne01, Q->ne[2], Q->ne[3], Q->nb[1], Q->nb[2], Q->nb[3], ++ K->ne[0], K->ne[1], K->ne[2], K->ne[3], K->nb[1], K->nb[2], K->nb[3], ++ V->nb[1], V->nb[2], V->nb[3], ++ mask->ne[1], mask->ne[2], mask->ne[3], mask->nb[1], mask->nb[2], mask->nb[3]); ++ CUDA_CHECK(cudaGetLastError()); ++ ++ // stream-k fixup: combine partial results across blocks for each output tile (verbatim dispatch). ++ if (nblocks % ntiles_dst == 0 && nblocks > ntiles_dst) { ++ const int bpt = nblocks / ntiles_dst; // uniform: equal blocks per tile ++ const uint3 fd0 = init_fastdiv_values(ntiles_x * ntiles_z_gqa * K->ne[2]); ++ const uint3 fd1 = init_fastdiv_values(ntiles_x * ntiles_z_gqa); ++ const uint3 fd2 = init_fastdiv_values(ntiles_x); ++ const dim3 bdc(DV, 1, 1); ++ const dim3 bnc((unsigned) ntiles_dst, ncols1, ncols2); ++ flash_attn_stream_k_fixup_uniform<<>> ++ ((float *) KQV->data, /*dst_lse=*/nullptr, dst_tmp_meta.ptr, ++ Q->ne[1], Q->ne[2], K->ne[2], nblocks, gqa_ratio, bpt, fd0, fd1, fd2); ++ CUDA_CHECK(cudaGetLastError()); ++ } else if (ntiles_dst % nblocks != 0) { ++ const int total_work = ntiles_KV * ntiles_dst; // general: fractional, unequal blocks per tile ++ const uint3 fd_k_j_z_ne12 = init_fastdiv_values(ntiles_KV * ntiles_x * ntiles_z_gqa * K->ne[2]); ++ const uint3 fd_k_j_z = init_fastdiv_values(ntiles_KV * ntiles_x * ntiles_z_gqa); ++ const uint3 fd_k_j = init_fastdiv_values(ntiles_KV * ntiles_x); ++ const uint3 fd_k = init_fastdiv_values(ntiles_KV); ++ const dim3 bdc(DV, 1, 1); ++ const dim3 bnc((unsigned) nblocks, ncols1, ncols2); ++ flash_attn_stream_k_fixup_general<<>> ++ ((float *) KQV->data, /*dst_lse=*/nullptr, 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); ++ CUDA_CHECK(cudaGetLastError()); ++ } ++ }; ++ if (logit_softcap == 0.0f) { ++ launch(flash_attn_ext_f16_dca_fused); ++ } else { ++ launch(flash_attn_ext_f16_dca_fused); ++ } ++} ++ + #define DECL_FATTN_MMA_F16_CASE(DKQ, DV, ncols1, ncols2) \ + template void ggml_cuda_flash_attn_ext_mma_f16_case \ + (ggml_backend_cuda_context & ctx, ggml_tensor * dst) \ +diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh +--- a/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh ++++ b/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh +@@ -794,6 +794,7 @@ static __global__ void flash_attn_tile( + const char * __restrict__ mask, + const char * __restrict__ sinks, + const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + float * __restrict__ dst, + float2 * __restrict__ dst_meta, + const float scale, +@@ -819,7 +820,7 @@ static __global__ void flash_attn_tile( + #endif // GGML_USE_WMMA_FATTN + (use_logit_softcap && !(DV == 128 || DV == 256 || DV == 512)) + ) { +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +@@ -946,17 +947,22 @@ static __global__ void flash_attn_tile( + + // Main loop over KV cache: + const int k_VKQ_max = KV_max ? KV_max[sequence*gridDim.x + blockIdx.x] : ne11; ++ // opencoti F5 dca perf (T87.pD-opt O4): low band edge — skip K/V tiles entirely below it (DCA ++ // INTRA/SUCC regimes, SWA windows). k_VKQ_min==0 for causal masks → the guards below never fire. ++ const int k_VKQ_min = KV_min ? KV_min[sequence*gridDim.x + blockIdx.x] : 0; + if (ncols2 == 1) { + // Branch with out-of-bounds checks. + int k_VKQ_0 = blockIdx.y*nbatch_fa; + while (k_VKQ_0 < k_VKQ_max - nbatch_fa) { +- constexpr bool oob_check = false; +- flash_attn_tile_iter +- (Q_tmp, K_h2, V_h2, maskh, ne01, logit_softcap, slope, KQ, KV_tmp, +- stride_K2, stride_V2, stride_mask, KQ_max, KQ_sum, VKQ, k_VKQ_0, k_VKQ_max, col_Q_0); ++ if (k_VKQ_0 + nbatch_fa > k_VKQ_min) { ++ constexpr bool oob_check = false; ++ flash_attn_tile_iter ++ (Q_tmp, K_h2, V_h2, maskh, ne01, logit_softcap, slope, KQ, KV_tmp, ++ stride_K2, stride_V2, stride_mask, KQ_max, KQ_sum, VKQ, k_VKQ_0, k_VKQ_max, col_Q_0); ++ } + k_VKQ_0 += gridDim.y*nbatch_fa; + } +- if (k_VKQ_0 < k_VKQ_max) { ++ if (k_VKQ_0 < k_VKQ_max && k_VKQ_0 + nbatch_fa > k_VKQ_min) { + constexpr bool oob_check = true; + flash_attn_tile_iter + (Q_tmp, K_h2, V_h2, maskh, ne01, logit_softcap, slope, KQ, KV_tmp, +@@ -965,6 +971,9 @@ static __global__ void flash_attn_tile( + } else { + // Branch without out-of-bounds checks. + for (int k_VKQ_0 = blockIdx.y*nbatch_fa; k_VKQ_0 < k_VKQ_max; k_VKQ_0 += gridDim.y*nbatch_fa) { ++ if (k_VKQ_0 + nbatch_fa <= k_VKQ_min) { ++ continue; ++ } + constexpr bool oob_check = false; + flash_attn_tile_iter + (Q_tmp, K_h2, V_h2, maskh, ne01, logit_softcap, slope, KQ, KV_tmp, +@@ -1126,7 +1135,7 @@ static __global__ void flash_attn_tile( + } + } + #else +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh +--- a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh ++++ b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh +@@ -25,6 +25,7 @@ static __global__ void flash_attn_ext_vec( + const char * __restrict__ mask, + const char * __restrict__ sinks, + const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + float * __restrict__ dst, + float2 * __restrict__ dst_meta, + const float scale, +@@ -45,7 +46,7 @@ static __global__ void flash_attn_ext_vec( + + // Skip unused kernel variants for faster compilation: + if (use_logit_softcap && !(D == 128 || D == 256)) { +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +@@ -246,12 +247,19 @@ static __global__ void flash_attn_ext_vec( + } + + const int k_VKQ_max = KV_max ? KV_max[sequence*gridDim.x + blockIdx.x] : ne11; ++ const int k_VKQ_min = KV_min ? KV_min[sequence*gridDim.x + blockIdx.x] : 0; // opencoti F5 dca perf (T87.pD-opt O4) + K += blockIdx.y*nthreads * nb11; + V += blockIdx.y*nthreads * nb21; + maskh += blockIdx.y*nthreads; + for (int k_VKQ_0 = blockIdx.y*nthreads; k_VKQ_0 < k_VKQ_max; k_VKQ_0 += gridDim.y*nthreads, + // Increment pointers after each loop: + K += gridDim.y*nthreads*nb11, V += gridDim.y*nthreads*nb21, maskh += gridDim.y*nthreads) { ++ // opencoti F5 dca perf (T87.pD-opt O4): skip K/V tiles entirely below the low band edge ++ // (DCA INTRA/SUCC regimes, SWA windows). k_VKQ_min==0 for causal masks → never skips. The ++ // for-increment above still advances the K/V/mask pointers, so they stay in sync. ++ if (k_VKQ_0 + nthreads <= k_VKQ_min) { ++ continue; ++ } + + // Calculate KQ tile and keep track of new maximum KQ values: + float KQ_reg[ncols]; // KQ in registers. +@@ -512,7 +520,7 @@ static __global__ void flash_attn_ext_vec( + dst_meta[((sequence*int(ne01.z) + ic0 + tid)*ne02 + head)*gridDim.y + blockIdx.y] = make_float2(KQ_max[tid], KQ_sum[tid]); + } + #else +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu b/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu +--- a/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu ++++ b/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu +@@ -30,6 +30,7 @@ static __global__ void flash_attn_ext_f16( + const char * __restrict__ mask, + const char * __restrict__ sinks, + const int * __restrict__ KV_max, ++ const int * __restrict__ KV_min, // opencoti F5 dca perf (T87.pD-opt O4): low-edge band trim + float * __restrict__ dst, + float2 * __restrict__ dst_meta, + const float scale, +@@ -190,7 +191,13 @@ static __global__ void flash_attn_ext_f16( + + // Iterate over ne11 == previous tokens: + const int k_VKQ_max = KV_max ? KV_max[sequence*gridDim.x + blockIdx.x] : ne11; ++ const int k_VKQ_min = KV_min ? KV_min[sequence*gridDim.x + blockIdx.x] : 0; // opencoti F5 dca perf (T87.pD-opt O4) + for (int k_VKQ_0 = blockIdx.y*FATTN_KQ_STRIDE; k_VKQ_0 < k_VKQ_max; k_VKQ_0 += gridDim.y*FATTN_KQ_STRIDE) { ++ // opencoti F5 dca perf (T87.pD-opt O4): skip K/V tiles entirely below the low band edge ++ // (DCA INTRA/SUCC regimes, SWA windows). k_VKQ_min==0 for causal masks → never skips. ++ if (k_VKQ_0 + FATTN_KQ_STRIDE <= k_VKQ_min) { ++ continue; ++ } + // Calculate tile of KQ: + #pragma unroll + for (int i_KQ_0 = 0; i_KQ_0 < FATTN_KQ_STRIDE; i_KQ_0 += KQ_stride_tc) { +@@ -494,7 +501,7 @@ static __global__ void flash_attn_ext_f16( + dst_meta[j_dst_unrolled] = dst_meta_val; + } + #else +- GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale, ++ GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + max_bias, m0, m1, n_head_log2, logit_softcap, + ne00, ne01, ne02, ne03, + nb01, nb02, nb03, +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 +@@ -10,6 +10,23 @@ + // 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; ++// opencoti F5 dca perf (T87.pD-opt O2): set by launch_fattn at its dispatch site to report whether a ++// finalize kernel (the populator of dst_lse) actually ran. The DCA _lse op reads it right after the FA ++// call to decide whether the LSE was emitted for free (common: large-KV prefill/decode takes a finalize ++// path) or must be recomputed by the streaming_lse_kernel fallback (rare: in-kernel single-block path). ++thread_local bool opencoti_fattn_dst_lse_written = false; ++ ++// opencoti F5 dca perf (T87.pD-opt O6 Tier B): analytical band-bounds side-channel. A DCA SUCC/INTER ++// regime is a CONTIGUOUS causal key-band per query ([(cq-1)c, cq*c) for SUCC, [0,(cq-1)c) for INTER, ++// cq = pos_q/c) that needs NO per-element causal — every band key is < pos_q for every query. So instead ++// of building+uploading a dense [n_kv x n_tokens] mask (the host O(n_kv*n_tokens) prefill bottleneck), ++// ggml_cuda_flash_attn_ext_lse sets these before the inner FA dispatch; launch_fattn then fills the ++// per-tile KV_min/KV_max ANALYTICALLY from pos_q (device, len n_q) + chunk and passes mask=nullptr. ++// Drained back to nullptr/0/-1 after the dispatch (mirrors opencoti_fattn_dst_lse). regime is ++// DCA_REGIME_SUCC(1) / _INTER(2); INTRA(0) keeps its native-causal mask and never sets these. ++thread_local const int * opencoti_fattn_dca_pos_q = nullptr; ++thread_local int opencoti_fattn_dca_chunk = 0; ++thread_local int opencoti_fattn_dca_regime = -1; + + template + static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { +@@ -375,7 +392,14 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const + + // The effective batch size for the kernel can be increased by gqa_ratio. + // The kernel versions without this optimization are also used for ALiBi, if there is no mask, or if the KV cache is not padded, +- bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B / Path B): SUCC/INTER analytical-band regimes pass mask=nullptr ++ // and mask purely via the on-GPU KV band (kb0_min/kb0_stop from KV_min/KV_max). ggml_cuda_flash_attn_ext_lse ++ // arms the thread_local opencoti_fattn_dca_pos_q around the inner FA call, so it is readable here at selection ++ // time → allow the gqa-batched MMA kernel for a null-mask band op. The band-aware D512 kernel (fattn-mma-f16.cuh) ++ // skips the mask load/add when mask_h==nullptr and relies solely on the tile band. In every reachable STOCK ++ // state ncols2>1 ⟹ mask!=null, so this disjunct changes nothing outside DCA. ++ const bool dca_band = (opencoti_fattn_dca_pos_q != nullptr); ++ bool gqa_opt_applies = gqa_ratio >= 2 && (mask || dca_band) && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; + for (const ggml_tensor * t : {Q, K, V, mask}) { + if (t == nullptr || ggml_is_quantized(t->type)) { + continue; +@@ -597,6 +621,192 @@ 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 F5 dca (#593) — fwd-decl: streaming_lse_kernel is defined ~line 645, but the LSE-op ++// forward below launches it from above its definition; declare it here so the call resolves. ++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, ++ // opencoti F5 dca perf (O6): analytical band for SUCC/INTER (defaults keep legacy 18-arg callers). ++ const int * __restrict__ dca_pos_q = nullptr, const int dca_chunk = 0, const int dca_regime = 0); ++ ++// opencoti F5 dca (#593) — CUDA forward for GGML_OP_FLASH_ATTN_EXT_LSE. Runs a standard ++// flash-attention for the NORMALIZED O (written into the leading DV*n_rows of dst, shaped ++// exactly like flash_attn_ext) via the proven streaming-op trick of an O-typed view of dst; ++// then computes the per-row LSE into the trailing n_rows of dst with streaming_lse_kernel — ++// the same softcap-aware recompute the rolling-KV streaming op uses for its no-finalize ++// path, so it works for ANY n_q (incl. 512k prefill) and avoids the decode_lse / bug-263 ++// finalize fragility. dst packs [O (DV*n_rows) | lse (n_rows)] (ne[0] = DV+1). The DCA ++// 3-regime graph-math combine reads O and lse back via offset views. ++// opencoti F5 dca perf (T87.pD Branch B): ncols1/ncols2 selection for the FUSED DCA kernel. DKQ=DV is the ++// model head dim (Gemma-4 26B-A4B = 512, qwen35moe = 256, qwen2/qwen3/qwen3moe = 128). Mirrors the stock ++// switch_ncols2/switch_ncols1 GQA-opt path (the fused op guarantees a mask, max_bias==0, aligned f16 K/V ⇒ ++// use_gqa_opt is always true here). ++// opencoti F5 dca (#623): templated on DKQ so each arch gets a head-dim-correct kernel. The earlier code ++// hardcoded <512,512> (Gemma's head dim) for EVERY model: supports_op (ggml-cuda.cu) admits head_dim ++// 128/256/512, so Qwen ops reached this fused path and launched a 512-wide kernel over 256/128-wide K/V/Q — ++// striding past each head into the next head's memory → multi-chunk garbage from token 1 (no crash). The ++// fused case + process_tile are already DKQ-templated (softcap whitelist DKQ∈{128,256,512}); the config ++// table has matching entries for all three. The single_chunk path (plain ggml_flash_attn_ext) was never ++// affected — only the analytical-bands fused dispatch went through here. ++template ++static void ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ++ const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; ++ const ggml_tensor * Q = dst->src[0]; ++ if constexpr (ncols2 <= 8) { ++ if (turing_mma_available(cc) && Q->ne[1] <= 8/ncols2) { ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case(ctx, dst); return; ++ } ++ } ++ if constexpr (ncols2 <= 16) { ++ if ((turing_mma_available(cc) || amd_wmma_available(cc)) && Q->ne[1] <= 16/ncols2) { ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case(ctx, dst); return; ++ } ++ } ++ if (ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_TURING || amd_wmma_available(cc) || Q->ne[1] <= 32/ncols2) { ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case(ctx, dst); return; ++ } ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_case(ctx, dst); ++} ++ ++// gqa_ratio → ncols2 selection for a fixed head dim DKQ (the old switch body, now DKQ-parameterized). ++template ++static void ggml_cuda_flash_attn_ext_mma_f16_dca_fused_dispatch_gqa(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ++ const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; ++ const ggml_tensor * Q = dst->src[0]; ++ const ggml_tensor * K = dst->src[1]; ++ GGML_ASSERT(Q->ne[2] % K->ne[2] == 0); ++ const int gqa_ratio = Q->ne[2] / K->ne[2]; ++ if (cc == GGML_CUDA_CC_VOLTA) { ++ if (gqa_ratio % 8 == 0) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ if (gqa_ratio % 4 == 0) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ if (gqa_ratio % 2 == 0) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; ++ } ++ if (gqa_ratio > 4) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ if (gqa_ratio > 2) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ if (gqa_ratio > 1) { ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); return; } ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch_ncols1(ctx, dst); ++} ++ ++// opencoti F5 dca (#623): head-dim switch mirroring the stock ggml_cuda_flash_attn_ext_mma_f16 path. DCA ++// only ever uses DKQ==DV (key_length==value_length on every supported arch), so dispatch on Q->ne[0] alone. ++static void ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ++ const ggml_tensor * Q = dst->src[0]; ++ const ggml_tensor * V = dst->src[2]; ++ switch (Q->ne[0]) { ++ case 128: ++ GGML_ASSERT(V->ne[0] == 128); ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_dispatch_gqa<128>(ctx, dst); break; ++ case 256: ++ GGML_ASSERT(V->ne[0] == 256); ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_dispatch_gqa<256>(ctx, dst); break; ++ case 512: ++ GGML_ASSERT(V->ne[0] == 512); ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_dispatch_gqa<512>(ctx, dst); break; ++ default: ++ GGML_ABORT("dca fused: unsupported head dim %d (expected 128/256/512)", (int) Q->ne[0]); ++ break; ++ } ++} ++ ++void ggml_cuda_flash_attn_ext_lse(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ++ ggml_cuda_set_device(ctx.device); ++ const ggml_tensor * Q = dst->src[0]; ++ const ggml_tensor * K = dst->src[1]; ++ const ggml_tensor * V = dst->src[2]; ++ const ggml_tensor * M = dst->src[3]; ++ ++ const int64_t DV = V->ne[0]; ++ const int64_t n_head = Q->ne[2]; ++ const int64_t n_q = Q->ne[1]; ++ const int64_t n_batch = Q->ne[3]; ++ const int64_t n_rows = n_head * n_q * n_batch; ++ GGML_ASSERT(dst->ne[0] == DV + 1 && "flash_attn_ext_lse dst must pack [O|lse] on ne[0]"); ++ GGML_ASSERT(Q->type == GGML_TYPE_F32 && "flash_attn_ext_lse expects f32 Q for the lse recompute"); ++ GGML_ASSERT(K->type == GGML_TYPE_F16 && "flash_attn_ext_lse expects f16 K for the lse recompute"); ++ ++ // (1) NORMALIZED O into the leading DV*n_rows of dst (contiguous [DV,n_head,n_q,n_batch]). ++ ggml_tensor o = *dst; ++ o.op = GGML_OP_FLASH_ATTN_EXT; ++ o.src[4] = nullptr; // no attention sinks for DCA ++ // opencoti F5 dca perf (O6): the DCA band metadata travels via the side-channel, not op_params; clear ++ // op_params[3..4] on the inner EXT temp so no downstream op_params reader can misinterpret them. ++ o.op_params[3] = 0; ++ o.op_params[4] = 0; ++ o.ne[0] = DV; o.ne[1] = n_head; o.ne[2] = n_q; o.ne[3] = n_batch; ++ o.nb[0] = sizeof(float); ++ o.nb[1] = o.nb[0] * o.ne[0]; ++ o.nb[2] = o.nb[1] * o.ne[1]; ++ o.nb[3] = o.nb[2] * o.ne[2]; ++ o.data = dst->data; ++ ++ // opencoti F5 dca perf (T87.pD-opt O2): the FA finalize kernels already compute the per-row LSE ++ // (max + log(sumexp)) and can write it for free via the rolling-kv side-channel. Point it at the ++ // SAME trailing region streaming_lse_kernel would fill — the finalize row order (h + n_head*qi + ++ // n_head*n_q*b) matches O's row order, which is exactly what the host 3-regime combine reads back. ++ // launch_fattn sets opencoti_fattn_dst_lse_written iff a finalize path ran; on the rare in-kernel ++ // single-block path it stays false and we fall back to the streaming recompute below. This removes ++ // one full O(L^2) pass per regime on the multi-chunk (>chunk) path. ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): when this is a SUCC/INTER regime (src[4]=pos_q set, ++ // mask=src[3]=nullptr), arm the analytical band side-channel so launch_fattn fills KV_min/KV_max from ++ // positions (no dense mask to scan). INTRA leaves src[4]=NULL ⇒ side-channel disarmed (legacy path). ++ const ggml_tensor * PQ = dst->src[4]; ++ const int dca_chunk = ggml_get_op_params_i32(dst, 3); ++ const int dca_regime = ggml_get_op_params_i32(dst, 4); ++ ++ // opencoti F5 dca perf (T87.pD Branch B): FUSED single-pass DCA. The three band passes (INTER/SUCC/ ++ // INTRA) are folded into ONE block-per-tile kernel that streams K/V once and swaps the pre-roped Q ++ // variant at the 2 phase boundaries — no per-regime LSE emit, no host combine, no streaming_lse ++ // fallback. The fused case reads its inputs straight off dst (src[0]=Q_intra, src[1]=K, src[2]=V, ++ // src[3]=mask_intra, src[4]=pos_q_real, src[5]=Q_succ, src[6]=Q_inter, op_params[3]=chunk) and writes ++ // normalized O into the leading DV*n_rows of dst's [O|lse] buffer (the trailing lse row is left ++ // untouched — dca.cpp slices it off). Bands partition the causal range [0,i] ⇒ one online-softmax ++ // pass is bit-faithful to the 3-op + host LSE combine it replaces. ++ if (dca_regime == 3) { ++ ggml_cuda_flash_attn_ext_mma_f16_dca_fused_switch(ctx, dst); ++ return; ++ } ++ ++ float * lse_out = (float *) dst->data + DV * n_rows; ++ opencoti_fattn_dst_lse = lse_out; ++ opencoti_fattn_dst_lse_written = false; ++ if (PQ) { ++ opencoti_fattn_dca_pos_q = (const int *) PQ->data; ++ opencoti_fattn_dca_chunk = dca_chunk; ++ opencoti_fattn_dca_regime = dca_regime; ++ } ++ ggml_cuda_flash_attn_ext(ctx, &o); ++ opencoti_fattn_dst_lse = nullptr; // defensive: don't leak the side-channels into the next FA op ++ opencoti_fattn_dca_pos_q = nullptr; ++ opencoti_fattn_dca_chunk = 0; ++ opencoti_fattn_dca_regime = -1; ++ ++ // (2) Fallback per-row LSE recompute — only when the finalize didn't emit it for free. ++ if (!opencoti_fattn_dst_lse_written) { ++ float scale = 1.0f, logit_softcap = 0.0f; ++ memcpy(&scale, (const float *) dst->op_params + 0, sizeof(float)); ++ memcpy(&logit_softcap, (const float *) dst->op_params + 2, sizeof(float)); ++ const int gqa_ratio = (int) (n_head / K->ne[2]); ++ cudaStream_t stream = ctx.stream(); ++ streaming_lse_kernel<<<(unsigned) n_rows, WARP_SIZE, 0, stream>>>( ++ (const char *) Q->data, (const char *) K->data, ++ M ? (const char *) M->data : nullptr, lse_out, ++ scale, logit_softcap, (int) K->ne[1], gqa_ratio, (int) K->ne[0], ++ (int) n_head, (int) n_q, ++ Q->nb[1], Q->nb[2], Q->nb[3], K->nb[1], K->nb[2], K->nb[3], ++ M ? M->nb[0] : 0, M ? M->nb[1] : 0, ++ // opencoti F5 dca perf (O6): band-aware LSE for SUCC/INTER when the FA took a no-finalize ++ // path (rare; e.g. decode n_q==1). PQ==NULL (INTRA/legacy) ⇒ regime 0 ⇒ full-KV scan. ++ PQ ? (const int *) PQ->data : nullptr, dca_chunk, PQ ? dca_regime : 0); ++ CUDA_CHECK(cudaGetLastError()); ++ } ++} ++ + // ============================================================================ + // opencoti-hook: f5-rolling-kv — W4 M7-D #304 / R2-b S2 (#311). + // Streaming flash-attention via key-axis tiling + LSE-weighted online-softmax. +@@ -636,7 +846,13 @@ static __global__ void streaming_lse_kernel( + 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 int64_t m_nb0, const int64_t m_nb1, ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): analytical band for SUCC/INTER. When dca_regime>0 ++ // there is no mask; this fallback (rare: no-finalize path, e.g. decode n_q==1) must compute the ++ // LSE over EXACTLY the regime's contiguous key band [lo,hi) so it matches the band-limited O the ++ // FA produced (KV_min/KV_max). Key index == position (single-stream contiguous prefill). ++ // (defaults are declared on the forward declaration above — C++ forbids repeating them here.) ++ const int * __restrict__ dca_pos_q, const int dca_chunk, const int dca_regime) { + const int row = blockIdx.x; // flattened (h, qi, b) + const int h = row % n_head; + const int qi = (row / n_head) % n_q; +@@ -648,9 +864,22 @@ static __global__ void streaming_lse_kernel( + 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; + ++ // opencoti F5 dca O6: derive this query's regime band [band_lo, band_hi) in key indices. ++ int band_lo = 0, band_hi = tile_kv; ++ if (dca_regime > 0) { ++ const int p1 = dca_pos_q[qi]; ++ const int cq = p1 / dca_chunk; ++ if (dca_regime == 1) { band_lo = (cq - 1) * dca_chunk; band_hi = cq * dca_chunk; } // SUCC ++ else { band_lo = 0; band_hi = (cq - 1) * dca_chunk; } // INTER ++ if (band_lo < 0) { band_lo = 0; } ++ if (band_hi < 0) { band_hi = 0; } ++ if (band_hi > tile_kv) { band_hi = tile_kv; } ++ if (band_lo > band_hi) { band_lo = band_hi; } ++ } ++ + float m = -INFINITY; + float l = 0.0f; +- for (int j = 0; j < tile_kv; ++j) { ++ for (int j = band_lo; j < band_hi; ++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) { +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 +@@ -4,6 +4,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 F5 dca — #593: flash-attention emitting normalized O + per-row lse (packed [DV+1]). ++void ggml_cuda_flash_attn_ext_lse(ggml_backend_cuda_context & ctx, 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 +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 +@@ -3146,6 +3146,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg + case GGML_OP_FLASH_ATTN_EXT: + ggml_cuda_flash_attn_ext(ctx, dst); + break; ++ case GGML_OP_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 ++ ggml_cuda_flash_attn_ext_lse(ctx, dst); ++ break; + case GGML_OP_CROSS_ENTROPY_LOSS: + ggml_cuda_cross_entropy_loss(ctx, dst); + break; +@@ -5615,6 +5618,23 @@ 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_FLASH_ATTN_EXT_LSE: { // opencoti F5 dca — #593 ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B / Path B): a SUCC/INTER analytical-band regime carries ++ // mask=src[3]=null and pos_q=src[4]!=null. It masks purely via the on-GPU KV band; the band-aware D512 ++ // MMA kernel handles a null mask in the gqa-batched path. The strict FA-supported check (below) rejects ++ // a null mask at D512, so short-circuit to supported here. The op is CUDA-only (no CPU impl) and its ++ // Q/K/V shapes are identical to the INTRA regime's op which strict-passes, so dims are valid by ++ // construction. INTRA keeps src[3]=mask, src[4]=null → falls through to the strict path below. ++ if (op->src[3] == nullptr && op->src[4] != nullptr) { ++ return true; ++ } ++ ggml_tensor o = *op; // dst packs [O|lse] (ne[0]=DV+1); check the O FA (ne[0]=DV) ++ o.ne[0] = op->src[2]->ne[0]; ++ o.src[4] = nullptr; // opencoti F5 dca perf (O6): src[4] is DCA pos_q (I32), NOT attention ++ // sinks — clear it so the FA-supported check (which reads src[4] as ++ // sinks) doesn't see a non-F32 tensor and report unsupported→CPU. ++ return ggml_cuda_flash_attn_ext_supported(dev_ctx->device, &o); ++ } + 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 +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 +@@ -1131,9 +1131,10 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { + "STREAMING_FLASH_ATTN", + "FLASH_ATTN_TAIL_PARTIAL", + "TURBO_WHT", ++ "FLASH_ATTN_EXT_LSE", + }; + +-static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); ++static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); + + static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { + "none", +@@ -1245,9 +1246,10 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { + "streaming_flash_attn(q,k,v)", + "flash_attn_tail_partial(q,k,v)", + "turbo_wht(x)", ++ "flash_attn_ext_lse(q,k,v)", + }; + +-static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); ++static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); + + static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); + +@@ -5628,6 +5630,137 @@ struct ggml_tensor * ggml_flash_attn_ext_tail_partial( + return result; + } + ++// opencoti F5 dca (#593) — flash-attention emitting NORMALIZED O + per-row LSE packed on ++// the head-dim axis: dst = [DV+1, n_head, n_q, n_batch]; row [0,DV) = O, row [DV] = lse ++// (= m + logf(sum exp); softcap + sinks folded in by the kernel). Same construction as ++// ggml_flash_attn_ext but the CUDA forward also writes the per-row LSE the DCA 3-regime ++// graph-math combine consumes (O = sum_r w_r O_r / sum_r w_r, w_r = exp(lse_r - max)). ++struct ggml_tensor * ggml_flash_attn_ext_lse( ++ 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, ++ struct ggml_tensor * pos_q, ++ int chunk_size, ++ int regime) { ++ 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); ++ } ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): analytical band-bounds form for SUCC/INTER. ++ // pos_q is the device I32 query-position vector (len n_q); the CUDA op derives the per-tile KV band ++ // from pos_q[0]/chunk_size for the regime, so NO dense mask is needed (and max_bias must be 0 — ++ // ALiBi would need the mask). INTRA passes pos_q==NULL and keeps its native-causal mask. ++ if (pos_q) { ++ GGML_ASSERT(mask == NULL && "dca O6: analytical (pos_q) regimes pass mask=NULL"); ++ GGML_ASSERT(pos_q->type == GGML_TYPE_I32); ++ GGML_ASSERT(pos_q->ne[0] == q->ne[1] && "dca O6: pos_q length must equal n_q"); ++ GGML_ASSERT(max_bias == 0.0f && "dca O6: analytical band path does not support ALiBi"); ++ GGML_ASSERT(chunk_size > 0 && (regime == 1 || regime == 2) && "dca O6: regime is SUCC(1)/INTER(2)"); ++ } ++ ++ // permute(0,2,1,3) output like flash_attn_ext, with ONE extra head-dim slot for lse. ++ int64_t ne[4] = { v->ne[0] + 1, 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)); ++ // op_params[3..4]: DCA analytical band metadata (0 for the legacy INTRA/mask form; ggml_new_tensor ++ // zero-inits op_params, so a pre-O6 call without pos_q leaves these at 0 ⇒ CUDA takes the mask path). ++ ggml_set_op_params_i32(result, 3, pos_q ? chunk_size : 0); ++ ggml_set_op_params_i32(result, 4, pos_q ? regime : 0); ++ ++ result->op = GGML_OP_FLASH_ATTN_EXT_LSE; ++ result->src[0] = q; ++ result->src[1] = k; ++ result->src[2] = v; ++ result->src[3] = mask; ++ result->src[4] = pos_q; // opencoti F5 dca perf (O6): device I32 query positions (NULL for INTRA) ++ ++ return result; ++} ++ ++// opencoti F5 dca perf (T87.pD Branch B) — FUSED single-pass DCA flash-attention. Encoded as a ++// GGML_OP_FLASH_ATTN_EXT_LSE node with regime=3 (FUSED); the CUDA forward walks the causal range in ++// 3 contiguous key-phases (INTER→SUCC→INTRA), reloading the matching pre-roped Q at each phase ++// boundary, with one continuous online-softmax — bit-faithful to the 3-op + host-combine path (the ++// regime bands partition [0,i]). Output keeps the [DV+1,...] LSE layout for table reuse; the lse row ++// is unused (the single pass needs no combine). ++struct ggml_tensor * ggml_flash_attn_ext_dca_fused( ++ struct ggml_context * ctx, ++ struct ggml_tensor * q_intra, ++ struct ggml_tensor * q_succ, ++ struct ggml_tensor * q_inter, ++ struct ggml_tensor * k, ++ struct ggml_tensor * v, ++ struct ggml_tensor * mask_intra, ++ struct ggml_tensor * mask_succ, ++ struct ggml_tensor * mask_inter, ++ struct ggml_tensor * pos_q_real, ++ float scale, ++ float max_bias, ++ float logit_softcap, ++ int chunk_size) { ++ GGML_ASSERT(ggml_can_mul_mat(k, q_intra)); ++ GGML_ASSERT(q_intra->ne[3] == k->ne[3]); ++ GGML_ASSERT(q_intra->ne[3] == v->ne[3]); ++ // the 3 Q variants are the same query tensor re-roped per regime: identical shape. ++ GGML_ASSERT(q_succ && ggml_are_same_shape(q_intra, q_succ)); ++ GGML_ASSERT(q_inter && ggml_are_same_shape(q_intra, q_inter)); ++ GGML_ASSERT(mask_intra && mask_intra->type == GGML_TYPE_F16); ++ GGML_ASSERT(ggml_is_contiguous(mask_intra)); ++ GGML_ASSERT(q_intra->ne[2] % mask_intra->ne[2] == 0); ++ GGML_ASSERT(q_intra->ne[3] % mask_intra->ne[3] == 0); ++ // opencoti F5 dca (#617B): FRAGMENTED-fallback masks travel together (both NULL on the contiguous ++ // fast path, both present on a shifted cache). When present they must match the INTRA mask's dtype ++ // + shape so the kernel can offset all three identically (same stride_mask). ++ GGML_ASSERT((mask_succ == NULL) == (mask_inter == NULL) && "dca fused: succ/inter masks are all-or-nothing"); ++ if (mask_succ) { ++ GGML_ASSERT(mask_succ->type == GGML_TYPE_F16 && ggml_is_contiguous(mask_succ)); ++ GGML_ASSERT(mask_inter->type == GGML_TYPE_F16 && ggml_is_contiguous(mask_inter)); ++ GGML_ASSERT(ggml_are_same_shape(mask_intra, mask_succ)); ++ GGML_ASSERT(ggml_are_same_shape(mask_intra, mask_inter)); ++ } ++ GGML_ASSERT(pos_q_real && pos_q_real->type == GGML_TYPE_I32); ++ GGML_ASSERT(pos_q_real->ne[0] == q_intra->ne[1] && "dca fused: pos_q_real length must equal n_q"); ++ GGML_ASSERT(max_bias == 0.0f && "dca fused: analytical band path does not support ALiBi"); ++ GGML_ASSERT(chunk_size > 0); ++ ++ // same [DV+1, n_head, n_q, n_batch] layout as ggml_flash_attn_ext_lse (lse row unused here). ++ int64_t ne[4] = { v->ne[0] + 1, q_intra->ne[2], q_intra->ne[1], q_intra->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)); ++ ggml_set_op_params_i32(result, 3, chunk_size); ++ ggml_set_op_params_i32(result, 4, 3); // regime==3 ⇒ FUSED (CUDA forward branches on this) ++ ++ result->op = GGML_OP_FLASH_ATTN_EXT_LSE; ++ result->src[0] = q_intra; ++ result->src[1] = k; ++ result->src[2] = v; ++ result->src[3] = mask_intra; ++ result->src[4] = pos_q_real; // device I32 real abs query positions (band channel) ++ result->src[5] = q_succ; // pre-roped Q for the SUCC phase ++ result->src[6] = q_inter; // pre-roped Q for the INTER phase ++ result->src[7] = mask_succ; // opencoti F5 dca (#617B): dense SUCC mask (NULL ⇒ analytical fast path) ++ result->src[8] = mask_inter; // opencoti F5 dca (#617B): dense INTER mask (NULL ⇒ analytical fast path) ++ ++ return result; ++} ++ + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. Walsh–Hadamard + // rotation op for the fused turbo FA-VEC path (Q forward-rotate / FA-output + // inverse-rotate). B-adapted vs the AtomicBot fork: `scale` is normally NULL and +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 +@@ -376,6 +376,11 @@ extern "C" { + // opencoti F5 W3 #292 fused-MoE — see docs/features/advanced_kv.md + // true = fuse decode MoE up+gate+GLU into one op. false = off (default). + bool fused_moe_up_gate; ++ // opencoti F5 dca — Dual Chunk Attention training-free long-context (docs/features/gemma4_dca.md) ++ // dca_enabled: false = off (default). dca_chunk_size: 0 = auto. dca_yarn_factor: 1.0 = none. ++ bool dca_enabled; ++ uint32_t dca_chunk_size; ++ float dca_yarn_factor; + uint32_t n_rs_seq; // number of recurrent-state snapshots per seq for rollback (0 = no rollback) [EXPERIMENTAL] + int32_t n_threads; // number of threads to use for generation + int32_t n_threads_batch; // number of threads to use for batch processing +diff --git a/llama.cpp/src/dca.cpp b/llama.cpp/src/dca.cpp +new file mode 100644 +--- /dev/null ++++ b/llama.cpp/src/dca.cpp +@@ -0,0 +1,327 @@ ++#include "dca.h" ++ ++#include "ggml.h" ++#include "llama-graph.h" ++#include "llama-kv-cache.h" ++#include "llama-kv-cache-iswa.h" ++ ++#include ++#include ++#include ++ ++// opencoti F5 dca (#593, T87.pD) — Section C host-fill + input builder. Design + rationale: ++// docs/features/gemma4_dca.md (Section C). This file lands the DCA 3-regime position-remap INPUT ++// (built once per ubatch). The build_attn_dca graph bodies that consume it rope the small Q per ++// regime and run the fused single-pass analytical-band FA over the PRE-roped cached K (#617: K is ++// roped once at insert, not here; #618: the legacy 3x flash_attn_ext_lse + host LSE-combine fallback ++// was removed). The host-side per-cell fill lives on the kv-cache (set_input_dca) because only it can ++// read v_cells positions; set_input() delegates. ++ ++uint32_t dca_resolve_chunk_size(const llama_cparams & cparams, const llama_hparams & hparams) { ++ if (cparams.dca_chunk_size != 0) { ++ return cparams.dca_chunk_size; ++ } ++ // auto: the model's original/pretrain window (so no intra/succ q-k distance exceeds it). ++ uint32_t c = cparams.n_ctx_orig_yarn; ++ if (c == 0) { ++ c = hparams.n_ctx_train; ++ } ++ GGML_ASSERT(c >= 1 && "dca: could not resolve a positive chunk size"); ++ return c; ++} ++ ++// opencoti F5 dca (#623): qwen35/qwen35moe carry the MROPE bit (rope_type IMROPE=40); the DCA per-regime ++// position tensors are then 4-wide (section split) and rope via ggml_rope_multi. NEOX/NORMAL archs are ++// 1-wide via ggml_rope_ext. See the build_dca_rope rationale block below. ++static inline bool dca_is_mrope(int rope_type) { ++ return (rope_type & GGML_ROPE_TYPE_MROPE) != 0; ++} ++// Position ids per token the DCA input must allocate/fill: 4 (MROPE section split) or 1 (NEOX scalar). ++static inline int dca_pos_per_token(int rope_type) { ++ return dca_is_mrope(rope_type) ? GGML_MROPE_SECTIONS : 1; ++} ++ ++void llm_graph_input_dca::set_input(const llama_ubatch * ubatch) { ++ if (!mctx) { ++ return; ++ } ++ // Delegate to the kv-cache: it owns v_cells, the only source of per-cell cached positions. ++ mctx->set_input_dca( ++ pos_k, ++ pos_q[DCA_REGIME_INTRA], pos_q[DCA_REGIME_SUCC], pos_q[DCA_REGIME_INTER], ++ mask[DCA_REGIME_INTRA], mask[DCA_REGIME_SUCC], mask[DCA_REGIME_INTER], ++ pos_q_real, // O6-fix: real abs query positions for the analytical band (null when not analytical) ++ ubatch, chunk_size); ++} ++ ++bool llm_graph_input_dca::can_reuse(const llm_graph_params & params) { ++ // pos_k length tracks the (growing) KV count and the per-regime masks are position-specific, so ++ // the DCA input is rebuilt every ubatch (conservative; set_input re-fills regardless). A future ++ // optimization could reuse when n_kv + n_tokens shapes match, like can_reuse_kq_mask. ++ (void) params; ++ return false; ++} ++ ++// Allocate the DCA input tensors and register the input. mctx_kv = the BASE kv-cache context: for ++// Gemma-4 globals that is build_attn_inp_kv_iswa()->mctx->get_base(); for Qwen it is the plain kv ++// context. n_kv / n_tokens follow build_attn_inp_kq_mask (single-stream; DCA asserts n_stream==1). ++llm_graph_input_dca * llm_graph_context::build_attn_inp_dca(const llama_kv_cache_context * mctx_kv) const { ++ const uint32_t c = dca_resolve_chunk_size(cparams, hparams); ++ ++ auto inp = std::make_unique(hparams, cparams, mctx_kv, c); ++ ++ const int64_t n_kv = mctx_kv->get_n_kv(); // used for the per-regime mask shapes below ++ ++ // opencoti F5 dca (#617B): is this ubatch's cache non-contiguous (cell index != absolute position, ++ // post context-shift / self-extend)? Read the STICKY latch at graph-build time. When true on the ++ // multi-chunk (analytical) path we additionally allocate the dense SUCC/INTER fallback masks so the ++ // host fill + fused kernel can select keys by actual position instead of the (now-invalid) index ++ // bands. ALWAYS false on the hot path (prefill, decode within n_ctx, all evals) ⇒ no extra masks, ++ // no extra VRAM, kernel gets nullptr ⇒ #617 analytical fast path stays byte-identical. ++ bool fragmented = mctx_kv->get_is_fragmented(); ++ ++ // opencoti F5 dca (#623): MROPE/IMROPE archs (qwen35/qwen35moe) need 4 position ids per token so the ++ // per-regime Q ropes through ggml_rope_multi with the model's section split; NEOX archs use 1. The ++ // mask + pos_q_real shapes are position-COUNT based (n_tokens), unaffected — only pos_q widens. ++ const int64_t npp = dca_pos_per_token(rope_type); ++ ++ // opencoti F5 dca perf (T87.pD-gen #617): pos_k (regime-invariant key positions, mod c) is NO ++ // LONGER allocated. The cache now stores PRE-roped K (roped once at insert in build_attn_dca), so ++ // build_attn_dca_core reads it directly and nothing consumes pos_k. Leaving it nullptr avoids an ++ // orphan graph input (registered-but-unread => null buffer => set_input writes null->data — the ++ // same DCA_INTRA_ONLY segfault class guarded below). set_input_dca null-skips the pos_k fill. ++ ++ static const char * rname[DCA_REGIME_COUNT] = { "dca_mask_intra", "dca_mask_succ", "dca_mask_inter" }; ++ for (int r = 0; r < DCA_REGIME_COUNT; ++r) { ++ // opencoti F5 dca perf (T87.pD-opt Stage 1): single-chunk serving (n_ctx <= c) only ever uses ++ // the INTRA regime — every query is in chunk 0 so SUCC/INTER masks are all -inf. Leave their ++ // pos_q/mask tensors null so they are NOT registered as graph inputs: a registered-but-unread ++ // input gets a null buffer from the allocator and set_input_dca would write to null->data ++ // (the DCA_INTRA_ONLY segfault). build_attn_dca_core runs the INTRA-only fast path to match. ++ if (inp->single_chunk && r != DCA_REGIME_INTRA) { ++ continue; ++ } ++ inp->pos_q[r] = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens * npp); ++ ggml_set_input(inp->pos_q[r]); ++ ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): SUCC/INTER under analytical_bands need NO dense ++ // mask — the CUDA _lse op derives their contiguous KV band on-GPU from pos_q + chunk. Leave ++ // mask[r]/mask_cnv[r] null (the FA call passes mask=nullptr for them), killing the host ++ // O(n_kv*n_tokens) fill + multi-GB H2D for 2 of 3 regimes. pos_q[r] above is still required: it ++ // ropes Q for the regime AND seeds the on-GPU band. INTRA always keeps its native-causal mask. ++ // opencoti F5 dca (#617B): on the FRAGMENTED multi-chunk path we DO allocate the SUCC/INTER ++ // dense masks (index-based bands are invalid when index != pos) so set_input_dca can fill them ++ // position-based and build_attn_dca_core feeds them to the fused kernel. Contiguous cache ++ // (the always-case) keeps skipping them ⇒ identical to #618/#617 fast path. ++ if (inp->analytical_bands && r != DCA_REGIME_INTRA && !fragmented) { ++ continue; ++ } ++ ++ // additive mask [n_kv, n_tokens, 1, 1] — single-stream. Mirrors build_attn_inp_kq_mask's ++ // unpadded shape; cast to F16 for the flash-attn op when FA is enabled. ++ inp->mask[r] = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, n_kv, n_tokens, 1, 1); ++ ggml_set_input(inp->mask[r]); ++ ggml_set_name(inp->mask[r], rname[r]); ++ ++ inp->mask_cnv[r] = cparams.flash_attn ++ ? ggml_cast(ctx0, inp->mask[r], GGML_TYPE_F16) ++ : inp->mask[r]; ++ } ++ ++ // opencoti F5 dca O6-fix (>=3 chunks): the analytical band needs the REAL absolute query position ++ // (cq = real_pos/chunk). The per-regime pos_q the kernel received was RoPE-remapped (succ->oq+c, ++ // inter->c const) so cq==1 for every query at chunk>=2 — SUCC collapsed to the chunk-0 band, INTER ++ // emptied. Allocate one regime-independent real-position channel; only the multi-chunk analytical ++ // path reads it (the SUCC + INTER _lse ops). single_chunk / dense-mask paths leave it nullptr (no ++ // orphan input — set_input_dca null-guards the fill). ++ if (inp->analytical_bands) { ++ inp->pos_q_real = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens); ++ ggml_set_input(inp->pos_q_real); ++ ggml_set_name(inp->pos_q_real, "dca_pos_q_real"); ++ } ++ ++ return (llm_graph_input_dca *) res->add_input(std::move(inp)); ++} ++ ++// Shared DCA attention core. The caller has stored PRE-roped K (#617: roped once at insert) + un-roped ++// V to the cache (mctx_cur). q_cur is PRE-rope (post-norm) [n_embd_head, n_head, n_tokens]. Reads the ++// cached K directly, ropes the small Q per regime, and runs the fused single-pass analytical-band FA ++// (multi-chunk) or plain FA (single_chunk) before the wo projection. The legacy 3x flash_attn_ext_lse ++// + host LSE-combine path was removed (#618). See docs/features/gemma4_dca.md Section C. ++ ++// opencoti F5 dca (T87.pD-qwen35 #623): DCA remaps the query position per regime, but qwen35/qwen35moe ++// carry rope_type = GGML_ROPE_TYPE_IMROPE (=40, MROPE bit set), whose rope op asserts 4 ids/token ++// (a->ne[2]*4 == b->ne[0]) and assigns INTERLEAVED per-section frequencies across head dims. The early ++// NEOX-collapse (rope MROPE as plain NEOX with a scalar position) shape-fixed the assert but produced ++// GARBAGE: IMRoPE's interleaved freq->dim map is NOT equal to NEOX even when every section sits at the ++// same text position (the single-chunk ±DCA gate diverged at char 21 into a token loop). The faithful ++// path below ropes the DCA-remapped position through ggml_rope_multi with the model's own section split ++// — for TEXT all spatial sections sit at the temporal position (sections 0-2 = remapped pos, section 3 ++// = 0), exactly the layout llm_graph_input_pos::set_input builds for native MROPE text. NEOX/NORMAL ++// archs (gemma4/qwen2/qwen3) keep ggml_rope_ext (1-wide pos) — (rope_type & MROPE)==0. The ++// dca_is_mrope / dca_pos_per_token predicates this relies on are defined near the top of the file. ++ ++// Rope the DCA per-regime Q (or insert-K) faithfully for the arch. `pos` must be the width ++// build_attn_inp_dca allocated: 4*n_tokens for MROPE (section split below), n_tokens for NEOX. ++ggml_tensor * llm_graph_context::build_dca_rope(ggml_tensor * x, ggml_tensor * pos, const dca_rope & rp) const { ++ if (dca_is_mrope(rope_type)) { ++ int sections[GGML_MROPE_SECTIONS] = { rp.sections[0], rp.sections[1], rp.sections[2], rp.sections[3] }; ++ return ggml_rope_multi(ctx0, x, pos, rp.freq_factors, ++ rp.n_rot, sections, rope_type, n_ctx_orig, rp.freq_base, rp.freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow); ++ } ++ return ggml_rope_ext(ctx0, x, pos, rp.freq_factors, ++ rp.n_rot, rope_type, n_ctx_orig, rp.freq_base, rp.freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow); ++} ++ ++ggml_tensor * llm_graph_context::build_attn_dca_core( ++ const llama_kv_cache_context * mctx_cur, ++ llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ++ ggml_tensor * q_cur, ++ const dca_rope & rp, ++ float kq_scale, int il) const { ++ const float max_bias = hparams.f_max_alibi_bias; ++ const float softcap = hparams.attn_soft_cap ? hparams.f_attn_logit_softcapping : 0.0f; ++ ++ // (1) read the cached K. opencoti F5 dca perf (T87.pD-gen #617): K is now PRE-ROPED at insert ++ // (rope-at-(pos mod c), applied once to the small new-token K in build_attn_dca below) and ++ // kept correct under context-shift by the DCA-aware modular K-shift. So read it DIRECTLY — no ++ // per-step full-cache rope. That re-rope of the entire cache every graph eval was the ++ // O(n_kv)/token decode bottleneck (gen halving per context-doubling). get_k returns ++ // [n_embd_head, n_head_kv, n_kv, 1]. DCA requires an f16 K cache (no q8_0/turbo on this path). ++ ggml_tensor * k_dca = mctx_cur->get_k(ctx0, il); ++ GGML_ASSERT(k_dca->type == GGML_TYPE_F16 && "dca: requires an f16 K cache (no q8_0/turbo KV)"); ++ ++ // (2) read V once; permute K + V to the flash-attn layout [n_embd_head, seq, n_head, n_stream]. ++ ggml_tensor * v = mctx_cur->get_v(ctx0, il); ++ const bool v_trans = v->nb[1] > v->nb[2]; ++ ggml_tensor * vp = ggml_permute(ctx0, v, 0, 2, 1, 3); ++ if (v_trans) { ++ vp = ggml_transpose(ctx0, vp); ++ } ++ if (vp->type == GGML_TYPE_F32) { ++ vp = ggml_cast(ctx0, vp, GGML_TYPE_F16); ++ } ++ ggml_tensor * kp = ggml_permute(ctx0, k_dca, 0, 2, 1, 3); ++ ++ // (3) Q-rope + attention. Two sub-paths (the legacy non-analytical multi-chunk fallback was removed ++ // in #618 — its precondition is enforced at context init): ++ // single_chunk → plain optimized FA (INTRA covers all keys, ordinary positions ⇒ exact full ++ // attention; GGML_PREC_F32 mirrors the native global-layer call). ++ // analytical bands → ONE fused single-pass FA (T87.pD Branch B): the 3 pre-roped Q variants feed ++ // a kernel that walks the partitioned causal range in INTER→SUCC→INTRA key- ++ // phases under one continuous online-softmax — NO per-regime ops, NO host LSE ++ // combine. analytical_bands == !single_chunk; requires cq uniform per ubatch ++ // (chunk_size % n_ubatch == 0), enforced at context init. ++ ggml_tensor * cur; ++ if (inp_dca->single_chunk) { ++ ggml_tensor * qr = build_dca_rope(q_cur, inp_dca->get_pos_q(DCA_REGIME_INTRA), rp); ++ ggml_tensor * qp = ggml_permute(ctx0, qr, 0, 2, 1, 3); ++ ggml_tensor * O = ggml_flash_attn_ext(ctx0, qp, kp, vp, ++ inp_dca->get_mask(DCA_REGIME_INTRA), kq_scale, max_bias, softcap); ++ ggml_flash_attn_ext_set_prec(O, GGML_PREC_F32); ++ cur = O; // [DV, H, NQ, 1] normalized ++ } else if (inp_dca->analytical_bands) { ++ // pre-rope the 3 Q variants (K/V already roped once above); permute to FA layout. ++ ggml_tensor * qp[DCA_REGIME_COUNT]; ++ for (int r = 0; r < DCA_REGIME_COUNT; ++r) { ++ ggml_tensor * qr = build_dca_rope(q_cur, inp_dca->get_pos_q((dca_regime) r), rp); ++ qp[r] = ggml_permute(ctx0, qr, 0, 2, 1, 3); ++ } ++ // one fused pass: INTER→SUCC→INTRA key-phases; INTRA mask bounds the diagonal; pos_q_real ++ // (REAL abs query pos) is the band channel the kernel reads to derive cq = real_pos/chunk. ++ // opencoti F5 dca (#617B): get_mask(SUCC)/get_mask(INTER) are nullptr on the contiguous fast ++ // path (analytical bands) and the dense position-based F16 masks when the cache is fragmented; ++ // when non-null the kernel switches each phase to full-range + dense-mask selection. ++ ggml_tensor * P = ggml_flash_attn_ext_dca_fused(ctx0, ++ qp[DCA_REGIME_INTRA], qp[DCA_REGIME_SUCC], qp[DCA_REGIME_INTER], ++ kp, vp, inp_dca->get_mask(DCA_REGIME_INTRA), ++ inp_dca->get_mask(DCA_REGIME_SUCC), inp_dca->get_mask(DCA_REGIME_INTER), ++ inp_dca->get_pos_q_real(), ++ kq_scale, max_bias, softcap, (int) inp_dca->chunk_size); ++ // dst is [DV+1, H, NQ, 1] (dense O block then UNUSED lse row); slice the dense O block. ++ const int64_t DV = P->ne[0] - 1; ++ const int64_t H = P->ne[1]; ++ const int64_t NQ = P->ne[2]; ++ const size_t ts = P->nb[0]; ++ cur = ggml_cont(ctx0, ggml_view_4d(ctx0, P, DV, H, NQ, 1, ++ DV*ts, DV*H*ts, DV*H*NQ*ts, 0)); ++ } else { ++ // opencoti F5 dca (#618): the legacy non-analytical multi-chunk path (3 per-regime ++ // ggml_flash_attn_ext_lse ops at regime 0 + a host f32 online-softmax combine) was REMOVED. It ++ // was a CUDA-only fallback for chunk_size % n_ubatch != 0, never ran in production, and degraded ++ // output. That precondition is now enforced at context init (llama_context::llama_context ++ // throws), so analytical_bands == !single_chunk and this branch is unreachable — abort ++ // defensively if the invariant is ever violated. ++ GGML_ABORT("dca: non-analytical multi-chunk reached — chunk_size must be a multiple of n_ubatch " ++ "(enforced at context init; legacy fallback removed, #618)"); ++ } ++ cb(cur, "kqv_out", il); ++ ++ // (5) merge heads -> [n_embd, n_tokens] and apply the output projection. ++ cur = ggml_reshape_2d(ctx0, cur, cur->ne[0]*cur->ne[1], cur->ne[2]); ++ if (wo) { ++ cur = build_lora_mm(wo, cur); ++ } ++ if (wo_b) { ++ cur = ggml_add(ctx0, cur, wo_b); ++ } ++ return cur; ++} ++ ++// iswa overload — Gemma-4 GLOBAL (non-SWA) layers. DCA never runs on SWA layers (local window <= c ++// already, no chunking needed). Stores PRE-rope K + un-roped V to the BASE cache; has_kv==false ++// shared-KV global layers pass k_cur==v_cur==nullptr and just read the earlier layer's cache. ++ggml_tensor * llm_graph_context::build_attn_dca( ++ llm_graph_input_attn_kv_iswa * inp, ++ llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ++ ggml_tensor * q_cur, ggml_tensor * k_cur, ggml_tensor * v_cur, ++ const dca_rope & rp, float kq_scale, int il) const { ++ GGML_ASSERT(!hparams.is_swa(il) && "dca: only global (non-SWA) layers route through build_attn_dca"); ++ ++ const auto * mctx_cur = inp->mctx->get_base(); ++ ++ ggml_build_forward_expand(gf, q_cur); ++ if (k_cur) { ggml_build_forward_expand(gf, k_cur); } ++ if (v_cur) { ggml_build_forward_expand(gf, v_cur); } ++ ++ // opencoti F5 dca perf (T87.pD-gen #617): PRE-ROPE the small new-token K once, HERE, with the ++ // regime-invariant key positions pos mod c, then store the roped K. pos_q[INTRA] is filled with ++ // exactly ubatch->pos[i] % c (set_input_dca), which equals get_pos_k() for a freshly inserted cell ++ // — so the cache holds bit-identically what build_attn_dca_core's (now dropped) full-cache K-rope ++ // would have produced. Decode then reads pre-roped K with no O(n_kv) per-step rope. ++ if (k_cur) { ++ ggml_tensor * k_ins = build_dca_rope(k_cur, inp_dca->get_pos_q(DCA_REGIME_INTRA), rp); ++ ggml_build_forward_expand(gf, mctx_cur->cpy_k(ctx0, k_ins, inp->get_k_idxs(), il)); ++ } ++ if (v_cur) { ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, inp->get_v_idxs(), il)); } ++ ++ return build_attn_dca_core(mctx_cur, inp_dca, wo, wo_b, q_cur, rp, kq_scale, il); ++} ++ ++// plain-kv overload — Qwen (all layers are full attention). ++ggml_tensor * llm_graph_context::build_attn_dca( ++ llm_graph_input_attn_kv * inp, ++ llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ++ ggml_tensor * q_cur, ggml_tensor * k_cur, ggml_tensor * v_cur, ++ const dca_rope & rp, float kq_scale, int il) const { ++ const auto * mctx_cur = inp->mctx; ++ ++ ggml_build_forward_expand(gf, q_cur); ++ if (k_cur) { ggml_build_forward_expand(gf, k_cur); } ++ if (v_cur) { ggml_build_forward_expand(gf, v_cur); } ++ ++ // opencoti F5 dca perf (T87.pD-gen #617): PRE-ROPE the new-token K once at insert (pos mod c via ++ // pos_q[INTRA]); see the iswa overload above for the full rationale. Decode reads pre-roped K. ++ if (k_cur) { ++ ggml_tensor * k_ins = build_dca_rope(k_cur, inp_dca->get_pos_q(DCA_REGIME_INTRA), rp); ++ ggml_build_forward_expand(gf, mctx_cur->cpy_k(ctx0, k_ins, inp->get_k_idxs(), il)); ++ } ++ if (v_cur) { ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, inp->get_v_idxs(), il)); } ++ ++ return build_attn_dca_core(mctx_cur, inp_dca, wo, wo_b, q_cur, rp, kq_scale, il); ++} +diff --git a/llama.cpp/src/dca.h b/llama.cpp/src/dca.h +new file mode 100644 +--- /dev/null ++++ b/llama.cpp/src/dca.h +@@ -0,0 +1,122 @@ ++#pragma once ++ ++// opencoti F5 dca (#593, T87.pD) — Dual Chunk Attention: training-free long-context. ++// Design + rationale: docs/features/gemma4_dca.md (Section C). This header declares the ++// graph input that carries the 3-regime DCA position remap (built once per ubatch, host-filled ++// from the batch + KV-cell positions) plus the small helpers shared by the two build_attn_dca ++// overloads (in dca.cpp). The LSE-emitting flash-attn op it feeds is ggml_flash_attn_ext_lse ++// (Section D, ggml.h). The build_attn_dca methods themselves are declared on llm_graph_context ++// in llama-graph.h (so they can reach build_lora_mm / ctx0 / gf like the other build_attn ops). ++ ++#include "llama-graph.h" // llm_graph_input_i, llm_graph_params ++#include "llama-hparams.h" ++#include "llama-cparams.h" ++ ++#include ++ ++class llama_kv_cache_context; ++ ++// DCA position regimes. The three keep-masks partition the causal lower-triangle exactly, so the ++// LSE-merge of the three per-regime softmaxes equals full attention under DCA-remapped positions. ++// intra: query/key in the same chunk — pos_q = i mod c (distance < c) ++// succ : key in the immediately-prior chunk — pos_q = (i mod c) + c (distance in (0, 2c)) ++// inter: key in any older chunk — pos_q = c (constant) (distance = c - ok) ++// All three share pos_k[j] = cellpos(j) mod c, so the (large) cached K is roped ONCE; only the ++// (small) Q is roped per regime. See docs/features/gemma4_dca.md "rope-K-ONCE, rope-Q-THRICE". ++enum dca_regime { ++ DCA_REGIME_INTRA = 0, ++ DCA_REGIME_SUCC = 1, ++ DCA_REGIME_INTER = 2, ++ DCA_REGIME_COUNT = 3, ++}; ++ ++// Per-layer rope parameters threaded into build_attn_dca. Section B fills these from the model ++// (freq_base = model.get_rope_freq_base(cparams, il), freq_scale likewise, n_rot = hparams.n_rot(il), ++// freq_factors = model.layers[il].rope_freqs). The global YaRN knobs (ext/attn/beta_*) + rope_type + ++// n_ctx_orig come from the llm_graph_context directly, matching the per-arch model rope call. ++struct dca_rope { ++ float freq_base = 0.0f; ++ float freq_scale = 1.0f; ++ int n_rot = 0; ++ ggml_tensor * freq_factors = nullptr; ++ // opencoti F5 dca (#623) — MROPE/IMROPE section split (hparams.rope_sections), e.g. [11,11,10,0] ++ // for qwen35/qwen35moe. Used only when rope_type carries the MROPE bit (build_dca_rope picks ++ // ggml_rope_multi); NEOX/NORMAL archs leave it {0,0,0,0} and rope through ggml_rope_ext. ++ int sections[4] = { 0, 0, 0, 0 }; ++}; ++ ++// Resolve the effective chunk length c: cparams.dca_chunk_size when non-zero, else (auto) the ++// model's original/pretrain context window (n_ctx_orig_yarn, falling back to n_ctx_train). c must ++// be >= 1; callers gate on cparams.dca_enabled before constructing the input. ++uint32_t dca_resolve_chunk_size(const llama_cparams & cparams, const llama_hparams & hparams); ++ ++// Per-ubatch DCA position input. Lives on the BASE kv-cache context (the non-SWA one for Gemma-4 ++// global layers via build_attn_inp_kv_iswa()->get_base(), or the plain kv context for Qwen). ++class llm_graph_input_dca : public llm_graph_input_i { ++public: ++ llm_graph_input_dca( ++ const llama_hparams & hparams, ++ const llama_cparams & cparams, ++ const llama_kv_cache_context * mctx, ++ uint32_t chunk_size) : ++ hparams(hparams), ++ cparams(cparams), ++ mctx(mctx), ++ chunk_size(chunk_size), ++ single_chunk(cparams.n_ctx <= chunk_size), ++ // opencoti F5 dca (#618): the multi-chunk path is ALWAYS the fused analytical-band kernel, so ++ // analytical_bands is simply !single_chunk. Its precondition — chunk_size % n_ubatch == 0, so ++ // every query in a ubatch shares cq = pos/c and each regime's KV band is one contiguous range ++ // (no per-element causal, no chunk-straddling tile) — is enforced at context init ++ // (llama_context::llama_context throws otherwise), so !single_chunk ⇒ analytical holds. The old ++ // dense-mask "legacy" 3x flash_attn_ext_lse + HOST f32 online-softmax combine fallback was ++ // REMOVED (#618): it never ran in prod, was CUDA-only (not a portability path), and degraded ++ // output. Future non-CUDA backends port the fused kernel, not the legacy host combine. ++ analytical_bands(!(cparams.n_ctx <= chunk_size)) { ++ } ++ ~llm_graph_input_dca() = default; ++ ++ void set_input(const llama_ubatch * ubatch) override; ++ ++ bool can_reuse(const llm_graph_params & params) override; ++ ++ // Regime-invariant key positions: pos_k[j] = cellpos(j) mod chunk_size. I32 [n_kv]. ++ ggml_tensor * get_pos_k() const { return pos_k; } ++ // Per-regime query positions. I32 [n_tokens]. ++ ggml_tensor * get_pos_q(dca_regime r) const { return pos_q[r]; } ++ // O6-fix (>=3 chunks): REAL absolute query positions (= ubatch->pos[i]), regime-independent. ++ // Passed as the analytical-band side-channel (NOT for RoPE) so the kernel derives the TRUE ++ // cq = real_pos/chunk. get_pos_q(r) is RoPE-remapped (succ -> oq+c, inter -> c const) and yields ++ // cq==1 for every query at chunk>=2 -> SUCC band collapses to chunk 0, INTER band empties. I32 [n_tokens]. ++ ggml_tensor * get_pos_q_real() const { return pos_q_real; } ++ // Per-regime additive mask, already cast to the flash-attn dtype. [n_kv, pad(n_tps), 1, n_stream]. ++ ggml_tensor * get_mask(dca_regime r) const { return mask_cnv[r]; } ++ ++ ggml_tensor * pos_k = nullptr; // I32 [n_kv] ++ ggml_tensor * pos_q [DCA_REGIME_COUNT] = { nullptr, nullptr, nullptr }; // I32 [n_tokens] ++ ggml_tensor * pos_q_real = nullptr; // I32 [n_tokens] real abs query pos (band channel, O6-fix) ++ ggml_tensor * mask [DCA_REGIME_COUNT] = { nullptr, nullptr, nullptr }; // F32 [n_kv, pad, 1, n_stream] ++ ggml_tensor * mask_cnv [DCA_REGIME_COUNT] = { nullptr, nullptr, nullptr }; // F16 (flash-attn) or = mask ++ ++ const llama_kv_cache_context * mctx = nullptr; ++ ++ // copies (graph-reuse safety, mirrors llm_graph_input_attn_kv) — see llama-graph.h note. ++ const llama_hparams hparams; ++ const llama_cparams cparams; ++ const uint32_t chunk_size; ++ ++ // opencoti F5 dca perf (T87.pD-opt Stage 1): n_ctx <= chunk_size => every query is in chunk 0, ++ // so SUCC/INTER regimes are always empty and the 3-regime result == INTRA-only (plain attention). ++ // When set, build_attn_inp_dca allocates only the INTRA pos_q/mask (SUCC/INTER stay nullptr — NOT ++ // orphan graph inputs), set_input_dca fills only INTRA, and build_attn_dca_core runs one regime ++ // and skips the LSE combine. Eliminates 2 of 3 flash-attn + lse-recompute passes per global layer. ++ const bool single_chunk; ++ ++ // opencoti F5 dca perf (T87.pD-opt O6 Tier B): when true, the SUCC/INTER regimes use ANALYTICAL ++ // per-tile KV band-bounds (computed on-GPU from pos_q + chunk) instead of a dense [n_kv x n_tokens] ++ // mask — killing 2 of 3 host O(n_kv*n_tokens) mask builds + their multi-GB H2D. build_attn_inp_dca ++ // then skips the SUCC/INTER mask allocation (pos_q stays — it ropes Q AND seeds the bounds), and ++ // build_attn_dca_core passes mask=nullptr + pos_q + chunk + regime to ggml_flash_attn_ext_lse for ++ // those two. INTRA always keeps its native-causal mask (the irreducible native-equivalent cost). ++ const bool analytical_bands; ++}; +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 +@@ -10,6 +10,7 @@ + #include "llama-mmap.h" + #include "llama-model.h" + #include "llama-ext.h" ++#include "dca.h" // opencoti F5 dca (#618): dca_resolve_chunk_size for the chunk % n_ubatch == 0 guard + #include "llama.h" + + #include +@@ -18,6 +19,12 @@ + #include + #include + ++#include ++#include ++#include ++#include ++#include ++ + // + // llama_context + // +@@ -74,6 +81,10 @@ llama_context::llama_context( + cparams.neo_pipeline_mode = params.neo_pipeline_mode; + // opencoti F5 W3 #292 fused-MoE — see docs/features/advanced_kv.md + cparams.fused_moe_up_gate = params.fused_moe_up_gate; ++ // opencoti F5 dca — see docs/features/gemma4_dca.md ++ cparams.dca_enabled = params.dca_enabled; ++ cparams.dca_chunk_size = params.dca_chunk_size; ++ cparams.dca_yarn_factor = params.dca_yarn_factor; + + cparams.n_threads = params.n_threads; + cparams.n_threads_batch = params.n_threads_batch; +@@ -199,6 +210,21 @@ llama_context::llama_context( + + cparams.n_ubatch = std::min(cparams.n_batch, params.n_ubatch == 0 ? params.n_batch : params.n_ubatch); + ++ // opencoti F5 dca (#618): the legacy non-analytical multi-chunk DCA fallback was removed, so DCA now ++ // REQUIRES chunk_size % n_ubatch == 0 — the fused analytical-band kernel needs every query in a ++ // ubatch to share cq = pos/chunk (one contiguous KV band per regime). Enforce it here: n_ubatch is ++ // final, and the effective chunk size is resolvable (dca_enabled/dca_chunk_size set above, ++ // n_ctx_orig_yarn at ~:81). Fail loud rather than silently degrade output (the removed fallback ++ // produced malformed generations on a chunk%ubatch != 0 config). ++ if (cparams.dca_enabled) { ++ const uint32_t dca_c = dca_resolve_chunk_size(cparams, hparams); ++ if (cparams.n_ubatch == 0 || dca_c % cparams.n_ubatch != 0) { ++ throw std::runtime_error( ++ "DCA: dca chunk size (" + std::to_string(dca_c) + ") must be a multiple of n_ubatch (" + ++ std::to_string(cparams.n_ubatch) + "); set -ub/--ubatch-size to a divisor of the chunk size"); ++ } ++ } ++ + cparams.op_offload = params.op_offload; + cparams.kv_unified = params.kv_unified; + +@@ -3385,6 +3411,9 @@ llama_context_params llama_context_default_params() { + /*.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 ++ /*.dca_enabled =*/ false, // opencoti F5 dca — off by default ++ /*.dca_chunk_size =*/ 0, // opencoti F5 dca — 0 = auto (<- n_ctx_orig_yarn) ++ /*.dca_yarn_factor =*/ 1.0f, // opencoti F5 dca — 1.0 = no YaRN stretch + /*.n_rs_seq =*/ 0, + /*.n_threads =*/ GGML_DEFAULT_N_THREADS, // TODO: better default + /*.n_threads_batch =*/ GGML_DEFAULT_N_THREADS, +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 +@@ -48,6 +48,14 @@ struct llama_cparams { + // true = build_moe_ffn collapses the up/gate expert matmuls + GLU into one + // GGML_OP_MOE_FUSED_UP_GATE op at decode (n_tokens==1). false = unfused path. + bool fused_moe_up_gate; ++ // opencoti F5 dca — Dual Chunk Attention training-free long-context (docs/features/gemma4_dca.md). ++ // dca_enabled: route full-attention layers (Gemma4 global / Qwen all) through build_attn_dca ++ // (cache K un-roped, 3-regime rope-on-read + streaming_combine_n). false = off (default). ++ // dca_chunk_size: chunk length c; 0 = auto (<- n_ctx_orig_yarn pretrain window). ++ // dca_yarn_factor: YaRN mscale stretch composed with DCA; 1.0 = none. ++ bool dca_enabled; ++ uint32_t dca_chunk_size; ++ float dca_yarn_factor; + uint32_t n_rs_seq; // number of recurrent-state snapshots per seq for rollback + int32_t n_threads; // number of threads to use for generation + int32_t n_threads_batch; // number of threads to use for batch processing +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 +@@ -107,7 +107,13 @@ bool llm_graph_input_embd::can_reuse(const llm_graph_params & params) { + } + + void llm_graph_input_pos::set_input(const llama_ubatch * ubatch) { +- if (ubatch->pos && pos) { ++ // opencoti F5 dca #623: in a DCA-on graph the full-attention layers rope Q from inp_dca's own ++ // per-regime positions (build_attn_dca), so the standard inp_pos is consumed by no node on archs ++ // where EVERY attention layer is DCA-routed (qwen2/qwen3/qwen3moe; qwen35/qwen35moe whose SSM ++ // layers are positionless). galloc then leaves `pos` with a NULL buffer and the unguarded ++ // ggml_backend_tensor_set below asserts buf!=NULL (ggml-backend.cpp:297). Guard on a live buffer. ++ // gemma4-iswa is unaffected: its SWA layers still rope inp_pos, so pos is always consumed. ++ if (ubatch->pos && pos && pos->buffer) { + const int64_t n_tokens = ubatch->n_tokens; + + if (ubatch->token && n_pos_per_embd == 4) { +@@ -454,16 +460,22 @@ void llm_graph_input_attn_no_cache::set_input(const llama_ubatch * ubatch) { + } + + void llm_graph_input_attn_kv::set_input(const llama_ubatch * ubatch) { +- mctx->set_input_k_idxs(self_k_idxs, ubatch); +- mctx->set_input_v_idxs(self_v_idxs, ubatch); ++ // opencoti F5 dca #623: when DCA diverts the full-attention layers to build_attn_dca (which builds ++ // its own dca_mask_* and pre-ropes K at insert), the standard K/V write-index and K-shift rot ++ // tensors are consumed by no compute node, so galloc leaves them with a NULL buffer. Writing them ++ // then asserts in set_input_k_idxs's is_host(dst->buffer) (bug-404). Guard every idx/rot setter on ++ // a live buffer — exactly the iswa (gemma4) path's guard. (set_input_kq_mask self-guards on a null ++ // buffer.) Non-DCA graphs: the tensors are always consumed, so the guards are no-ops. ++ if (self_k_idxs && self_k_idxs->buffer) { mctx->set_input_k_idxs(self_k_idxs, ubatch); } ++ if (self_v_idxs && self_v_idxs->buffer) { mctx->set_input_v_idxs(self_v_idxs, ubatch); } + + mctx->set_input_kq_mask(self_kq_mask, ubatch, cparams.causal_attn); + +- if (self_k_rot) { ++ if (self_k_rot && self_k_rot->buffer) { + mctx->set_input_k_rot(self_k_rot); + } + +- if (self_v_rot) { ++ if (self_v_rot && self_v_rot->buffer) { + mctx->set_input_v_rot(self_v_rot); + } + } +@@ -593,16 +605,21 @@ void llm_graph_input_attn_cross::set_input(const llama_ubatch * ubatch) { + } + + void llm_graph_input_mem_hybrid::set_input(const llama_ubatch * ubatch) { +- mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch); +- mctx->get_attn()->set_input_v_idxs(inp_attn->self_v_idxs, ubatch); ++ // opencoti F5 dca #623: same orphan-buffer guard as the iswa path. DCA diverts the hybrid's ++ // full-attention layers to build_attn_dca, leaving the standard idx/rot tensors unconsumed ++ // (NULL buffer); the unguarded setters then assert is_host(NULL). Guard each on a live buffer. ++ // (kq_mask self-guards.) qwen35moe routes EVERY full-attention layer through DCA, so unlike ++ // gemma4-iswa (SWA layers keep the masks live) nothing else consumes these — the guard is load-bearing. ++ if (inp_attn->self_k_idxs && inp_attn->self_k_idxs->buffer) { mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch); } ++ if (inp_attn->self_v_idxs && inp_attn->self_v_idxs->buffer) { mctx->get_attn()->set_input_v_idxs(inp_attn->self_v_idxs, ubatch); } + + mctx->get_attn()->set_input_kq_mask(inp_attn->self_kq_mask, ubatch, cparams.causal_attn); + +- if (inp_attn->self_k_rot) { ++ if (inp_attn->self_k_rot && inp_attn->self_k_rot->buffer) { + mctx->get_attn()->set_input_k_rot(inp_attn->self_k_rot); + } + +- if (inp_attn->self_v_rot) { ++ if (inp_attn->self_v_rot && inp_attn->self_v_rot->buffer) { + mctx->get_attn()->set_input_v_rot(inp_attn->self_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 +@@ -27,6 +27,9 @@ class llama_memory_recurrent_context; + class llama_memory_hybrid_context; + class llama_memory_hybrid_iswa_context; + ++class llm_graph_input_dca; // opencoti F5 dca #593 (defined in dca.h) ++struct dca_rope; // opencoti F5 dca #593 (defined in dca.h) ++ + // certain models (typically multi-modal) can produce different types of graphs + enum llm_graph_type { + LLM_GRAPH_TYPE_DEFAULT, +@@ -998,6 +1001,34 @@ struct llm_graph_context { + + llm_graph_input_attn_kv * build_attn_inp_kv() const; + ++ // opencoti F5 dca (#593) — DCA 3-regime position-remap input (src/dca.cpp). mctx_kv = the ++ // BASE kv-cache context (iswa get_base() for Gemma-4 globals / the plain kv context for Qwen). ++ llm_graph_input_dca * build_attn_inp_dca(const llama_kv_cache_context * mctx_kv) const; ++ ++ // opencoti F5 dca (#593) — DCA attention (3-regime chunked-position) for one layer (src/dca.cpp). ++ // Two overloads mirror the build_attn input types: iswa (Gemma-4 globals) + plain-kv (Qwen). ++ // q_cur and k_cur are PRE-rope (Section B passes them un-roped); rope happens per regime inside. ++ ggml_tensor * build_attn_dca( ++ llm_graph_input_attn_kv_iswa * inp, llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ++ ggml_tensor * q_cur, ggml_tensor * k_cur, ggml_tensor * v_cur, ++ const dca_rope & rp, float kq_scale, int il) const; ++ ggml_tensor * build_attn_dca( ++ llm_graph_input_attn_kv * inp, llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ++ ggml_tensor * q_cur, ggml_tensor * k_cur, ggml_tensor * v_cur, ++ const dca_rope & rp, float kq_scale, int il) const; ++ // shared core: K/V already stored by the caller; mctx_cur is the BASE kv-cache context. ++ ggml_tensor * build_attn_dca_core( ++ const llama_kv_cache_context * mctx_cur, llm_graph_input_dca * inp_dca, ++ ggml_tensor * wo, ggml_tensor * wo_b, ggml_tensor * q_cur, ++ const dca_rope & rp, float kq_scale, int il) const; ++ // opencoti F5 dca (#623) — rope the DCA per-regime Q (or the insert-K) with the op that matches ++ // the arch: ggml_rope_multi (MROPE/IMROPE — qwen35/qwen35moe; pos is 4-wide, model section split) ++ // or ggml_rope_ext (NEOX/NORMAL — gemma4/qwen2/qwen3; pos 1-wide). `pos` MUST be the width ++ // build_attn_inp_dca allocated for this arch (n_tokens or 4*n_tokens). See src/dca.cpp. ++ ggml_tensor * build_dca_rope(ggml_tensor * x, ggml_tensor * pos, const dca_rope & rp) const; ++ + ggml_tensor * build_attn( + llm_graph_input_attn_kv * inp, + ggml_tensor * wo, +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 +@@ -4,6 +4,7 @@ + #include "llama-io.h" + #include "llama-model.h" + #include "llama-context.h" ++#include "dca.h" // opencoti F5 dca (#593) — dca_resolve_chunk_size for the DCA-aware K-shift + + #include + #include +@@ -2142,6 +2143,20 @@ bool llama_kv_cache::get_has_shift() const { + return result; + } + ++// opencoti F5 dca (#617B): true if ANY stream's cells have been position-relabeled in place (a ++// classic context-shift / self-extend), so cell index != absolute position. The DCA host-fill reads ++// this at graph-build time to gate the position-based dense-mask fallback (the analytical band ++// selection is index-based and silently wrong on a non-contiguous cache). See llama_kv_cells. ++bool llama_kv_cache::get_is_fragmented() const { ++ bool result = false; ++ ++ for (uint32_t s = 0; s < n_stream; ++s) { ++ result |= v_cells[s].get_is_fragmented(); ++ } ++ ++ return result; ++} ++ + ggml_type llama_kv_cache::type_k() const { + // opencoti F4 M3 Phase 4: all per-stream tensors of a layer carry the + // same type — read it from stream 0. Works in both unified and +@@ -3547,6 +3562,36 @@ void llama_kv_cache::set_input_k_shift(ggml_tensor * dst) const { + } + } + ++// opencoti F5 dca (#593, T87.pD-gen #617) — modular K-shift delta for DCA global layers. They cache ++// PRE-roped K at rope-at-(pos mod c) (roped once on the new-token K at insert, in build_attn_dca). A ++// context-shift moves a cell from pos_old -> pos_new, so its K rotation must move (pos_old mod c) -> ++// (pos_new mod c). RoPE composes additively, so build_rope_shift applies rope-by-delta to the cached ++// K; for DCA the delta is the MODULAR difference (pos_new mod c) - (pos_old mod c), NOT the native ++// absolute pos_new - pos_old. Both are host-side here: get_shift(i) == pos_new - pos_old (the native ++// channel), so pos_old = pos_new - get_shift(i). Empty cells contribute 0 (no rotation). ++void llama_kv_cache::set_input_k_shift_dca(ggml_tensor * dst, uint32_t chunk_size) const { ++ GGML_ASSERT(ggml_backend_buffer_is_host(dst->buffer)); ++ const int64_t c = (int64_t) chunk_size; ++ GGML_ASSERT(c >= 1 && "dca: K-shift needs a positive chunk size"); ++ ++ int32_t * data = (int32_t *) dst->data; ++ ++ for (uint32_t s = 0; s < n_stream; ++s) { ++ const auto & cells = v_cells[s]; ++ ++ for (uint32_t i = 0; i < cells.size(); ++i) { ++ if (cells.is_empty(i)) { ++ data[s*cells.size() + i] = 0; ++ continue; ++ } ++ const llama_pos pnew = cells.pos_get(i); ++ const llama_pos pold = pnew - cells.get_shift(i); // position the cached K is currently roped at ++ // pnew, pold are >= 0, so the C++ % is non-negative; delta range is (-(c-1), c-1). ++ data[s*cells.size() + i] = (int32_t) (pnew % c) - (int32_t) (pold % c); ++ } ++ } ++} ++ + struct args_set_input_kq_mask { + const llama_hparams & hparams; + const llama_ubatch * ubatch; +@@ -3741,6 +3786,15 @@ static void set_input_kq_mask_impl(const args_set_input_kq_mask & args, float * + void llama_kv_cache::set_input_kq_mask(ggml_tensor * dst, const llama_ubatch * ubatch, bool causal_attn) const { + const uint32_t n_tokens = ubatch->n_tokens; + ++ // opencoti F5 dca #593: when DCA is active the global/full-attention layers are diverted to ++ // build_attn_dca (which builds its own dca_mask_* tensors), so the standard kq_mask for those ++ // layers (gemma4 self_kq_mask; qwen self_kq_mask) is consumed by no compute node and the graph ++ // allocator leaves it with a NULL buffer. It is a pure orphan — nothing reads it — so there is ++ // nothing to fill. Skip it. (In non-DCA graphs this never triggers: the mask is always used.) ++ if (dst->buffer == NULL) { ++ return; ++ } ++ + GGML_ASSERT(ggml_backend_buffer_is_host(dst->buffer)); + float * data = (float *) dst->data; + +@@ -3802,6 +3856,185 @@ void llama_kv_cache::set_input_pos_bucket(ggml_tensor * dst, const llama_ubatch + } + } + ++// opencoti F5 dca (#593) — host-fill the DCA 3-regime position remap. Single-stream only (a 1M ++// context is single-sequence; multi-stream is a TODO, mirroring set_input_pos_bucket). Fills: ++// pos_k[j] = cellpos(j) mod c (regime-invariant key positions; empty cell -> 0) ++// pos_q_*[i] per regime: intra=oq, succ=oq+c, inter=c (oq = pos(i) mod c) ++// mask_*[i,j] = 0 where (causal p0<=p1, same-seq, nonempty) AND chunk(j) in the regime band, ++// else -inf. The three masks partition the causal lower-triangle exactly, so the ++// LSE-merge of the three softmaxes equals full attention under DCA positions. ++void llama_kv_cache::set_input_dca( ++ ggml_tensor * pos_k, ++ ggml_tensor * pos_q_intra, ggml_tensor * pos_q_succ, ggml_tensor * pos_q_inter, ++ ggml_tensor * mask_intra, ggml_tensor * mask_succ, ggml_tensor * mask_inter, ++ ggml_tensor * pos_q_real, ++ const llama_ubatch * ubatch, uint32_t chunk_size) const { ++ GGML_ASSERT(n_stream == 1 && "dca: single-stream only (1M context is single-sequence); TODO multi-stream"); ++ GGML_ASSERT(chunk_size >= 1); ++ ++ const auto & cells = v_cells[0]; ++ ++ const int64_t n_tokens = ubatch->n_tokens; ++ // opencoti F5 dca perf (T87.pD-gen #617): pos_k is no longer allocated (K is PRE-roped at insert, ++ // build_attn_dca_core no longer ropes the cache), so derive n_kv from the INTRA mask (always ++ // present, mask_intra->ne[0] == n_kv) and skip the pos_k fill. pos_k arrives nullptr; the ++ // parameter is retained only to keep the shared set_input_dca signature stable. ++ (void) pos_k; ++ const int64_t n_kv = mask_intra->ne[0]; ++ const int64_t c = (int64_t) chunk_size; ++ ++ GGML_ASSERT(ggml_backend_buffer_is_host(mask_intra->buffer)); ++ GGML_ASSERT(mask_intra->ne[1] >= n_tokens); ++ ++ // pos_q per regime. opencoti F5 dca perf (T87.pD-opt Stage 1): in single-chunk serving ++ // (n_ctx <= c) only INTRA is allocated/consumed; SUCC/INTER tensors are null here — guard every ++ // SUCC/INTER write so set_input never touches an unallocated input's null ->data. ++ int32_t * pqi = (int32_t *) pos_q_intra->data; ++ int32_t * pqs = pos_q_succ ? (int32_t *) pos_q_succ->data : nullptr; ++ int32_t * pqn = pos_q_inter ? (int32_t *) pos_q_inter->data : nullptr; ++ // opencoti F5 dca (#623): MROPE/IMROPE archs (qwen35/qwen35moe) allocate pos_q 4-wide (section ++ // split) so the per-regime Q ropes via ggml_rope_multi; NEOX archs keep it 1-wide. Derive the per- ++ // token id count from the tensor shape and write the remapped position into sections 0-2 with ++ // section 3 = 0 — the exact layout llm_graph_input_pos::set_input builds for native MROPE TEXT. ++ // npp==1 reduces to the previous scalar write byte-identically. pos_q_real stays scalar (band channel). ++ const int64_t npp = pos_q_intra->ne[0] / n_tokens; ++ // O6-fix: real abs query positions for the analytical band (kernel: cq = real_pos/chunk = the TRUE ++ // chunk index). Regime-independent; null unless the multi-chunk analytical path allocated it. The ++ // RoPE-remapped pqs (oq+c) / pqn (c const) would pin cq==1 for every query at chunk>=2 — SUCC band ++ // collapses to chunk 0, INTER band empties — so they cannot double as the band channel. ++ int32_t * pqr = pos_q_real ? (int32_t *) pos_q_real->data : nullptr; ++ for (int64_t i = 0; i < n_tokens; ++i) { ++ const llama_pos p1 = ubatch->pos[i]; ++ const int32_t oq = (int32_t) (p1 % c); ++ for (int64_t s = 0; s < npp; ++s) { ++ const int64_t off = s * n_tokens + i; // contiguous per-section blocks (native MROPE layout) ++ const bool spatial = (s < 3); // sections 0-2 carry the position; section 3 = 0 for text ++ pqi[off] = spatial ? oq : 0; ++ if (pqs) { pqs[off] = spatial ? (oq + (int32_t) c) : 0; } ++ // opencoti DCA #635 FIX: INTER-Q position = 2c, NOT c. K is pre-roped ONCE at j%c (#617), ++ // so the INTER relative distance is q_inter - (j%c). With q_inter = c that is c-(j%c) in ++ // [1,c], which OVERLAPS the recent/SUCC window (SUCC true rel runs up to 2c-1): a chunk-end ++ // far key lands at rel~1 and masquerades as the most-recent token, so at the FIRST cq>=2 ++ // query the older chunks flood ~50% of the attention mass and the model loses the thread ++ // (verified: count-up collapses exactly at pos==2c; true-ref relerr 1.1; INTER mass 0.50). ++ // q_inter = 2c gives rel = 2c-(j%c) in [c+1,2c]: every far key is strictly BEYOND the window ++ // edge c (canonical DCA places them AT c), so they stay retrievable but cannot out-compete ++ // genuine recent keys. Offline-validated on captured vectors: INTER mass 0.50 -> 0.06, ++ // recent-argmax 0.28-0.81 -> 0.91-1.00. (This is an approximation forced by #617's ++ // pre-roped-once K, which cannot realize canonical DCA's constant rel==c for INTER.) ++ if (pqn) { pqn[off] = spatial ? (int32_t)(2 * c) : 0; } ++ } ++ if (pqr) { pqr[i] = (int32_t) p1; } // analytical band channel: always scalar, never roped ++ } ++ ++ // masks: init -inf, then open the INTRA band for causal same-seq nonempty cells. ++ float * mi = (float *) mask_intra->data; ++ ++ // opencoti F5 dca (#617B): FRAGMENTED fallback. After a classic context-shift the cache is ++ // position-relabeled in place (cell index != absolute position, with holes), so the analytical ++ // chunk bands — which assume index == pos — select the WRONG physical cells and the model ++ // collapses. build_attn_inp_dca detects this via cells.get_is_fragmented() and allocates the dense ++ // SUCC/INTER masks (mask_succ/mask_inter non-null here ⇒ this branch). Fill all THREE regimes by ++ // reading each cell's ACTUAL position over the FULL [0,n_kv): for query chunk cq=p1/c and key ++ // chunk ck=p0/c — INTRA iff ck==cq, SUCC iff ck==cq-1, INTER iff ckbuffer)); ++ GGML_ASSERT(ggml_backend_buffer_is_host(mask_inter->buffer)); ++ float * ms = (float *) mask_succ->data; ++ float * mn = (float *) mask_inter->data; ++ const int64_t ne1 = mask_intra->ne[1]; ++ for (int64_t t = 0; t < n_kv*ne1; ++t) { mi[t] = -INFINITY; ms[t] = -INFINITY; mn[t] = -INFINITY; } ++ std::vector kpos0(n_kv); // cell position, or -1 for an empty/hole cell (skip) ++ for (int64_t j = 0; j < n_kv; ++j) { ++ kpos0[j] = cells.is_empty(j) ? -1 : (int32_t) cells.pos_get(j); ++ } ++ for (int64_t i = 0; i < n_tokens; ++i) { ++ const llama_seq_id seq_id = ubatch->seq_id[i][0]; ++ const llama_pos p1 = ubatch->pos[i]; ++ const int64_t cq = (int64_t) (p1 / c); ++ const uint64_t idst = (uint64_t) n_kv * i; ++ for (int64_t j = 0; j < n_kv; ++j) { ++ const int32_t p0 = kpos0[j]; ++ if (p0 < 0) { continue; } // empty / hole ++ if (p0 > p1) { continue; } // causal ++ if (!cells.seq_has(j, seq_id)) { continue; } // same seq ++ const int64_t ck = (int64_t) (p0 / c); // ck <= cq under causal p0 <= p1 ++ if (ck == cq) { mi[idst + j] = 0.0f; } // INTRA: same chunk ++ else if (ck == cq - 1) { ms[idst + j] = 0.0f; } // SUCC: immediately-prior chunk ++ else { mn[idst + j] = 0.0f; } // INTER: any older chunk (ck < cq-1) ++ } ++ } ++ return; ++ } ++ ++ // opencoti F5 dca perf (O9 analytical-INTRA host window). Under analytical_bands (signalled by a ++ // non-null pos_q_real; SUCC/INTER masks are null and derived on-GPU) the fused kernel reads the INTRA ++ // mask ONLY in the single-chunk index band [cq*c, (cq+1)*c) — its INTRA phase starts at ++ // intra_lo = cq*dca_c_blk (fattn-mma-f16.cuh) and cq is UNIFORM per ubatch (chunk_size % n_ubatch == 0). ++ // Every row below cq*c (5/6 of n_kv at 256k, 121/122 at 1M) is filled+copied today but NEVER read. ++ // Narrowing the fill to the chunk window makes the host mask cost O(c*n_tokens) — INDEPENDENT of n_kv — ++ // so DCA-on prefill stops starving the GPU as context grows. The SKIP-maskfill A/B proved the ++ // full-width fill was the ENTIRE residual 256k bottleneck (SM-util p10 58->100, 3403->4014 t/s ~= the ++ // native-off 4034). Only the j-RANGE is narrowed to provably-unread bounds; the per-cell causal ++ // (p0<=p1) + seq_has checks are byte-identical to the full path. Keys in [cq*c,(cq+1)*c) are all in ++ // chunk cq, so only the INTRA mask is written (no ck compare). single_chunk / legacy-3-mask paths ++ // (pos_q_real==null) keep the proven full-width fill below. ++ if (pos_q_real) { ++ const int64_t cq = (int64_t) (ubatch->pos[0] / c); // uniform per ubatch under analytical bands ++ const int64_t lo = cq * c; ++ const int64_t hi = (cq + 1) * c < n_kv ? (cq + 1) * c : n_kv; ++ std::vector wpos(hi > lo ? (size_t) (hi - lo) : 0); // window cell positions (empty => -1), reused per query ++ for (int64_t j = lo; j < hi; ++j) { ++ wpos[j - lo] = cells.is_empty(j) ? -1 : (int32_t) cells.pos_get(j); ++ } ++ for (int64_t i = 0; i < n_tokens; ++i) { ++ float * const row = mi + (uint64_t) n_kv * i; ++ for (int64_t j = lo; j < hi; ++j) { row[j] = -INFINITY; } // vectorizable window memset ++ const llama_seq_id seq_id = ubatch->seq_id[i][0]; ++ const llama_pos p1 = ubatch->pos[i]; ++ for (int64_t j = lo; j < hi; ++j) { ++ const int32_t p0 = wpos[j - lo]; ++ if (p0 < 0) { continue; } // empty cell ++ if (p0 > p1) { continue; } // causal ++ if (!cells.seq_has(j, seq_id)) { continue; } // per-query seq match (cheap bitset, as native) ++ row[j] = 0.0f; // window ⊂ chunk cq ⇒ INTRA ++ } ++ } ++ return; ++ } ++ ++ // opencoti F5 dca (#618): the legacy dense 3-regime (INTRA/SUCC/INTER) full-width fill was REMOVED. ++ // This path is reached ONLY for single_chunk serving (n_ctx <= chunk ⇒ pos_q_real == null and the ++ // SUCC/INTER masks unallocated): every query and key live in chunk 0, so the partition collapses to ++ // INTRA-only with no chunk-index compare. Fill the single native-causal INTRA mask. Keep the two ++ // passes (vectorizable -inf memset, then sparse causal open) — fusing them regressed prefill ~20% ++ // (O8, reverted): a per-cell store guarded by a conditional cannot vectorize. ++ const int64_t ne1 = mask_intra->ne[1]; ++ for (int64_t t = 0; t < n_kv*ne1; ++t) { ++ mi[t] = -INFINITY; ++ } ++ std::vector kpos0(n_kv); // cell position, or -1 for an empty cell (skip) ++ for (int64_t j = 0; j < n_kv; ++j) { ++ kpos0[j] = cells.is_empty(j) ? -1 : (int32_t) cells.pos_get(j); ++ } ++ for (int64_t i = 0; i < n_tokens; ++i) { ++ const llama_seq_id seq_id = ubatch->seq_id[i][0]; ++ const llama_pos p1 = ubatch->pos[i]; ++ const uint64_t idst = (uint64_t) n_kv * i; ++ for (int64_t j = 0; j < n_kv; ++j) { ++ const int32_t p0 = kpos0[j]; ++ if (p0 < 0) { continue; } // empty cell ++ if (p0 > p1) { continue; } // causal ++ if (!cells.seq_has(j, seq_id)) { continue; } // per-query seq match (cheap bitset, as native) ++ mi[idst + j] = 0.0f; // single chunk ⇒ every causal key is INTRA ++ } ++ } ++} ++ + void llama_kv_cache::set_input_k_rot(ggml_tensor * dst) const { + GGML_ASSERT(ggml_backend_buffer_is_host(dst->buffer)); + +@@ -3917,7 +4150,14 @@ public: + + void set_input(const llama_ubatch * ubatch) override; + +- ggml_tensor * k_shift; // I32 [kv_size*n_stream] ++ ggml_tensor * k_shift = nullptr; // I32 [kv_size*n_stream] ++ ++ // opencoti F5 dca (#593, T87.pD-gen #617): parallel shift channel for DCA global layers, which ++ // cache PRE-roped K (rope-at-(pos mod c)) and so need the per-cell MODULAR rotation delta ++ // (pos_new mod c) - (pos_old mod c), not the native absolute pos_new - pos_old. Null (and unread) ++ // unless this cache owns a DCA layer. dca_chunk = the resolved chunk size c. ++ ggml_tensor * k_shift_dca = nullptr; // I32 [kv_size*n_stream] ++ uint32_t dca_chunk = 0; + + // note: assumes k_rot^2 == I + ggml_tensor * k_rot = nullptr; +@@ -3932,6 +4172,10 @@ void llm_graph_input_k_shift::set_input(const llama_ubatch * ubatch) { + kv_self->set_input_k_shift(k_shift); + } + ++ if (k_shift_dca) { ++ kv_self->set_input_k_shift_dca(k_shift_dca, dca_chunk); ++ } ++ + if (k_rot) { + kv_self->set_input_k_rot(k_rot); + } +@@ -3943,12 +4187,41 @@ ggml_cgraph * llama_kv_cache::build_graph_shift(llm_graph_result * res, llama_co + + auto inp = std::make_unique(this); + +- inp->k_shift = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, (int64_t) get_size()*n_stream); +- ggml_set_input(inp->k_shift); ++ const auto & cparams = lctx->get_cparams(); ++ ++ // opencoti F5 dca (#593, T87.pD-gen #617): the native k_shift channel is consumed only by layers ++ // that take the native ABSOLUTE shift — SWA layers, or any layer when DCA is off. When DCA is on ++ // AND every layer is full-attention (e.g. Qwen: all layers route to k_shift_dca in the loop below), ++ // the native k_shift is an ORPHAN graph input: the allocator assigns it no buffer, yet set_input ++ // would still write to it -> GGML_ASSERT(buffer) at the first context shift. Allocate it ONLY when ++ // some layer actually uses it (mirrors the has_dca_layer guard for k_shift_dca). k_shift defaults ++ // to nullptr so set_input's `if (k_shift)` guard correctly skips the fill when it's not allocated. ++ bool needs_native_shift = false; ++ for (const auto & layer : layers) { ++ if (!(cparams.dca_enabled && !hparams.is_swa(layer.il))) { needs_native_shift = true; break; } ++ } ++ if (needs_native_shift) { ++ inp->k_shift = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, (int64_t) get_size()*n_stream); ++ ggml_set_input(inp->k_shift); ++ } + + inp->k_rot = build_input_k_rot(ctx); + +- const auto & cparams = lctx->get_cparams(); ++ // opencoti F5 dca (#593, T87.pD-gen #617): DCA global layers cache PRE-roped K (rope-at-(pos mod ++ // c)), so their KV-shift needs a per-cell MODULAR delta, not the native absolute pos_new-pos_old. ++ // Allocate a parallel shift channel + record the chunk size c, but ONLY when this cache actually ++ // owns a DCA (non-SWA) layer — otherwise it'd be an orphan graph input (null buffer -> null write). ++ if (cparams.dca_enabled) { ++ bool has_dca_layer = false; ++ for (const auto & layer : layers) { ++ if (!hparams.is_swa(layer.il)) { has_dca_layer = true; break; } ++ } ++ if (has_dca_layer) { ++ inp->k_shift_dca = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, (int64_t) get_size()*n_stream); ++ ggml_set_input(inp->k_shift_dca); ++ inp->dca_chunk = dca_resolve_chunk_size(cparams, hparams); ++ } ++ } + + // opencoti F4 M3 Phase 4-C: k-shift loops over per-stream tensors. + // With per-stream owning tensors, each stream's K is a separate buffer, +@@ -3987,10 +4260,17 @@ ggml_cgraph * llama_kv_cache::build_graph_shift(llm_graph_result * res, llama_co + ggml_row_size(layer_k->type, n_embd_k_gqa), + ggml_row_size(layer_k->type, n_embd_nope)); + ++ // opencoti F5 dca (#593, T87.pD-gen #617): DCA global layers store PRE-roped K, so they ++ // shift via the modular-delta channel (k_shift_dca); SWA / non-DCA layers stay on the ++ // native absolute k_shift. The DCA channel is only allocated when this cache owns a DCA ++ // layer — exactly when the predicate below can be true. ++ ggml_tensor * k_shift_src = (cparams.dca_enabled && !hparams.is_swa(il)) ++ ? inp->k_shift_dca ++ : inp->k_shift; + ggml_tensor * k_shift_s = (n_stream == 1) +- ? inp->k_shift +- : ggml_view_1d(ctx, inp->k_shift, get_size(), +- (size_t) s * get_size() * inp->k_shift->nb[0]); ++ ? k_shift_src ++ : ggml_view_1d(ctx, k_shift_src, get_size(), ++ (size_t) s * get_size() * k_shift_src->nb[0]); + + ggml_tensor * cur = build_rope_shift(cparams, ctx, k, k_shift_s, inp->k_rot, rope_factors, freq_base_l, freq_scale_l, il); + +@@ -4803,6 +5083,10 @@ uint32_t llama_kv_cache_context::get_n_kv() const { + return n_kv; + } + ++bool llama_kv_cache_context::get_is_fragmented() const { ++ return kv->get_is_fragmented(); ++} ++ + ggml_type llama_kv_cache_context::type_k() const { + return kv->type_k(); + } +@@ -4919,6 +5203,17 @@ void llama_kv_cache_context::set_input_pos_bucket(ggml_tensor * dst, const llama + kv->set_input_pos_bucket(dst, ubatch); + } + ++// opencoti F5 dca (#593) — per-batch wrapper around the DCA host-fill. ++void llama_kv_cache_context::set_input_dca( ++ ggml_tensor * pos_k, ++ ggml_tensor * pos_q_intra, ggml_tensor * pos_q_succ, ggml_tensor * pos_q_inter, ++ ggml_tensor * mask_intra, ggml_tensor * mask_succ, ggml_tensor * mask_inter, ++ ggml_tensor * pos_q_real, ++ const llama_ubatch * ubatch, uint32_t chunk_size) const { ++ kv->set_input_dca(pos_k, pos_q_intra, pos_q_succ, pos_q_inter, ++ mask_intra, mask_succ, mask_inter, pos_q_real, ubatch, chunk_size); ++} ++ + void llama_kv_cache_context::set_input_k_rot(ggml_tensor * dst) const { + kv->set_input_k_rot(dst); + } +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 +@@ -185,6 +185,11 @@ public: + + bool get_has_shift() const; + ++ // opencoti F5 dca (#617B): true if any stream's cells have been position-relabeled in place ++ // (context-shift / self-extend) ⇒ cell index != absolute position. Read at DCA graph-build time ++ // to gate the position-based dense-mask fallback. See llama_kv_cells::get_is_fragmented. ++ bool get_is_fragmented() const; ++ + // opencoti F4 M3 Phase 1+2 — see docs/decisions/0001-lazy-slot-context.md + // Idempotent. Zeroes cells [n_cells_cleared, up_to_cells) of every KV + // layer's K and V tensors across every stream. The buffers are malloc'd +@@ -307,9 +312,17 @@ public: + void set_input_v_idxs(ggml_tensor * dst, const llama_ubatch * ubatch, const slot_info & sinfo) const; + + void set_input_k_shift(ggml_tensor * dst) const; ++ // opencoti F5 dca (#593, T87.pD-gen #617) — modular K-shift delta for DCA global layers (which ++ // cache PRE-roped K at pos mod c). Fills dst[i] = (pos_new % c) - (pos_old % c) per cell. ++ void set_input_k_shift_dca(ggml_tensor * dst, uint32_t chunk_size) const; + + void set_input_kq_mask (ggml_tensor * dst, const llama_ubatch * ubatch, bool causal_attn) const; + void set_input_pos_bucket(ggml_tensor * dst, const llama_ubatch * ubatch) const; ++ // opencoti F5 dca (#593) — DCA 3-regime position-remap host-fill (single-stream). ++ void set_input_dca(ggml_tensor * pos_k, ggml_tensor * pos_q_intra, ggml_tensor * pos_q_succ, ++ ggml_tensor * pos_q_inter, ggml_tensor * mask_intra, ggml_tensor * mask_succ, ++ ggml_tensor * mask_inter, ggml_tensor * pos_q_real, ++ const llama_ubatch * ubatch, uint32_t chunk_size) const; + + void set_input_k_rot(ggml_tensor * dst) const; + void set_input_v_rot(ggml_tensor * dst) const; +@@ -593,6 +606,11 @@ public: + + uint32_t get_n_kv() const; + ++ // opencoti F5 dca (#617B): delegate the cache's non-contiguity latch (index != pos after a ++ // context-shift / self-extend). build_attn_inp_dca reads this at graph-build to decide whether to ++ // allocate the dense SUCC/INTER fallback masks for this ubatch. ++ bool get_is_fragmented() const; ++ + ggml_type type_k() const; + ggml_type type_v() const; + +@@ -650,6 +668,11 @@ public: + void set_input_k_shift (ggml_tensor * dst) const; + void set_input_kq_mask (ggml_tensor * dst, const llama_ubatch * ubatch, bool causal_attn) const; + void set_input_pos_bucket(ggml_tensor * dst, const llama_ubatch * ubatch) const; ++ // opencoti F5 dca (#593) — DCA 3-regime position-remap host-fill (single-stream). ++ void set_input_dca(ggml_tensor * pos_k, ggml_tensor * pos_q_intra, ggml_tensor * pos_q_succ, ++ ggml_tensor * pos_q_inter, ggml_tensor * mask_intra, ggml_tensor * mask_succ, ++ ggml_tensor * mask_inter, ggml_tensor * pos_q_real, ++ const llama_ubatch * ubatch, uint32_t chunk_size) const; + + void set_input_k_rot(ggml_tensor * dst) const; + void set_input_v_rot(ggml_tensor * dst) const; +diff --git a/llama.cpp/src/llama-kv-cells.h b/llama.cpp/src/llama-kv-cells.h +--- a/llama.cpp/src/llama-kv-cells.h ++++ b/llama.cpp/src/llama-kv-cells.h +@@ -41,6 +41,10 @@ public: + + has_shift = false; + ++ // opencoti F5 dca (#617B): the non-contiguity latch (index!=pos) clears ONLY on a full reset. ++ // reset_shift() must NOT touch it — see the member comment below. ++ is_fragmented = false; ++ + used.clear(); + + for (uint32_t s = 0; s < LLAMA_MAX_SEQ; ++s) { +@@ -96,6 +100,16 @@ public: + return has_shift; + } + ++ // opencoti F5 dca (#617B): true once any operation has relabeled a cell's position in place ++ // (pos_add / pos_div, i.e. a classic context-shift or self-extend), breaking the cell-index == ++ // absolute-position invariant that DCA's analytical band selection depends on. STICKY across ++ // reset_shift() (which only clears the transient rope-delta queue, leaving pos[] relabeled) — ++ // cleared ONLY on a full reset(). The DCA host-fill gates a position-based dense-mask fallback on ++ // this so a shifted (non-contiguous) cache stays correct; a never-shifted cache keeps the fast path. ++ bool get_is_fragmented() const { ++ return is_fragmented; ++ } ++ + // move cell isrc to idst (used during defrag) + //void mv(uint32_t isrc, uint32_t idst) { + // assert(isrc < pos.size()); +@@ -420,6 +434,7 @@ public: + shift[i] += d; + + has_shift = true; ++ is_fragmented = true; // opencoti F5 dca (#617B): pos[i] relabeled in place ⇒ index != pos + + if (pos[i] < 0) { + seq[i].reset(); +@@ -453,11 +468,20 @@ public: + seq_pos_add(i); + + has_shift = true; ++ is_fragmented = true; // opencoti F5 dca (#617B): pos[i] relabeled in place ⇒ index != pos + } + + private: + bool has_shift = false; + ++ // opencoti F5 dca (#617B): sticky "cell index != absolute position somewhere" latch. Set true by ++ // pos_add / pos_div (the only ops that relabel a live cell's position in place — context-shift's ++ // seq_add and self-extend's seq_div); NOT cleared by reset_shift() (positions stay relabeled after ++ // the rope-shift applies), only by reset(). Plain rm / seq_rm are NOT flagged: a removal leaves ++ // surviving cells at pos==index (holes are just empty slots the dense mask skips), and prompt-cache ++ // reuse calls seq_rm every request — flagging it would needlessly force the slow path on the hot flow. ++ bool is_fragmented = false; ++ + // set of indices of used cells (i.e. pos[i] != -1, allowed to not have any seq_id) + std::set used; + +diff --git a/llama.cpp/src/models/models.h b/llama.cpp/src/models/models.h +--- a/llama.cpp/src/models/models.h ++++ b/llama.cpp/src/models/models.h +@@ -1747,6 +1747,7 @@ struct llama_model_qwen35 : public llama_model_base { + private: + ggml_tensor * build_layer_attn( + llm_graph_input_attn_kv * inp_attn, ++ llm_graph_input_dca * inp_dca, // opencoti F5 dca #623 (nullptr => --dca off, unchanged) + ggml_tensor * cur, + ggml_tensor * inp_pos, + int * sections, +@@ -1793,6 +1794,7 @@ struct llama_model_qwen35moe : public llama_model_base { + private: + ggml_tensor * build_layer_attn( + llm_graph_input_attn_kv * inp_attn, ++ llm_graph_input_dca * inp_dca, // opencoti F5 dca #623 (nullptr => --dca off, unchanged) + ggml_tensor * cur, + ggml_tensor * inp_pos, + int * sections, +diff --git a/llama.cpp/src/models/qwen2.cpp b/llama.cpp/src/models/qwen2.cpp +--- a/llama.cpp/src/models/qwen2.cpp ++++ b/llama.cpp/src/models/qwen2.cpp +@@ -1,4 +1,5 @@ + #include "models.h" ++#include "dca.h" // opencoti F5 dca #593 + + void llama_model_qwen2::load_arch_hparams(llama_model_loader & ml) { + ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); +@@ -65,6 +66,11 @@ llama_model_qwen2::graph::graph(const llama_model & model, const llm_graph_param + + auto * inp_attn = build_attn_inp_kv(); + ++ // opencoti F5 dca #593 — Qwen: every layer is full attention, so DCA applies to all of ++ // them. Build the 3-regime position remap once on the plain-kv context. ++ const bool use_dca = cparams.dca_enabled; ++ llm_graph_input_dca * inp_dca = use_dca ? build_attn_inp_dca(inp_attn->mctx) : nullptr; ++ + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + for (int il = 0; il < n_layer; ++il) { +@@ -82,25 +88,35 @@ llama_model_qwen2::graph::graph(const llama_model & model, const llm_graph_param + auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur, + n_embd_head, n_head, n_head_kv, il); + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope Q; rope per regime in build_attn_dca + Qcur = ggml_rope_ext( + ctx0, Qcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope K + Kcur = ggml_rope_ext( + ctx0, Kcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + + cb(Qcur, "Qcur", il); + cb(Kcur, "Kcur", il); + cb(Vcur, "Vcur", il); + ++ if (use_dca) { // opencoti F5 dca #593 ++ const dca_rope rp = { (float) freq_base, (float) freq_scale, (int) n_rot, nullptr }; ++ cur = build_attn_dca(inp_attn, inp_dca, model.layers[il].wo, model.layers[il].wo_b, ++ Qcur, Kcur, Vcur, rp, 1.0f/sqrtf(float(n_embd_head)), il); ++ } else { + cur = build_attn(inp_attn, + model.layers[il].wo, model.layers[il].wo_b, model.layers[il].wo_s, + Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il); ++ } + } + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); +diff --git a/llama.cpp/src/models/qwen3.cpp b/llama.cpp/src/models/qwen3.cpp +--- a/llama.cpp/src/models/qwen3.cpp ++++ b/llama.cpp/src/models/qwen3.cpp +@@ -1,4 +1,5 @@ + #include "models.h" ++#include "dca.h" // opencoti F5 dca #593 + + void llama_model_qwen3::load_arch_hparams(llama_model_loader & ml) { + ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); +@@ -65,6 +66,11 @@ llama_model_qwen3::graph::graph(const llama_model & model, const llm_graph_param + + auto * inp_attn = build_attn_inp_kv(); + ++ // opencoti F5 dca #593 — Qwen: every layer is full attention, so DCA applies to all of ++ // them. Build the 3-regime position remap once on the plain-kv context. ++ const bool use_dca = cparams.dca_enabled; ++ llm_graph_input_dca * inp_dca = use_dca ? build_attn_inp_dca(inp_attn->mctx) : nullptr; ++ + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + for (int il = 0; il < n_layer; ++il) { +@@ -85,28 +91,38 @@ llama_model_qwen3::graph::graph(const llama_model & model, const llm_graph_param + Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, NULL, LLM_NORM_RMS, il); + cb(Qcur, "Qcur_normed", il); + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope Q; rope per regime in build_attn_dca + Qcur = ggml_rope_ext( + ctx0, Qcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + + Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, NULL, LLM_NORM_RMS, il); + cb(Kcur, "Kcur_normed", il); + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope K + Kcur = ggml_rope_ext( + ctx0, Kcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + + cb(Qcur, "Qcur", il); + cb(Kcur, "Kcur", il); + cb(Vcur, "Vcur", il); + ++ if (use_dca) { // opencoti F5 dca #593 ++ const dca_rope rp = { (float) freq_base, (float) freq_scale, (int) n_rot, nullptr }; ++ cur = build_attn_dca(inp_attn, inp_dca, model.layers[il].wo, model.layers[il].wo_b, ++ Qcur, Kcur, Vcur, rp, 1.0f/sqrtf(float(n_embd_head)), il); ++ } else { + cur = build_attn(inp_attn, + model.layers[il].wo, model.layers[il].wo_b, model.layers[il].wo_s, + Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il); ++ } + } + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); +diff --git a/llama.cpp/src/models/qwen35.cpp b/llama.cpp/src/models/qwen35.cpp +--- a/llama.cpp/src/models/qwen35.cpp ++++ b/llama.cpp/src/models/qwen35.cpp +@@ -1,4 +1,5 @@ + #include "models.h" ++#include "dca.h" // opencoti F5 dca #623 — qwen35 (Qwen3.5-Next hybrid) DCA on the full-attn layers + #include "llama-memory-recurrent.h" + + void llama_model_qwen35::load_arch_hparams(llama_model_loader & ml) { +@@ -154,6 +155,12 @@ llama_model_qwen35::graph::graph(const llama_model & model, const llm_graph_para + + auto * inp = build_inp_mem_hybrid(); + ++ // opencoti F5 dca #623 — DCA applies ONLY to the full-attention layers (is_recurrent==false); ++ // the gated-delta-net (SSM) layers are recurrent and never chunked. Build the 3-regime position ++ // remap once on the attn kv-context. inp_dca==nullptr (default --dca off) => unchanged behaviour. ++ const bool use_dca = cparams.dca_enabled; ++ llm_graph_input_dca * inp_dca = use_dca ? build_attn_inp_dca(inp->get_attn()->mctx) : nullptr; ++ + ggml_tensor * inp_pos = build_inp_pos(); + ggml_tensor * inp_out_ids = build_inp_out_ids(); + +@@ -173,7 +180,7 @@ llama_model_qwen35::graph::graph(const llama_model & model, const llm_graph_para + cur = build_layer_attn_linear(inp->get_recr(), cur, il); + } else { + // Full attention layer +- cur = build_layer_attn(inp->get_attn(), cur, inp_pos, sections, il); ++ cur = build_layer_attn(inp->get_attn(), inp_dca, cur, inp_pos, sections, il); + } + + if (il == n_transformer_layers - 1 && inp_out_ids && cparams.embeddings_pre_norm_masked) { +@@ -259,6 +266,7 @@ ggml_tensor * llama_model_qwen35::graph::build_norm_gated( + + ggml_tensor * llama_model_qwen35::graph::build_layer_attn( + llm_graph_input_attn_kv * inp, ++ llm_graph_input_dca * inp_dca, // opencoti F5 dca #623 (nullptr => --dca off) + ggml_tensor * cur, + ggml_tensor * inp_pos, + int * sections, +@@ -301,29 +309,41 @@ ggml_tensor * llama_model_qwen35::graph::build_layer_attn( + + Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head, n_head_kv, n_tokens); + +- // Apply MRoPE +- Qcur = ggml_rope_multi( +- ctx0, Qcur, inp_pos, nullptr, +- n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, +- ext_factor, attn_factor, beta_fast, beta_slow +- ); +- +- Kcur = ggml_rope_multi( +- ctx0, Kcur, inp_pos, nullptr, +- n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, +- ext_factor, attn_factor, beta_fast, beta_slow +- ); +- +- cb(Qcur, "Qcur", il); +- cb(Kcur, "Kcur", il); +- cb(Vcur, "Vcur", il); +- + // Attention computation + const float kq_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale; + +- cur = build_attn(inp, +- nullptr, nullptr, nullptr, +- Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); ++ if (inp_dca) { ++ // opencoti F5 dca #623 — DCA path: cache PRE-rope Q/K and rope per-regime INSIDE build_attn_dca ++ // (so SKIP the MRoPE here). wo=nullptr returns the pre-projection attention output [n_embd, ++ // n_tokens] (build_attn_dca_core skips the proj on null wo) — exactly like the stock gated path, ++ // so the gate_sigmoid multiply + wo below run unchanged. Pass the model's IMRoPE section split so ++ // build_dca_rope ropes the DCA-remapped position via ggml_rope_multi with 4-wide positions — the ++ // NEOX-collapse produced garbage (gate diverged @char 21, #623). DCA needs an f16 K cache (default). ++ const dca_rope rp = { (float) freq_base, (float) freq_scale, (int) n_rot, nullptr, ++ { sections[0], sections[1], sections[2], sections[3] } }; ++ cur = build_attn_dca(inp, inp_dca, nullptr, nullptr, Qcur, Kcur, Vcur, rp, kq_scale, il); ++ } else { ++ // Apply MRoPE ++ Qcur = ggml_rope_multi( ++ ctx0, Qcur, inp_pos, nullptr, ++ n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow ++ ); ++ ++ Kcur = ggml_rope_multi( ++ ctx0, Kcur, inp_pos, nullptr, ++ n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow ++ ); ++ ++ cb(Qcur, "Qcur", il); ++ cb(Kcur, "Kcur", il); ++ cb(Vcur, "Vcur", il); ++ ++ cur = build_attn(inp, ++ nullptr, nullptr, nullptr, ++ Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); ++ } + cb(cur, "attn_pregate", il); + + ggml_tensor * gate_sigmoid = ggml_sigmoid(ctx0, gate); +diff --git a/llama.cpp/src/models/qwen35moe.cpp b/llama.cpp/src/models/qwen35moe.cpp +--- a/llama.cpp/src/models/qwen35moe.cpp ++++ b/llama.cpp/src/models/qwen35moe.cpp +@@ -1,4 +1,5 @@ + #include "models.h" ++#include "dca.h" // opencoti F5 dca #623 — qwen35moe (Qwen3.5-Next hybrid MoE) DCA on full-attn layers + #include "llama-memory-recurrent.h" + + void llama_model_qwen35moe::load_arch_hparams(llama_model_loader & ml) { +@@ -177,6 +178,12 @@ llama_model_qwen35moe::graph::graph(const llama_model & model, const llm_graph_p + + auto * inp = build_inp_mem_hybrid(); + ++ // opencoti F5 dca #623 — DCA applies ONLY to the full-attention layers (is_recurrent==false); ++ // the gated-delta-net (SSM) layers are recurrent and never chunked. Build the 3-regime position ++ // remap once on the attn kv-context. inp_dca==nullptr (default --dca off) => unchanged behaviour. ++ const bool use_dca = cparams.dca_enabled; ++ llm_graph_input_dca * inp_dca = use_dca ? build_attn_inp_dca(inp->get_attn()->mctx) : nullptr; ++ + ggml_tensor * inp_pos = build_inp_pos(); + ggml_tensor * inp_out_ids = build_inp_out_ids(); + +@@ -196,7 +203,7 @@ llama_model_qwen35moe::graph::graph(const llama_model & model, const llm_graph_p + cur = build_layer_attn_linear(inp->get_recr(), cur, il); + } else { + // Full attention layer +- cur = build_layer_attn(inp->get_attn(), cur, inp_pos, sections, il); ++ cur = build_layer_attn(inp->get_attn(), inp_dca, cur, inp_pos, sections, il); + } + + if (il == n_transformer_layers - 1 && inp_out_ids && cparams.embeddings_pre_norm_masked) { +@@ -282,6 +289,7 @@ ggml_tensor * llama_model_qwen35moe::graph::build_norm_gated( + + ggml_tensor * llama_model_qwen35moe::graph::build_layer_attn( + llm_graph_input_attn_kv * inp, ++ llm_graph_input_dca * inp_dca, // opencoti F5 dca #623 (nullptr => --dca off) + ggml_tensor * cur, + ggml_tensor * inp_pos, + int * sections, +@@ -324,29 +332,41 @@ ggml_tensor * llama_model_qwen35moe::graph::build_layer_attn( + + Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head, n_head_kv, n_tokens); + +- // Apply IMRoPE +- Qcur = ggml_rope_multi( +- ctx0, Qcur, inp_pos, nullptr, +- n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, +- ext_factor, attn_factor, beta_fast, beta_slow +- ); +- +- Kcur = ggml_rope_multi( +- ctx0, Kcur, inp_pos, nullptr, +- n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, +- ext_factor, attn_factor, beta_fast, beta_slow +- ); +- +- cb(Qcur, "Qcur", il); +- cb(Kcur, "Kcur", il); +- cb(Vcur, "Vcur", il); +- + // Attention computation + const float kq_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale; + +- cur = build_attn(inp, +- nullptr, nullptr, nullptr, +- Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); ++ if (inp_dca) { ++ // opencoti F5 dca #623 — DCA path: cache PRE-rope Q/K and rope per-regime INSIDE build_attn_dca ++ // (so SKIP the IMRoPE here). wo=nullptr returns the pre-projection attention output [n_embd, ++ // n_tokens] (build_attn_dca_core skips the proj on null wo) — exactly like the stock gated path, ++ // so the gate_sigmoid multiply + wo below run unchanged. Pass the model's IMRoPE section split so ++ // build_dca_rope ropes the DCA-remapped position via ggml_rope_multi with 4-wide positions — the ++ // NEOX-collapse produced garbage (gate diverged @char 21, #623). DCA needs an f16 K cache (default). ++ const dca_rope rp = { (float) freq_base, (float) freq_scale, (int) n_rot, nullptr, ++ { sections[0], sections[1], sections[2], sections[3] } }; ++ cur = build_attn_dca(inp, inp_dca, nullptr, nullptr, Qcur, Kcur, Vcur, rp, kq_scale, il); ++ } else { ++ // Apply IMRoPE ++ Qcur = ggml_rope_multi( ++ ctx0, Qcur, inp_pos, nullptr, ++ n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow ++ ); ++ ++ Kcur = ggml_rope_multi( ++ ctx0, Kcur, inp_pos, nullptr, ++ n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, ++ ext_factor, attn_factor, beta_fast, beta_slow ++ ); ++ ++ cb(Qcur, "Qcur", il); ++ cb(Kcur, "Kcur", il); ++ cb(Vcur, "Vcur", il); ++ ++ cur = build_attn(inp, ++ nullptr, nullptr, nullptr, ++ Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); ++ } + cb(cur, "attn_pregate", il); + + ggml_tensor * gate_sigmoid = ggml_sigmoid(ctx0, gate); +diff --git a/llama.cpp/src/models/qwen3moe.cpp b/llama.cpp/src/models/qwen3moe.cpp +--- a/llama.cpp/src/models/qwen3moe.cpp ++++ b/llama.cpp/src/models/qwen3moe.cpp +@@ -1,4 +1,5 @@ + #include "models.h" ++#include "dca.h" // opencoti F5 dca #593 + + void llama_model_qwen3moe::load_arch_hparams(llama_model_loader & ml) { + ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false); +@@ -75,6 +76,11 @@ llama_model_qwen3moe::graph::graph(const llama_model & model, const llm_graph_pa + + auto * inp_attn = build_attn_inp_kv(); + ++ // opencoti F5 dca #593 — Qwen: every layer is full attention, so DCA applies to all of ++ // them. Build the 3-regime position remap once on the plain-kv context. ++ const bool use_dca = cparams.dca_enabled; ++ llm_graph_input_dca * inp_dca = use_dca ? build_attn_inp_dca(inp_attn->mctx) : nullptr; ++ + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + for (int il = 0; il < n_layer; ++il) { +@@ -95,28 +101,38 @@ llama_model_qwen3moe::graph::graph(const llama_model & model, const llm_graph_pa + Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, NULL, LLM_NORM_RMS, il); + cb(Qcur, "Qcur_normed", il); + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope Q; rope per regime in build_attn_dca + Qcur = ggml_rope_ext( + ctx0, Qcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + + Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, NULL, LLM_NORM_RMS, il); + cb(Kcur, "Kcur_normed", il); + ++ if (!use_dca) { // opencoti F5 dca #593 — cache PRE-rope K + Kcur = ggml_rope_ext( + ctx0, Kcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow + ); ++ } + + cb(Qcur, "Qcur", il); + cb(Kcur, "Kcur", il); + cb(Vcur, "Vcur", il); + ++ if (use_dca) { // opencoti F5 dca #593 ++ const dca_rope rp = { (float) freq_base, (float) freq_scale, (int) n_rot, nullptr }; ++ cur = build_attn_dca(inp_attn, inp_dca, model.layers[il].wo, model.layers[il].wo_b, ++ Qcur, Kcur, Vcur, rp, 1.0f/sqrtf(float(n_embd_head)), il); ++ } else { + cur = build_attn(inp_attn, + model.layers[il].wo, model.layers[il].wo_b, model.layers[il].wo_s, + Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il); ++ } + } + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); +diff --git a/llama.cpp/tools/server/server-context.cpp b/llama.cpp/tools/server/server-context.cpp +--- a/llama.cpp/tools/server/server-context.cpp ++++ b/llama.cpp/tools/server/server-context.cpp +@@ -889,6 +889,17 @@ private: + return false; + } + ++ // opencoti F5 dca (#618): a context-init failure (e.g. the DCA `chunk %% n_ubatch == 0` ++ // precondition, or the pre-existing `n_ctx_seq == 0` / OOM / control-vector paths) makes ++ // common_init_from_params return with a loaded model but a NULL context. Without this guard ++ // the very next line (llama_n_ctx(ctx_tgt)) dereferences null -> SIGSEGV/coredump, masking the ++ // clear init error logged above. Mirror the model-null check so any such failure exits ++ // cleanly (false) with the error already on the log, instead of crashing. ++ if (ctx_tgt == nullptr) { ++ SRV_ERR("failed to create context (see initialization error above), '%s'\n", params_base.model.path.c_str()); ++ return false; ++ } ++ + vocab = llama_model_get_vocab(model_tgt); + + n_ctx = llama_n_ctx(ctx_tgt); +@@ -1019,8 +1030,18 @@ private: + + int n_ctx_slot = llama_n_ctx_seq(ctx_tgt); + if (n_ctx_slot > n_ctx_train) { +- SRV_WRN("the slot context (%d) exceeds the training context of the model (%d) - capping\n", n_ctx_slot, n_ctx_train); +- n_ctx_slot = n_ctx_train; ++ // opencoti F5 dca (T87.pD): Dual Chunk Attention is a TRAINING-FREE long-context method — ++ // it remaps query/key positions so no rope position ever exceeds the chunk size (<= the ++ // pretrain window), letting the model serve sequences far beyond n_ctx_train. The stock ++ // guard below caps the slot at n_ctx_train (correct for vanilla models, where positions past ++ // train produce garbage), but that cap is exactly wrong under DCA and silently blocked the ++ // >256k path. When --dca is on, keep the full requested slot context. ++ if (params_base.dca_enabled) { ++ SRV_INF("the slot context (%d) exceeds the training context (%d) — allowed: DCA extends context\n", n_ctx_slot, n_ctx_train); ++ } else { ++ SRV_WRN("the slot context (%d) exceeds the training context of the model (%d) - capping\n", n_ctx_slot, n_ctx_train); ++ n_ctx_slot = n_ctx_train; ++ } + } + + slots.clear();