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();