From: opencoti Subject: [PATCH 0089] sparse-attn vslash + uncaptured foundation (#551/#557-568, #591/#592) Consolidated capture of tree work that predated the 0090 (compute-reserve) snapshot but was never emitted as a patch: the sparse-attention vertical-slash feature (ggml_sparse_attn_vslash op, sparse-attn.cu/.cuh estimate+block-mass kernels, llama-graph wiring, sparse_attn_block_size plumbed through the KV-cache ctors) plus the #591/#592 decode-overhead work (async input H2D upload, on-device greedy draft argmax) and assorted fattn/turbo/tcq touch-ups from that window. Reconstructed as diff(chain-0088 baseline, vendor-backup 20260701-123044) after bug-2108 (see buglog) showed 0090+ were captured against this uncaptured state and the committed chain no longer strict-applied from a clean 0.10.3 tree. diff --git a/llama.cpp/common/arg.cpp b/llama.cpp/common/arg.cpp index cfd3b29..89bd9dd 100644 --- a/llama.cpp/common/arg.cpp +++ b/llama.cpp/common/arg.cpp @@ -1539,6 +1539,72 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.dca_yarn_factor = std::stof(value); } ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DCA_YARN_FACTOR")); + // opencoti #551 sparse-attn — Quest-style block-selector sparse attention — see docs/features/sparse_attn.md + add_opt(common_arg( + {"--sparse-attn"}, "MODE", + "Block-selector sparse attention for long-context decode (Quest-style, training-free): off (default), on. " + "Per-query, keeps only the top-K KV blocks by an exact min/max upper bound (+ recent/sink window) and " + "skips the rest before the QK dot. Composes with --dca and quantized KV. Default off = byte-identical.", + [](common_params & params, const std::string & value) { + if (value == "off" || value == "false" || value == "0") params.sparse_attn_enabled = false; + else if (value == "on" || value == "true" || value == "1") params.sparse_attn_enabled = true; + else throw std::invalid_argument("--sparse-attn must be on|off"); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN")); + add_opt(common_arg( + {"--sparse-attn-block-size"}, "N", + "Sparse-attn KV-block granularity B_SEL in tokens (default 64). Min/max key bounds are kept per block; " + "align to a multiple of the quant block size (q5_0/q4_0 blck=32) for clean skipping.", + [](common_params & params, const std::string & value) { + params.sparse_attn_block_size = (uint32_t) std::stoul(value); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_BLOCK_SIZE")); + add_opt(common_arg( + {"--sparse-attn-topk"}, "N", + "Sparse-attn: number of KV blocks to keep per query by upper-bound score (0 = all blocks, i.e. " + "off; 'auto' = adaptive ¼ of blocks, floor 64).", + [](common_params & params, const std::string & value) { + if (value == "auto") params.sparse_attn_topk = 0xFFFFFFFFu; // adaptive: max(64, n_blocks/4) + else params.sparse_attn_topk = (uint32_t) std::stoul(value); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_TOPK")); + add_opt(common_arg( + {"--sparse-attn-recent"}, "N", + "Sparse-attn: always-keep recent window in tokens (streaming-LLM safety; 0 = none).", + [](common_params & params, const std::string & value) { + params.sparse_attn_recent = (uint32_t) std::stoul(value); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_RECENT")); + add_opt(common_arg( + {"--sparse-attn-sink"}, "N", + "Sparse-attn: number of leading 'sink' KV blocks always kept (attention-sink safety; default 1).", + [](common_params & params, const std::string & value) { + params.sparse_attn_sink = (uint32_t) std::stoul(value); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_SINK")); + add_opt(common_arg( + {"--sparse-attn-refresh"}, "N", + "Sparse-attn PERF: re-run the block selector every N decode tokens and reuse the cached selection in " + "between (amortizes the per-step selector cost; default 8, 1 = every step). The sink block is always " + "kept so a reused selection never empties.", + [](common_params & params, const std::string & value) { + params.sparse_attn_refresh = (uint32_t) std::stoul(value); + if (params.sparse_attn_refresh < 1) params.sparse_attn_refresh = 1; + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_REFRESH")); + add_opt(common_arg( + // opencoti-hook: sparse-attn-vslash — #551 selector mechanism. + {"--sparse-attn-mode"}, "MODE", + "Sparse-attn block-selector mechanism: 'quest' (default) = geometric per-dim kmin/kmax spread bound; " + "'vslash' = vertical-slash content-driven selector (per-block max attention logit over the recent " + "query window — picks the blocks the query actually attends to). vslash succeeds at long context " + "where quest's lossless fraction rises with ctx. Only meaningful with --sparse-attn on.", + [](common_params & params, const std::string & value) { + if (value == "quest" || value == "0") params.sparse_attn_mode = 0; + else if (value == "vslash" || value == "1") params.sparse_attn_mode = 1; + else throw std::invalid_argument("--sparse-attn-mode must be quest|vslash"); + } + ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPARSE_ATTN_MODE")); // 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 index 7503291..190682c 100644 --- a/llama.cpp/common/common.cpp +++ b/llama.cpp/common/common.cpp @@ -1649,6 +1649,14 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.dca_enabled = params.dca_enabled; cparams.dca_chunk_size = params.dca_chunk_size; cparams.dca_yarn_factor = params.dca_yarn_factor; + // opencoti #551 sparse-attn — see docs/features/sparse_attn.md + cparams.sparse_attn_enabled = params.sparse_attn_enabled; + cparams.sparse_attn_block_size = params.sparse_attn_block_size; + cparams.sparse_attn_topk = params.sparse_attn_topk; + cparams.sparse_attn_recent = params.sparse_attn_recent; + cparams.sparse_attn_sink = params.sparse_attn_sink; + cparams.sparse_attn_refresh = params.sparse_attn_refresh; // opencoti #551 PERF (was omitted — flag was inert) + cparams.sparse_attn_mode = params.sparse_attn_mode; // opencoti-hook: sparse-attn-vslash — #551 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 index 205ec67..b688492 100644 --- a/llama.cpp/common/common.h +++ b/llama.cpp/common/common.h @@ -469,6 +469,19 @@ struct common_params { bool dca_enabled = false; uint32_t dca_chunk_size = 0; float dca_yarn_factor = 1.0f; + // opencoti #551 sparse-attn — Quest-style block-selector sparse attention (docs/features/sparse_attn.md) + // off by default; block_size 128 (B_SEL == FA-VEC tile width nthreads; clamped to a multiple of 128 in + // llama-graph so the compacted-list consumer's tiles align to blocks); topk 0 = all blocks (off), + // 0xFFFFFFFF = auto ¼; recent 0 = no forced window (tokens); sink = leading sink blocks always kept (1). + bool sparse_attn_enabled = false; + uint32_t sparse_attn_block_size = 128; + uint32_t sparse_attn_topk = 0; + uint32_t sparse_attn_recent = 0; + uint32_t sparse_attn_sink = 1; + uint32_t sparse_attn_refresh = 8; // #551 PERF: re-run the block selector every N decode tokens (amortize; 1 = every step) + // opencoti-hook: sparse-attn-vslash — #551 selector mechanism: 0 = Quest geometric kmin/kmax bound + // (default, byte-identical to prior behavior); 1 = vertical-slash content-driven max-logit selector. + uint32_t sparse_attn_mode = 0; // opencoti F5-opt W1 (#293) pcie profile — see docs/features/rolling_kv.md // Boot-time host<->device bandwidth + PCIe link detection, consumed by M7 // Rolling KV (#296) tile sizing. autodetect reads the rebar-probe JSON / diff --git a/llama.cpp/ggml/include/ggml.h b/llama.cpp/ggml/include/ggml.h index 3c72dd4..614da7f 100644 --- a/llama.cpp/ggml/include/ggml.h +++ b/llama.cpp/ggml/include/ggml.h @@ -612,6 +612,10 @@ extern "C" { GGML_OP_FLASH_ATTN_EXT_LSE, // opencoti F5 dca — #593 (O+lse packed) + GGML_OP_SPARSE_ATTN_FILL, // opencoti-hook: sparse-attn — #551 Quest per-block kmin/kmax fill + GGML_OP_SPARSE_ATTN_SELECT, // opencoti-hook: sparse-attn — #551 Quest per-(seq,qtile) block bitmask + GGML_OP_SPARSE_ATTN_VSLASH, // opencoti-hook: sparse-attn-vslash — #551 content-driven block selector + GGML_OP_COUNT, }; @@ -2577,6 +2581,67 @@ extern "C" { int group_size, struct ggml_tensor * scale); + // opencoti-hook: sparse-attn — #551 Quest block-selector. Reduce the freshly-written + // K (pre-quant f16) into per-block channel-wise min/max, merged in-place into the + // kbounds side-cache. kbounds: [n_embd_k, n_block, 2] f16 ([..,0]=min, [..,1]=max); + // the op writes into kbounds' buffer (ggml_cpy idiom) so it carries the write edge to + // the selector. k_cur = new tokens' K; k_idxs = their cache positions. block_size = B_SEL. + // See docs/features/sparse_attn.md (S2/S3 design). + GGML_API struct ggml_tensor * ggml_sparse_attn_fill( + struct ggml_context * ctx, + struct ggml_tensor * kbounds, + struct ggml_tensor * k_cur, + struct ggml_tensor * k_idxs, + int block_size); + + // opencoti-hook: sparse-attn — #551 Quest block selector. Emits a SORTED, −1-padded compacted + // list of selected block indices per sequence (union over KV-heads): block b is kept iff any + // KV-head's Quest upper bound for b is in the top-K, OR b is in the recent/sink window. The + // FA-VEC consumer walks the list and stops at the first −1 (O(selected), not O(capacity)). + // + // #551 PERF (amortized selection): when block_sel != nullptr the op writes IN-PLACE into that + // persistent per-stream buffer ([n_block+1, n_seq] I32) and the kernel re-runs the selection + // ONLY every `refresh` decode tokens — reading the decode position from k_idxs[0] on the device + // so it stays correct under CUDA-graph replay — reusing the cached list on the in-between steps. + // The sink block is always kept, so a reused (stale) list is never empty and never OOB (the + // FA-VEC loop skips block starts >= n_kv). block_sel == nullptr ⇒ legacy transient per-step list + // (dst is a fresh [n_block, n_seq]; k_idxs/refresh ignored). op_params = + // {block_size, n_block, k_sel, recent_blocks, sink_blocks, refresh}. + GGML_API struct ggml_tensor * ggml_sparse_attn_select( + struct ggml_context * ctx, + struct ggml_tensor * q, // rotated FA query [n_embd_k, n_head, n_qtile, n_seq] + struct ggml_tensor * kbounds, // [n_embd_k, n_block, 2] from the fill op + struct ggml_tensor * block_sel, // persistent [n_block+1, n_seq] I32 write target, or NULL (transient) + struct ggml_tensor * k_idxs, // KV write indices — decode position source (NULL when block_sel NULL) + int block_size, + int n_block, + int k_sel, + int recent_blocks, + int sink_blocks, + int refresh, // re-select every N decode tokens (>=1; ignored when block_sel NULL) + int sub_per_block); // #551: kbounds sub-blocks per walk-block (finer Quest bound); 1 = none + + // opencoti-hook: sparse-attn-vslash — #551 content-driven block selector (MInference-style). Scores + // each block by the EXACT per-block max attention logit over the last `q_win` query rows + // (score(b) = max over (q in window, key in block b, head) of q·key), instead of the Quest geometric + // kmin/kmax spread bound (which over-ranks high-spread filler and degrades with context). Reuses the + // select kernel's top-K + recent/sink sorted-compaction emit verbatim; src[1] is the actual cache K + // (f16), not a kbounds side-cache. block_sel/k_idxs/refresh semantics identical to ggml_sparse_attn_select. + // op_params = {block_size, n_block, k_sel, recent_blocks, sink_blocks, refresh, q_win}. + GGML_API struct ggml_tensor * ggml_sparse_attn_vslash( + struct ggml_context * ctx, + struct ggml_tensor * q, // rotated FA query [n_embd_k, n_head, n_q, n_seq] (f32) + struct ggml_tensor * k_full, // cache K [head_dim, n_head_kv, n_kv, 1] (f16) + struct ggml_tensor * block_sel, // persistent [n_block+1, n_seq] I32 write target, or NULL (transient) + struct ggml_tensor * k_idxs, // KV write indices — decode position source (NULL when block_sel NULL) + int block_size, + int n_block, + int k_sel, + int recent_blocks, + int sink_blocks, + int refresh, // re-select every N decode tokens (>=1; ignored when block_sel NULL) + int q_win); // # of recent query rows that define the importance window (>=1) + GGML_API enum ggml_prec ggml_flash_attn_ext_get_prec( const struct ggml_tensor * a); diff --git a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c index 468f762..697bf8d 100644 --- a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c +++ b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c @@ -2124,6 +2124,9 @@ 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_SPARSE_ATTN_FILL: // opencoti-hook: sparse-attn — #551 (CUDA-only, shares abort) + case GGML_OP_SPARSE_ATTN_SELECT: // opencoti-hook: sparse-attn — #551 (CUDA-only, shares abort) + case GGML_OP_SPARSE_ATTN_VSLASH: // opencoti-hook: sparse-attn-vslash — #551 (CUDA-only, shares abort) case GGML_OP_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 (CUDA-only, shares abort) case GGML_OP_STREAMING_FLASH_ATTN: { @@ -2417,6 +2420,9 @@ 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_SPARSE_ATTN_FILL: // opencoti-hook: sparse-attn — #551 (CUDA-only) + case GGML_OP_SPARSE_ATTN_SELECT: // opencoti-hook: sparse-attn — #551 (CUDA-only) + case GGML_OP_SPARSE_ATTN_VSLASH: // opencoti-hook: sparse-attn-vslash — #551 (CUDA-only) case GGML_OP_FLASH_ATTN_EXT_LSE: // opencoti F5 dca — #593 (CUDA-only) case GGML_OP_STREAMING_FLASH_ATTN: { diff --git a/llama.cpp/ggml/src/ggml-cuda/cpy.cu b/llama.cpp/ggml/src/ggml-cuda/cpy.cu index e769b7b..7aa2140 100644 --- a/llama.cpp/ggml/src/ggml-cuda/cpy.cu +++ b/llama.cpp/ggml/src/ggml-cuda/cpy.cu @@ -602,7 +602,10 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg // set-rows.cu) is fetched by the cast's row width (ne00 = n_embd_gqa); nullptr ⇒ identity // (default-off / no slot). See docs/features/poly_kv.md (M6-S2 perf-path Level A). } else if ((src0->type == GGML_TYPE_TURBO2_0 || src0->type == GGML_TYPE_TURBO3_0 || - src0->type == GGML_TYPE_TURBO4_0 || src0->type == GGML_TYPE_TURBO8_0) && + src0->type == GGML_TYPE_TURBO4_0 || src0->type == GGML_TYPE_TURBO8_0 || + // opencoti-hook #533: TCQ→f16 dequant-on-lift (prefill-MMA + DCA-compose). The TCQ + // kernels ignore d_scale_inv (no InnerQ) — scale_inv=nullptr handled inside the kernel. + src0->type == GGML_TYPE_TURBO3_TCQ || src0->type == GGML_TYPE_TURBO2_TCQ) && src1->type == GGML_TYPE_F16) { const float * d_scale_inv = nullptr; int head_dim = 0; diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh index e5181d7..26cd885 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh @@ -26,6 +26,9 @@ typedef void (* fattn_kernel_t)( 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 + const int * __restrict__ block_sel, // opencoti-hook: sparse-attn #551 — sorted −1-padded selected-block list (nullptr=dense) + const int32_t sparse_B, // opencoti-hook: sparse-attn #551 — block size (multiple of nthreads) + const int32_t sparse_nwords, // opencoti-hook: sparse-attn #551 — list stride == n_block per sequence float * __restrict__ dst, float2 * __restrict__ dst_meta, const float scale, @@ -59,6 +62,17 @@ 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; +// opencoti-hook: sparse-attn-vslash side-output (#551 S1) — consume-once side channel for emitting the +// per-block question-window attention MASS as a FREE byproduct of the prefill MMA FA (the softmax is +// computed anyway; #551 vslash separate-pass died on prefill cost it couldn't amortize — Quest wins by +// computing its bound for free, so vslash must too). The vslash graph site sets this immediately before +// the prefill FA dispatch (mode==1, n_q>1); launch_fattn drains it. The MMA kernel (S2) accumulates +// per-block window mass into block_mass when non-null. nullptr = byte-identical legacy path. Def in fattn.cu. +extern thread_local float * opencoti_fattn_block_mass; // [n_block] out, PRE-ZEROED by the caller +extern thread_local int opencoti_fattn_block_mass_nblock; +extern thread_local int opencoti_fattn_block_mass_B; // keys per block +extern thread_local int opencoti_fattn_block_mass_qwin; // # of trailing query rows in the window + 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); @@ -1187,6 +1201,19 @@ void launch_fattn( } opencoti_fattn_dst_lse = nullptr; + // opencoti-hook: sparse-attn-vslash side-output (#551 S1) — drain the consume-once block-mass channel. + // S1 only drains (no kernel consumer yet) so this is byte-identical; S2 threads these into the kernel. + float * block_mass = opencoti_fattn_block_mass; + const int block_mass_nblock = opencoti_fattn_block_mass_nblock; + const int block_mass_B = opencoti_fattn_block_mass_B; + const int block_mass_qwin = opencoti_fattn_block_mass_qwin; + opencoti_fattn_block_mass = nullptr; + opencoti_fattn_block_mass_nblock = 0; + opencoti_fattn_block_mass_B = 0; + opencoti_fattn_block_mass_qwin = 0; + GGML_UNUSED(block_mass); GGML_UNUSED(block_mass_nblock); + GGML_UNUSED(block_mass_B); GGML_UNUSED(block_mass_qwin); + const ggml_tensor * Q = dst->src[0]; const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; @@ -1195,6 +1222,11 @@ void launch_fattn( const ggml_tensor * mask = dst->src[3]; const ggml_tensor * sinks = dst->src[4]; + // opencoti-hook: sparse-attn #551 — block_sel rides src[7] (src[5]/src[6] are the streaming/ + // position-window k_tail/v_tail). Its op_params[0]=block_size; ne[0]=n_block (the −1-padded list stride). + const ggml_tensor * block_sel = dst->src[7]; + const int sparse_B = block_sel ? ((const int32_t *) block_sel->op_params)[0] : 0; + const int sparse_nwords = block_sel ? (int) block_sel->ne[0] : 0; ggml_tensor * KQV = dst; @@ -1398,6 +1430,9 @@ void launch_fattn( } } + // opencoti #551 bug-743: parallel_blocks=1 discriminator (2026-06-26) PROVED the gridDim.y combine + // is NOT the cause — sparse collapsed to niah=0 even with combine removed. Reverted; the bug is in + // the produce side (select cnt) or per-tile walk, now under select-kernel instrumentation. blocks_num.x = ntiles_x; blocks_num.y = parallel_blocks; blocks_num.z = ntiles_z_gqa*K->ne[2]*Q->ne[3]; @@ -1440,6 +1475,8 @@ void launch_fattn( 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) + block_sel ? (const int *) block_sel->data : nullptr, // opencoti-hook: sparse-attn #551 + sparse_B, sparse_nwords, // opencoti-hook: sparse-attn #551 !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], diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh index d12da11..185eaef 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -1936,6 +1936,9 @@ static __global__ void flash_attn_ext_f16( 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 + const int * __restrict__ block_sel, // opencoti-hook: sparse-attn #551 — unused in MMA (decode uses VEC) + const int32_t sparse_B, // opencoti-hook: sparse-attn #551 + const int32_t sparse_nwords, // opencoti-hook: sparse-attn #551 float * __restrict__ dst, float2 * __restrict__ dst_meta, const float scale, @@ -2103,7 +2106,7 @@ static __global__ void flash_attn_ext_f16( (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, kb0_min); #else - GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, 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-tile.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh index c0f030b..533d5c6 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-tile.cuh @@ -795,6 +795,9 @@ static __global__ void flash_attn_tile( 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 + const int * __restrict__ block_sel, // opencoti-hook: sparse-attn #551 — unused in tile (decode uses VEC) + const int32_t sparse_B, // opencoti-hook: sparse-attn #551 + const int32_t sparse_nwords, // opencoti-hook: sparse-attn #551 float * __restrict__ dst, float2 * __restrict__ dst_meta, const float scale, @@ -820,7 +823,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, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, dst, dst_meta, scale, max_bias, m0, m1, n_head_log2, logit_softcap, ne00, ne01, ne02, ne03, nb01, nb02, nb03, @@ -1135,7 +1138,7 @@ static __global__ void flash_attn_tile( } } #else - GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, 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-turbo.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh index a4ab82b..0a2615e 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh @@ -28,18 +28,85 @@ static __constant__ float FATTN_TURBO2_CENTROIDS[4] = { -0.133462f, -0.039994f, 0.039994f, 0.133462f }; // ---- per-element rotated-value decode (centroid[code]×norm). p = position within the 128-block. ---- -static __device__ __forceinline__ float fattn_turbo3_val(const block_turbo3_0 * blk, int p, float norm) { +// opencoti-hook: #535b turbo decode smem-codebook — `cb` is the centroid table; callers pass either the +// __constant__ FATTN_TURBO*_CENTROIDS symbol (legacy/func-ptr path) or a per-block smem copy (the fast +// path, mirroring the #534 TCQ treatment) to avoid __constant__ serializing the data-dependent per-element +// lookups up to 32-way per warp. Byte-identical either way (same values/index). +static __device__ __forceinline__ float fattn_turbo3_val(const block_turbo3_0 * blk, int p, float norm, const float * __restrict__ cb) { const uint8_t lo = (blk->qs[p >> 2] >> ((p & 3) * 2)) & 0x3; const uint8_t hi = (blk->signs[p >> 3] >> (p & 7)) & 0x1; - return FATTN_TURBO3_CENTROIDS[lo | (hi << 2)] * norm; + return cb[lo | (hi << 2)] * norm; } -static __device__ __forceinline__ float fattn_turbo4_val(const block_turbo4_0 * blk, int p, float norm) { +static __device__ __forceinline__ float fattn_turbo4_val(const block_turbo4_0 * blk, int p, float norm, const float * __restrict__ cb) { const uint8_t idx = (blk->qs[p >> 1] >> ((p & 1) * 4)) & 0xF; - return FATTN_TURBO4_CENTROIDS[idx] * norm; + return cb[idx] * norm; } -static __device__ __forceinline__ float fattn_turbo2_val(const block_turbo2_0 * blk, int p, float norm) { +static __device__ __forceinline__ float fattn_turbo2_val(const block_turbo2_0 * blk, int p, float norm, const float * __restrict__ cb) { const uint8_t idx = (blk->qs[p >> 2] >> ((p & 3) * 2)) & 0x3; - return FATTN_TURBO2_CENTROIDS[idx] * norm; + return cb[idx] * norm; +} + +// opencoti-hook: fattn-vec turbo K-read vectorized loads (#535 — see docs/evaluations/context.md). +// The KQ body is global-load-latency-bound (council csl-...0805-230a): the SM stalls ~400cyc on the +// per-element qs/signs byte loads, not on compute (ncu 4.5%/4.5%) nor the __constant__ LUT. Each lane +// owns ONE block-aligned, NON-straddling group of 2*NE2 (=8) consecutive elements per outer iter, so we +// can load the whole group's qs/signs as ONE wide aligned word + hoist norm ONCE, instead of ~8 scattered +// byte loads + 4 norm reloads. BYTE-IDENTICAL to the fattn_turboN_val path (same bits → same centroid×norm +// → same MAD order); only the load width changes. ib is 2*NE2-aligned ⇒ qs+(ib>>2) is uint16-aligned +// (turbo2/3) and qs+(ib>>1) is uint32-aligned (turbo4); blk base nb (14/10/68) preserves it. NE2!=4 keeps +// the scalar fallback (non-sm_86 cpy widths). norm is the hoisted __half2float(blk->norm) for this block. +template +static __device__ __forceinline__ void fattn_turbo3_decode_group( + const block_turbo3_0 * __restrict__ blk, int ib, float norm, half2 * __restrict__ tmp, const float * __restrict__ cb) { + if constexpr (NE2 == 4) { + const uint16_t qs2 = *(const uint16_t *)(blk->qs + (ib >> 2)); // 8 × 2-bit (low), aligned + const uint8_t s1 = blk->signs[ib >> 3]; // 8 × 1-bit (sign) + #pragma unroll + for (int k1 = 0; k1 < 4; ++k1) { + const int j0 = 2*k1, j1 = j0 + 1; + tmp[k1] = make_half2( + cb[((qs2 >> (j0*2)) & 0x3) | (((s1 >> j0) & 0x1) << 2)] * norm, + cb[((qs2 >> (j1*2)) & 0x3) | (((s1 >> j1) & 0x1) << 2)] * norm); + } + } else { + #pragma unroll + for (int k1 = 0; k1 < NE2; ++k1) + tmp[k1] = make_half2(fattn_turbo3_val(blk, ib + 2*k1, norm, cb), fattn_turbo3_val(blk, ib + 2*k1 + 1, norm, cb)); + } +} +template +static __device__ __forceinline__ void fattn_turbo4_decode_group( + const block_turbo4_0 * __restrict__ blk, int ib, float norm, half2 * __restrict__ tmp, const float * __restrict__ cb) { + if constexpr (NE2 == 4) { + const uint32_t qs4 = *(const uint32_t *)(blk->qs + (ib >> 1)); // 8 × 4-bit nibbles, aligned + #pragma unroll + for (int k1 = 0; k1 < 4; ++k1) { + const int j0 = 2*k1, j1 = j0 + 1; + tmp[k1] = make_half2(cb[(qs4 >> (j0*4)) & 0xF] * norm, + cb[(qs4 >> (j1*4)) & 0xF] * norm); + } + } else { + #pragma unroll + for (int k1 = 0; k1 < NE2; ++k1) + tmp[k1] = make_half2(fattn_turbo4_val(blk, ib + 2*k1, norm, cb), fattn_turbo4_val(blk, ib + 2*k1 + 1, norm, cb)); + } +} +template +static __device__ __forceinline__ void fattn_turbo2_decode_group( + const block_turbo2_0 * __restrict__ blk, int ib, float norm, half2 * __restrict__ tmp, const float * __restrict__ cb) { + if constexpr (NE2 == 4) { + const uint16_t qs2 = *(const uint16_t *)(blk->qs + (ib >> 2)); // 8 × 2-bit, aligned + #pragma unroll + for (int k1 = 0; k1 < 4; ++k1) { + const int j0 = 2*k1, j1 = j0 + 1; + tmp[k1] = make_half2(cb[(qs2 >> (j0*2)) & 0x3] * norm, + cb[(qs2 >> (j1*2)) & 0x3] * norm); + } + } else { + #pragma unroll + for (int k1 = 0; k1 < NE2; ++k1) + tmp[k1] = make_half2(fattn_turbo2_val(blk, ib + 2*k1, norm, cb), fattn_turbo2_val(blk, ib + 2*k1 + 1, norm, cb)); + } } // Q·K MAD against the float Q slice — Q_reg is half2 with V_DOT2_F32_F16_AVAILABLE, else float2 @@ -55,23 +122,19 @@ static __device__ __forceinline__ void fattn_turbo_qmad(float & sum, half2 kval, // ===================== K: vec_dot (centroid×norm · float Q), f16-structured ===================== // K_c → key row in turbo format; Q_v → this lane's float Q slice (Q_reg, already ×scale); the dot // covers this thread's D-slice, warp_reduce_sum combines lanes (done by the caller). -#define FATTN_TURBO_KQ_BODY(BLOCK_T, VAL_FN) \ +#define FATTN_TURBO_KQ_BODY(BLOCK_T, DECODE_GROUP, CB) \ GGML_UNUSED(Q_q8); GGML_UNUSED(Q_ds_v); \ constexpr int cpy_nb = ggml_cuda_get_max_cpy_bytes(); \ constexpr int cpy_ne = cpy_nb / 4; \ float sum = 0.0f; \ _Pragma("unroll") \ for (int k_KQ_0 = 0; k_KQ_0 < D/2; k_KQ_0 += nthreads*cpy_ne) { \ + /* this lane's group = 2*cpy_ne consecutive, block-aligned (no-straddle) elements */ \ + const int e_base = 2*(k_KQ_0 + (int)(threadIdx.x % nthreads)*cpy_ne); \ + const BLOCK_T * blk = (const BLOCK_T *) K_c + (e_base >> 7); \ + const float norm = __half2float(blk->norm); /* hoisted: once per group */ \ half2 tmp[cpy_ne]; \ - _Pragma("unroll") \ - for (int k1 = 0; k1 < cpy_ne; ++k1) { \ - const int e2 = k_KQ_0 + (int)(threadIdx.x % nthreads)*cpy_ne + k1; \ - const int e = 2*e2; \ - const BLOCK_T * blk = (const BLOCK_T *) K_c + (e >> 7); \ - const float norm = __half2float(blk->norm); \ - const int ib = e & 127; \ - tmp[k1] = make_half2(VAL_FN(blk, ib, norm), VAL_FN(blk, ib + 1, norm)); \ - } \ + DECODE_GROUP(blk, e_base & 127, norm, tmp, (CB)); /* one wide qs/signs load */ \ _Pragma("unroll") \ for (int k1 = 0; k1 < cpy_ne; ++k1) { \ fattn_turbo_qmad(sum, tmp[k1], Q_v, k_KQ_0/nthreads + k1); \ @@ -79,32 +142,61 @@ static __device__ __forceinline__ void fattn_turbo_qmad(float & sum, half2 kval, } \ return sum; +// Legacy __constant__-codebook K readers (kept for the function-pointer dispatch typedef; behaviour +// unchanged — they pass the __constant__ symbol as the codebook). template static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo3( const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { - FATTN_TURBO_KQ_BODY(block_turbo3_0, fattn_turbo3_val) + FATTN_TURBO_KQ_BODY(block_turbo3_0, fattn_turbo3_decode_group, FATTN_TURBO3_CENTROIDS) } template static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo4( const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { - FATTN_TURBO_KQ_BODY(block_turbo4_0, fattn_turbo4_val) + FATTN_TURBO_KQ_BODY(block_turbo4_0, fattn_turbo4_decode_group, FATTN_TURBO4_CENTROIDS) } template static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2( const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { - FATTN_TURBO_KQ_BODY(block_turbo2_0, fattn_turbo2_val) + FATTN_TURBO_KQ_BODY(block_turbo2_0, fattn_turbo2_decode_group, FATTN_TURBO2_CENTROIDS) +} + +// opencoti-hook: #535b — smem-codebook K readers (the fast path). `cb` is the per-block smem centroid +// copy staged once at the kernel head; avoids the __constant__ per-warp serialization. Byte-identical. +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo3_smem( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v, const float * __restrict__ cb) { + FATTN_TURBO_KQ_BODY(block_turbo3_0, fattn_turbo3_decode_group, cb) +} +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo4_smem( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v, const float * __restrict__ cb) { + FATTN_TURBO_KQ_BODY(block_turbo4_0, fattn_turbo4_decode_group, cb) +} +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2_smem( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v, const float * __restrict__ cb) { + FATTN_TURBO_KQ_BODY(block_turbo2_0, fattn_turbo2_decode_group, cb) +} + +// Compile-time tier dispatch for the smem K read (mirrors tcq_vec_dot_KQ_smem in fattn-vec.cuh). +template +static __device__ __forceinline__ float turbo_vec_dot_KQ_smem( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v, const float * __restrict__ cb) { + if constexpr (type_K == GGML_TYPE_TURBO3_0) return vec_dot_fattn_vec_KQ_turbo3_smem(K_c, Q_v, Q_q8, Q_ds_v, cb); + else if constexpr (type_K == GGML_TYPE_TURBO4_0) return vec_dot_fattn_vec_KQ_turbo4_smem(K_c, Q_v, Q_q8, Q_ds_v, cb); + else return vec_dot_fattn_vec_KQ_turbo2_smem(K_c, Q_v, Q_q8, Q_ds_v, cb); } // ===================== V: dequant (centroid×norm), ne consecutive (i0 4-aligned) ===================== // Reads ne consecutive rotated V values at element offset i0; writes T (half/float) as half2/float2. -#define FATTN_TURBO_V_BODY(BLOCK_T, VAL_FN) \ +#define FATTN_TURBO_V_BODY(BLOCK_T, VAL_FN, CB) \ static_assert(ne % 2 == 0, "bad ne"); \ const BLOCK_T * blk = (const BLOCK_T *) vx + (i0 >> 7); \ const float norm = __half2float(blk->norm); \ const int ib = (int)(i0 & 127); \ float vals[ne]; \ _Pragma("unroll") \ - for (int l = 0; l < ne; ++l) vals[l] = VAL_FN(blk, ib + l, norm); \ + for (int l = 0; l < ne; ++l) vals[l] = VAL_FN(blk, ib + l, norm, (CB)); \ if constexpr (std::is_same_v) { \ half2 * d2 = (half2 *) dst; \ _Pragma("unroll") \ @@ -115,15 +207,38 @@ static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2( for (int l = 0; l < ne; l += 2) d2[l/2] = make_float2(vals[l], vals[l+1]); \ } +// Legacy __constant__-codebook V dequant (function-pointer typedef path; behaviour unchanged). template static __device__ __forceinline__ void dequantize_V_turbo3(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { - FATTN_TURBO_V_BODY(block_turbo3_0, fattn_turbo3_val) + FATTN_TURBO_V_BODY(block_turbo3_0, fattn_turbo3_val, FATTN_TURBO3_CENTROIDS) } template static __device__ __forceinline__ void dequantize_V_turbo4(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { - FATTN_TURBO_V_BODY(block_turbo4_0, fattn_turbo4_val) + FATTN_TURBO_V_BODY(block_turbo4_0, fattn_turbo4_val, FATTN_TURBO4_CENTROIDS) } template static __device__ __forceinline__ void dequantize_V_turbo2(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { - FATTN_TURBO_V_BODY(block_turbo2_0, fattn_turbo2_val) + FATTN_TURBO_V_BODY(block_turbo2_0, fattn_turbo2_val, FATTN_TURBO2_CENTROIDS) +} + +// opencoti-hook: #535b — smem-codebook V dequant (fast path). `cb` = per-block smem centroid copy. +template +static __device__ __forceinline__ void dequantize_V_turbo3_smem(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0, const float * __restrict__ cb) { + FATTN_TURBO_V_BODY(block_turbo3_0, fattn_turbo3_val, cb) +} +template +static __device__ __forceinline__ void dequantize_V_turbo4_smem(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0, const float * __restrict__ cb) { + FATTN_TURBO_V_BODY(block_turbo4_0, fattn_turbo4_val, cb) +} +template +static __device__ __forceinline__ void dequantize_V_turbo2_smem(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0, const float * __restrict__ cb) { + FATTN_TURBO_V_BODY(block_turbo2_0, fattn_turbo2_val, cb) +} + +// Compile-time tier dispatch for the smem V dequant (mirrors tcq_dequantize_V_smem in fattn-vec.cuh). +template +static __device__ __forceinline__ void turbo_dequantize_V_smem(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0, const float * __restrict__ cb) { + if constexpr (type_V == GGML_TYPE_TURBO3_0) dequantize_V_turbo3_smem(vx, dst, i0, cb); + else if constexpr (type_V == GGML_TYPE_TURBO4_0) dequantize_V_turbo4_smem(vx, dst, i0, cb); + else dequantize_V_turbo2_smem(vx, dst, i0, cb); } diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh index abcce88..949db5e 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh @@ -1,5 +1,24 @@ #include "common.cuh" #include "fattn-common.cuh" +#include // opencoti: getenv/atof for sparse-v (#546) + +// opencoti-hook: sparse-v (#546) — attention-gated V-dequant skip (theTom/turboquant_plus port). +// Per-compilation-unit copy (no device RDC); broadcast from the host in case_impl below. Default 0.0 +// ⇒ the skip predicate (weight < tau) is never true ⇒ byte-identical to the dense V-read. +static __constant__ float d_sparse_v_tau = 0.0f; +// opencoti-hook: sparse-v adaptive-τ (#565) — mass-budget knob. When >0 it OVERRIDES the static τ with a +// runtime, distribution-aware threshold θ_j = (eps/n_kv)·KQ_sum[j] (eps × mean post-exp weight per col): +// skip key k iff its normalized contribution exp(S−m)/L < eps/n_kv for EVERY query column ⇒ total dropped +// softmax mass < eps by construction, and θ auto-tracks ctx length + attention diffuseness (the long-ctx +// failure of a fixed peak-relative τ). eps==0 ⇒ static-τ path, byte-identical. +static __constant__ float d_sparse_v_eps = 0.0f; +// opencoti-hook: sparse-v combined (#565) — sink + recent ALWAYS-KEEP guards. The eps/tau skip is a +// peak/mean-referenced predicate that (a) under attention-sink wrongly drops mid-ctx peaks, and (b) drops +// recent tokens that MTP draft/verify depend on. Always keep pos=n_kv-recent +// (recent window incl. spec-draft tokens) so the skip is lossless on the needle AND preserves draft acceptance. +// Both default 0 ⇒ no protection (byte-identical to pre-#565 when eps==tau==0). +static __constant__ int d_sparse_v_sink = 0; +static __constant__ int d_sparse_v_recent = 0; static int ggml_cuda_fattn_vec_get_nthreads_host(const int cc) { return 128; @@ -10,14 +29,64 @@ static constexpr __device__ int ggml_cuda_fattn_vec_get_nthreads_device() { return 128; } +// opencoti-hook: fattn-vec occupancy lift (#533/#535) — see docs/evaluations/context.md §FA-VEC-occupancy. +// The decode FA-VEC kernel was register-bound: nvcc grabbed 255 regs/thread (the hardware ceiling — it was +// already SPILLING to local memory, STACK:80) because __launch_bounds__'s min-blocks-per-SM was 1. That caps +// the kernel at 2 blocks/SM → 16.7% theoretical / 8.3% achieved occupancy on sm_86 (48 warps/SM max), so the +// GPU sat ~90% idle and the kernel ran LATENCY-bound (ncu: compute 4.5%, mem 4.5% — neither saturated). With +// only ~4 active warps there was nothing to hide the per-key K/V read latency. Forcing a higher min-blocks +// makes nvcc cap registers (65536/(128*N)) so N blocks co-reside, raising occupancy to hide that latency. +// Helps ALL KV types, and ESPECIALLY turbo/TCQ whose centroid/codebook unpack has the longest dep-chain (the +// most latency to hide). NUMERICS ARE BYTE-IDENTICAL — this only changes block scheduling, not the math. +// Tune here (sm_86, 128 thr): 2→≤256regs/16.7%occ · 3→≤170regs/25% · 4→≤128regs/33% · 6→≤85regs/50%. +#ifndef GGML_CUDA_FATTN_VEC_MIN_BLOCKS_PER_SM +#define GGML_CUDA_FATTN_VEC_MIN_BLOCKS_PER_SM 3 +#endif +// min-blocks is D-dependent (measured ptxas register/spill curve, sm_86, turbo3 D-rungs): +// D<=256 (A4B local layers, Qwen D128, the COMMON decode path): min-blocks=3 → nvcc caps at 168 regs +// and finds a CLEAN allocation — ~0 spill (8B stack / 4B spill vs 64B/72B at min-blocks=1) AND +// occupancy 16.7%→25%. A genuine free win: more warps to hide read latency, less local-mem traffic. +// D==512 (Gemma-global / 27B head_dim): already register-thrashing at the old budget (240B spill at +// 255 regs — the root of its pathological decode cost); a tighter budget only worsens it +// (mb=3 → 756B spill / 2236B loads). No clean launch-bounds point exists for D512; it needs a +// structural register cut (turbo-read restructure), tracked separately. So D512 keeps min-blocks=2. +static constexpr __device__ int ggml_cuda_fattn_vec_min_blocks(int D) { + // D<=256 → 3 blocks/SM (168 regs, ~0 spill, 25% occupancy vs the old 16.7%). D512 keeps the upstream + // min_blocks=2: at 168 regs it would spill 240B (the D-aware nthreads_KQ that avoided this regressed + // the kernel, see above), and decode tps was flat regardless of D512 occupancy. NOTE (#535): this + // occupancy lift is byte-identical but only marginally moved tps (decode is FFN/MoE-bound + the turbo + // gap is the read dep-chain, not occupancy); kept because it's free, not because it closed the gap. + return D <= 256 ? GGML_CUDA_FATTN_VEC_MIN_BLOCKS_PER_SM : 2; +} + // Currently llvm with the amdgcn target does not support unrolling loops // that contain a break that can not be resolved at compile time. #ifdef __clang__ #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wpass-failed" #endif // __clang__ +// opencoti-hook: #533 TCQ decode smem-codebook dispatch — the FA-VEC TCQ readers normally index a +// 512/256-entry codebook in __constant__ memory; with data-dependent (per-element) indices that +// serializes up to 32-way per warp. These thin helpers call the _cb reader variants (fattn-tcq.cuh, +// which take an explicit codebook pointer) so the kernel can point them at a shared-memory copy. Picks +// the tier at compile time; only instantiated under `if constexpr (is_tcq_*)`, so non-TCQ kernels never +// see them. See docs/features/poly_kv.md / UPSTREAM_SYNC.md. +template +static __device__ __forceinline__ float tcq_vec_dot_KQ_smem( + const char * __restrict__ K_c, const void * __restrict__ Q_v, + const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v, const float * __restrict__ cb) { + if constexpr (tK == GGML_TYPE_TURBO3_TCQ) return vec_dot_fattn_vec_KQ_turbo3_tcq_cb(K_c, Q_v, Q_q8, Q_ds_v, cb); + else return vec_dot_fattn_vec_KQ_turbo2_tcq_cb(K_c, Q_v, Q_q8, Q_ds_v, cb); +} +template +static __device__ __forceinline__ void tcq_dequantize_V_smem( + const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0, const float * __restrict__ cb) { + if constexpr (tV == GGML_TYPE_TURBO3_TCQ) dequantize_V_turbo3_tcq_cb(vx, dst, i0, cb); + else dequantize_V_turbo2_tcq_cb(vx, dst, i0, cb); +} + template // D == head size -__launch_bounds__(ggml_cuda_fattn_vec_get_nthreads_device(), 1) +__launch_bounds__(ggml_cuda_fattn_vec_get_nthreads_device(), ggml_cuda_fattn_vec_min_blocks(D)) static __global__ void flash_attn_ext_vec( const char * __restrict__ Q, const char * __restrict__ K, @@ -26,6 +95,9 @@ static __global__ void flash_attn_ext_vec( 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 + const int * __restrict__ block_sel, // opencoti-hook: sparse-attn #551 — sorted −1-padded selected-block list (nullptr=dense) + const int32_t sparse_B, // opencoti-hook: sparse-attn #551 — block size (multiple of nthreads) + const int32_t sparse_nwords, // opencoti-hook: sparse-attn #551 — list stride == n_block per sequence float * __restrict__ dst, float2 * __restrict__ dst_meta, const float scale, @@ -46,7 +118,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, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, dst, dst_meta, scale, max_bias, m0, m1, n_head_log2, logit_softcap, ne00, ne01, ne02, ne03, nb01, nb02, nb03, @@ -85,6 +157,10 @@ static __global__ void flash_attn_ext_vec( // WHT-rotation precision the trellis reconstruction needs. V uses our quantized split (ne=4), same as turbo_0. || type_K == GGML_TYPE_TURBO3_TCQ || type_K == GGML_TYPE_TURBO2_TCQ); constexpr int nthreads = ggml_cuda_fattn_vec_get_nthreads_device(); + // opencoti #535 NOTE: tried D-aware nthreads_KQ for turbo D512 (32 lanes → Q_reg 64→16, 168reg/0spill), + // but it REGRESSED the D512 decode kernel 287µs→338µs (+18%): the benign L1-cached spill was cheaper + // than the 32-lane warp reduction it introduced, and decode tps was flat regardless (D512 attn is a + // small decode slice). Reverted to the f16-style fixed split. See docs/evaluations/context.md. constexpr int nthreads_KQ = (type_K == GGML_TYPE_F16 || type_K == GGML_TYPE_BF16 || is_turbo_K) ? 128 / cpy_nb : nthreads_KQ_q; constexpr int nthreads_V = (type_V == GGML_TYPE_F16 || type_V == GGML_TYPE_BF16) ? 128 / cpy_nb : nthreads_V_q; @@ -250,20 +326,90 @@ static __global__ void flash_attn_ext_vec( #endif // V_DOT2_F32_F16_AVAILABLE } + // opencoti-hook: #533 TCQ decode smem-codebook — copy the per-tier codebook into shared memory once + // (cooperative over all nthreads), so the K-dot / V-dequant below read it from smem (broadcast, no + // bank conflict on a hit) instead of __constant__ (which serializes the data-dependent per-element + // lookups up to 32-way per warp). Also lets the TURBO_TCQ_CB runtime override reach decode (it writes + // the __constant__ symbol we copy from). ≤2 KB smem + a single block barrier at the kernel head. + constexpr bool is_tcq_K = (type_K == GGML_TYPE_TURBO3_TCQ || type_K == GGML_TYPE_TURBO2_TCQ); + constexpr bool is_tcq_V = (type_V == GGML_TYPE_TURBO3_TCQ || type_V == GGML_TYPE_TURBO2_TCQ); + constexpr int cb_K_n = type_K == GGML_TYPE_TURBO3_TCQ ? 512 : (type_K == GGML_TYPE_TURBO2_TCQ ? 256 : 1); + constexpr int cb_V_n = type_V == GGML_TYPE_TURBO3_TCQ ? 512 : (type_V == GGML_TYPE_TURBO2_TCQ ? 256 : 1); + __shared__ float tcq_cb_K[cb_K_n]; + __shared__ float tcq_cb_V[cb_V_n]; + if constexpr (is_tcq_K) { + const float * src_K = (type_K == GGML_TYPE_TURBO3_TCQ) ? d_turbo3_tcq_codebook_fattn : d_turbo2_tcq_codebook_fattn; + for (int i = tid; i < cb_K_n; i += nthreads) tcq_cb_K[i] = src_K[i]; + } + if constexpr (is_tcq_V) { + const float * src_V = (type_V == GGML_TYPE_TURBO3_TCQ) ? d_turbo3_tcq_codebook_fattn : d_turbo2_tcq_codebook_fattn; + for (int i = tid; i < cb_V_n; i += nthreads) tcq_cb_V[i] = src_V[i]; + } + // opencoti-hook: #535b plain-turbo decode smem-codebook — same treatment as the TCQ tiers above, for + // the non-trellis turbo2/3/4_0 centroid tables (4/8/16 entries). __constant__ serializes the + // data-dependent per-element centroid lookups up to 32-way/warp; smem broadcasts. This is the proven + // root of the turbo3_tcq(49µs) < turbo3(63µs) D256 K-read gap (#534 fixed it for TCQ only). Byte-identical. + constexpr bool is_turbo0_K = (type_K == GGML_TYPE_TURBO2_0 || type_K == GGML_TYPE_TURBO3_0 || type_K == GGML_TYPE_TURBO4_0); + constexpr bool is_turbo0_V = (type_V == GGML_TYPE_TURBO2_0 || type_V == GGML_TYPE_TURBO3_0 || type_V == GGML_TYPE_TURBO4_0); + constexpr int cb_t0_K_n = type_K == GGML_TYPE_TURBO3_0 ? 8 : (type_K == GGML_TYPE_TURBO4_0 ? 16 : (type_K == GGML_TYPE_TURBO2_0 ? 4 : 1)); + constexpr int cb_t0_V_n = type_V == GGML_TYPE_TURBO3_0 ? 8 : (type_V == GGML_TYPE_TURBO4_0 ? 16 : (type_V == GGML_TYPE_TURBO2_0 ? 4 : 1)); + __shared__ float turbo_cb_K[cb_t0_K_n]; + __shared__ float turbo_cb_V[cb_t0_V_n]; + if constexpr (is_turbo0_K) { + const float * src_K = (type_K == GGML_TYPE_TURBO3_0) ? FATTN_TURBO3_CENTROIDS : (type_K == GGML_TYPE_TURBO4_0 ? FATTN_TURBO4_CENTROIDS : FATTN_TURBO2_CENTROIDS); + for (int i = tid; i < cb_t0_K_n; i += nthreads) turbo_cb_K[i] = src_K[i]; + } + if constexpr (is_turbo0_V) { + const float * src_V = (type_V == GGML_TYPE_TURBO3_0) ? FATTN_TURBO3_CENTROIDS : (type_V == GGML_TYPE_TURBO4_0 ? FATTN_TURBO4_CENTROIDS : FATTN_TURBO2_CENTROIDS); + for (int i = tid; i < cb_t0_V_n; i += nthreads) turbo_cb_V[i] = src_V[i]; + } + if constexpr (is_tcq_K || is_tcq_V || is_turbo0_K || is_turbo0_V) { __syncthreads(); } + 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-hook: sparse-attn #551 — compacted selected-block iteration (council Rank-1, csl-2026-06-26-0516). + // Decode is weight-bound; the old "iterate every tile, continue on unselected" loop paid the per-tile branch + + // K/V/mask pointer-advance on the ~3/4 skipped tiles with no compensating saving (−6% vs dense). Here the + // selector emits a SORTED, −1-padded list of selected block indices (sparse_nwords == n_block entries/seq); we + // iterate ONLY those, jumping K/V/mask directly to each tile — O(selected), not O(capacity). The cumulative + // offset of the original strided loop at tile k_VKQ_0 is exactly base + k_VKQ_0 (mask elements / K,V bytes via + // nb11,nb21), so a from-base recompute is identical. Dense (block_sel==nullptr) walks every tile in that same + // strided order ⇒ byte-identical; topk=0 ⇒ list is [0,1,2,…] ⇒ also byte-identical. A gridDim.y block assigned + // zero tiles falls through to the identity partial write (KQ_max=−FLT_MAX/2, KQ_sum=0) below, leaving the + // online-softmax combine unaffected. sparse_B is a multiple of nthreads (enforced graph-side), so sparse_m≥1. + const char * const K_base = K; + const char * const V_base = V; + const auto maskh_base = maskh; + const int * const sparse_sel = block_sel ? block_sel + sequence*sparse_nwords : nullptr; // sorted, −1-padded + const int sparse_m = block_sel ? (sparse_B / nthreads) : 1; // tiles per selected block (≥1) + const int sparse_nw = block_sel ? sparse_nwords*sparse_m // upper bound: n_block*m tiles + : (k_VKQ_max + nthreads - 1)/nthreads; + for (int kv_it = blockIdx.y; kv_it < sparse_nw; kv_it += gridDim.y) { + int k_VKQ_0; + if (sparse_sel) { + const int blk = sparse_sel[kv_it / sparse_m]; + if (blk < 0) { // −1 sentinel: past the last selected block + break; + } + k_VKQ_0 = blk*sparse_B + (kv_it % sparse_m)*nthreads; + if (k_VKQ_0 >= k_VKQ_max) { // selected block beyond valid KV / DCA high edge + continue; + } + } else { + k_VKQ_0 = kv_it * nthreads; + if (k_VKQ_0 >= k_VKQ_max) { + break; + } + } // 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. + // (DCA INTRA/SUCC regimes, SWA windows). k_VKQ_min==0 for causal masks → never skips. if (k_VKQ_0 + nthreads <= k_VKQ_min) { continue; } + // Jump K/V/mask directly to this tile (char* byte strides nb11/nb21; maskh element stride). + K = K_base + (size_t) k_VKQ_0 * nb11; + V = V_base + (size_t) k_VKQ_0 * nb21; + maskh = maskh_base + k_VKQ_0; // Calculate KQ tile and keep track of new maximum KQ values: float KQ_reg[ncols]; // KQ in registers. @@ -280,7 +426,14 @@ static __global__ void flash_attn_ext_vec( #pragma unroll for (int j = 0; j < ncols; ++j) { - float sum = vec_dot_KQ(K + i_KQ*nb11, Q_reg[j], Q_i32[j], Q_ds[j]); + float sum; + if constexpr (is_tcq_K) { // opencoti #533: read codebook from smem, not __constant__ + sum = tcq_vec_dot_KQ_smem(K + i_KQ*nb11, Q_reg[j], Q_i32[j], Q_ds[j], tcq_cb_K); + } else if constexpr (is_turbo0_K) { // opencoti #535b: plain-turbo smem codebook + sum = turbo_vec_dot_KQ_smem(K + i_KQ*nb11, Q_reg[j], Q_i32[j], Q_ds[j], turbo_cb_K); + } else { + sum = vec_dot_KQ(K + i_KQ*nb11, Q_reg[j], Q_i32[j], Q_ds[j]); + } sum = warp_reduce_sum(sum); if (use_logit_softcap) { @@ -341,6 +494,30 @@ static __global__ void flash_attn_ext_vec( for (int j = 0; j < ncols; ++j) { KQ_k[j] = __half2half2(KQ[j*nthreads + k]); } + // opencoti-hook: sparse-v (#546) — skip V-dequant+accumulate for negligible-weight positions. + // Quantized-V only (f16/bf16 dequant is trivial). Divergent per-thread continue is safe here: no + // smem write / __syncwarp in this k0 iteration, and a skipped position's weighted V ≈ 0 (its tiny + // weight stays in the softmax normalizer, so the error is bounded by tau ⇒ below f16 noise). + if constexpr (type_V != GGML_TYPE_F16 && type_V != GGML_TYPE_BF16) { + // opencoti-hook: sparse-v combined skip (#546/#565). pos = absolute KV index of this key. + const int pos = k_VKQ_0 + k; + const int n_kv = (int) ne11; + if (pos >= d_sparse_v_sink && pos < n_kv - d_sparse_v_recent) { + // #565 partial-L FIX: threshold = eps × RUNNING MEAN (KQ_sum / keys-processed), NOT + // eps/n_kv·KQ_sum (partial-sum ÷ total ⇒ too small early ⇒ skip vanishes as ctx grows ⇒ + // the observed inversion). keys summed into KQ_sum so far = k_VKQ_0 + nthreads. eps==0 ⇒ + // thr=d_sparse_v_tau (∀-col 0.0f) ? d_sparse_v_eps * d_inv_nproc * KQ_sum[jj] : d_sparse_v_tau; + if (w >= thr) { opencoti_skip = false; break; } + } + if (opencoti_skip) { continue; } + } + } #pragma unroll for (int i_VKQ_0 = 0; i_VKQ_0 < D/2; i_VKQ_0 += nthreads_V*V_rows_per_thread/2) { half2 tmp[V_rows_per_thread/2]; @@ -352,6 +529,12 @@ static __global__ void flash_attn_ext_vec( for (int i_VKQ_1 = 0; i_VKQ_1 < V_rows_per_thread/2; ++i_VKQ_1) { tmp[i_VKQ_1] = __float22half2_rn(tmp_f[i_VKQ_1]); } + } else if constexpr (is_tcq_V) { // opencoti #533: read codebook from smem, not __constant__ + tcq_dequantize_V_smem(V + k*nb21, tmp, + 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread, tcq_cb_V); + } else if constexpr (is_turbo0_V) { // opencoti #535b: plain-turbo smem codebook + turbo_dequantize_V_smem(V + k*nb21, tmp, + 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread, turbo_cb_V); } else { dequantize_V(V + k*nb21, tmp, 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread); @@ -370,11 +553,35 @@ static __global__ void flash_attn_ext_vec( for (int j = 0; j < ncols; ++j) { KQ_k[j] = KQ[j*nthreads + k]; } + // opencoti-hook: sparse-v combined skip (#546/#565) — see half2 path above. eps==tau==0 byte-identical. + if constexpr (type_V != GGML_TYPE_F16 && type_V != GGML_TYPE_BF16) { + const int pos = k_VKQ_0 + k; + const int n_kv = (int) ne11; + if (pos >= d_sparse_v_sink && pos < n_kv - d_sparse_v_recent) { + const float d_inv_nproc = 1.0f / (float) (k_VKQ_0 + nthreads); + bool opencoti_skip = true; +#pragma unroll + for (int jj = 0; jj < ncols; ++jj) { + const float w = KQ_k[jj]; + const float thr = (d_sparse_v_eps > 0.0f) ? d_sparse_v_eps * d_inv_nproc * KQ_sum[jj] : d_sparse_v_tau; + if (w >= thr) { opencoti_skip = false; break; } + } + if (opencoti_skip) { continue; } + } + } #pragma unroll for (int i_VKQ_0 = 0; i_VKQ_0 < D/2; i_VKQ_0 += nthreads_V*V_rows_per_thread/2) { float2 tmp[V_rows_per_thread/2]; - dequantize_V(V + k*nb21, tmp, - 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread); + if constexpr (is_tcq_V) { // opencoti #533: read codebook from smem, not __constant__ + tcq_dequantize_V_smem(V + k*nb21, tmp, + 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread, tcq_cb_V); + } else if constexpr (is_turbo0_V) { // opencoti #535b: plain-turbo smem codebook + turbo_dequantize_V_smem(V + k*nb21, tmp, + 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread, turbo_cb_V); + } else { + dequantize_V(V + k*nb21, tmp, + 2*i_VKQ_0 + (nthreads_V == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads_V)*V_rows_per_thread); + } #pragma unroll for (int i_VKQ_1 = 0; i_VKQ_1 < V_rows_per_thread/2; ++i_VKQ_1) { #pragma unroll @@ -524,7 +731,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, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, dst, dst_meta, scale, max_bias, m0, m1, n_head_log2, logit_softcap, ne00, ne01, ne02, ne03, nb01, nb02, nb03, @@ -542,6 +749,26 @@ static __global__ void flash_attn_ext_vec( template void ggml_cuda_flash_attn_ext_vec_case_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + // opencoti-hook: sparse-v (#546) — broadcast the host threshold into THIS instance TU's __constant__ + // copy once (no device RDC ⇒ per-TU symbol). TURBO_SPARSE_V_TAU unset/0 ⇒ dense V-read (byte-identical). + static bool opencoti_sparse_v_set = false; + if (!opencoti_sparse_v_set) { + const char * e = getenv("TURBO_SPARSE_V_TAU"); + const float tau = e ? (float) atof(e) : 0.0f; + CUDA_CHECK(cudaMemcpyToSymbol(d_sparse_v_tau, &tau, sizeof(float))); + // opencoti-hook: sparse-v adaptive-τ (#565) — mass-budget override; unset/0 ⇒ static-τ, byte-identical. + const char * ee = getenv("TURBO_SPARSE_V_EPS"); + const float eps = ee ? (float) atof(ee) : 0.0f; + CUDA_CHECK(cudaMemcpyToSymbol(d_sparse_v_eps, &eps, sizeof(float))); + // opencoti-hook: sparse-v combined (#565) — sink/recent always-keep guards (unset/0 ⇒ no protection). + const char * es = getenv("TURBO_SPARSE_V_SINK"); + const int sink = es ? atoi(es) : 0; + CUDA_CHECK(cudaMemcpyToSymbol(d_sparse_v_sink, &sink, sizeof(int))); + const char * er = getenv("TURBO_SPARSE_V_RECENT"); + const int recent = er ? atoi(er) : 0; + CUDA_CHECK(cudaMemcpyToSymbol(d_sparse_v_recent, &recent, sizeof(int))); + opencoti_sparse_v_set = true; + } const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int nthreads = ggml_cuda_fattn_vec_get_nthreads_host(cc); diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu b/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu index 0f582fc..85dac93 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-wmma-f16.cu @@ -31,6 +31,9 @@ static __global__ void flash_attn_ext_f16( 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 + const int * __restrict__ block_sel, // opencoti-hook: sparse-attn #551 — unused in WMMA (legacy path) + const int32_t sparse_B, // opencoti-hook: sparse-attn #551 + const int32_t sparse_nwords, // opencoti-hook: sparse-attn #551 float * __restrict__ dst, float2 * __restrict__ dst_meta, const float scale, @@ -501,7 +504,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, KV_min, dst, dst_meta, scale, + GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, KV_min, block_sel, sparse_B, sparse_nwords, 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 index 23d90a3..e53ac9c 100644 --- a/llama.cpp/ggml/src/ggml-cuda/fattn.cu +++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cu @@ -28,6 +28,15 @@ 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; +// opencoti-hook: sparse-attn-vslash side-output (#551 S1) — definitions of the block-mass side channel +// declared in fattn-common.cuh. Set by the vslash prefill graph site before the FA dispatch; drained by +// launch_fattn. The MMA kernel emits per-block question-window mass here (S2) when non-null; until the +// S2 kernel emit lands these are set/drained but unconsumed (inert, byte-identical). +thread_local float * opencoti_fattn_block_mass = nullptr; +thread_local int opencoti_fattn_block_mass_nblock = 0; +thread_local int opencoti_fattn_block_mass_B = 0; +thread_local int opencoti_fattn_block_mass_qwin = 0; + template static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; @@ -574,6 +583,16 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const return BEST_FATTN_KERNEL_VEC; } + // opencoti-hook: sparse-attn #551 (bug-742) — the block-selector skip walk lives ONLY in the FA-VEC + // kernel (fattn-vec.cuh compacted −1-padded list). High-GQA + long-KV decode (gqa>4 && K->ne[1]>=8192, + // the exact 1M-serving regime) otherwise falls to the MMA path below, which has no block_sel reader and + // streams the FULL KV — silently dropping the skip (measured: tps bit-invariant across k_sel 32..512). + // When the selector attaches a block_sel list (dst->src[7] != nullptr, i.e. sparse-attn active with + // topk!=0), force VEC so the skip is honored. Dense/topk0 leaves src[7]==nullptr ⇒ routing unchanged. + if (dst->src[7] != nullptr && can_use_vector_kernel) { + return BEST_FATTN_KERNEL_VEC; + } + // If Turing tensor cores are available, use them: if (turing_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) { if (can_use_vector_kernel) { diff --git a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu index a6ef9ee..4db22f3 100644 --- a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -69,6 +69,7 @@ #include "ggml-cuda/cumsum.cuh" #include "ggml-cuda/fill.cuh" #include "ggml-cuda/turbo-wht.cuh" // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391 +#include "ggml-cuda/sparse-attn.cuh" // opencoti-hook: sparse-attn — #551 Quest block-selector #include "ggml.h" #include @@ -3047,6 +3048,21 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg // turbo-wht.cu. ggml_cuda_turbo_wht(ctx, dst); break; + case GGML_OP_SPARSE_ATTN_FILL: + // opencoti-hook: sparse-attn — #551 B1. Quest per-block kmin/kmax fill from the + // freshly-written cache K. See sparse-attn.cu. + ggml_cuda_sparse_attn_fill(ctx, dst); + break; + case GGML_OP_SPARSE_ATTN_SELECT: + // opencoti-hook: sparse-attn — #551 B2. Quest per-(seq) block bitmask from Q·kbounds + // upper bound + top-K + recent/sink. See sparse-attn.cu. + ggml_cuda_sparse_attn_select(ctx, dst); + break; + case GGML_OP_SPARSE_ATTN_VSLASH: + // opencoti-hook: sparse-attn-vslash — #551 content-driven per-(seq) block list from the + // actual per-block max logit (last q_win queries · cache K) + top-K + recent/sink. See sparse-attn.cu. + ggml_cuda_sparse_attn_vslash(ctx, dst); + break; case GGML_OP_OUT_PROD: ggml_cuda_out_prod(ctx, dst); break; @@ -5421,7 +5437,9 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de // build_attn_mha's ggml_cast(turboN→f16) on CUDA instead of spilling to the CPU // backend. See docs/features/poly_kv.md (M6-S2 perf-path Level A). if ((src0_type == GGML_TYPE_TURBO2_0 || src0_type == GGML_TYPE_TURBO3_0 || - src0_type == GGML_TYPE_TURBO4_0 || src0_type == GGML_TYPE_TURBO8_0) && + src0_type == GGML_TYPE_TURBO4_0 || src0_type == GGML_TYPE_TURBO8_0 || + // opencoti-hook #533: TCQ→f16 dequant-on-lift (prefill-MMA + DCA-compose) + src0_type == GGML_TYPE_TURBO3_TCQ || src0_type == GGML_TYPE_TURBO2_TCQ) && src1_type == GGML_TYPE_F16) { return true; } @@ -5687,6 +5705,41 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de memcpy(&group_size, (const char *) op->op_params + sizeof(int), sizeof(int)); return group_size == 128 && a->ne[0] % 128 == 0; } + case GGML_OP_SPARSE_ATTN_FILL: { + // opencoti-hook: sparse-attn — #551 B1 / bug-744. f16 kbounds dst [n_embd,n_block,2] + + // k_cur (src[1], the NEW tokens, pre-quantization f16 OR f32) + I64-or-I32 k_idxs (src[2]). + // cpy_k scatters via ggml_set_rows whose indices are I64, so accept both. Reading k_cur + // (not the quantized cache view) is what lets the Quest bound compose with EVERY cache KV + // type — the cache quantization happens downstream in cpy_k; the bound is the exact pre-quant + // envelope (Quest's upper bound stays conservative). + const ggml_tensor * k = op->src[1]; + const ggml_tensor * k_idxs = op->src[2]; + if (op->type != GGML_TYPE_F16 || op->ne[2] != 2) return false; + if (!k || (k->type != GGML_TYPE_F16 && k->type != GGML_TYPE_F32)) return false; + if (!k_idxs || + (k_idxs->type != GGML_TYPE_I64 && k_idxs->type != GGML_TYPE_I32)) return false; + return true; + } + case GGML_OP_SPARSE_ATTN_SELECT: { + // opencoti-hook: sparse-attn — #551 B2. I32 bitmask dst; f32 query (src[0]); + // f16 kbounds [n_embd,n_block,2] (src[1]). + const ggml_tensor * q = op->src[0]; + const ggml_tensor * kbounds = op->src[1]; + if (op->type != GGML_TYPE_I32) return false; + if (!q || q->type != GGML_TYPE_F32) return false; + if (!kbounds || kbounds->type != GGML_TYPE_F16 || kbounds->ne[2] != 2) return false; + return true; + } + case GGML_OP_SPARSE_ATTN_VSLASH: { + // opencoti-hook: sparse-attn-vslash — #551. I32 list dst; f32 query (src[0]); + // f16 cache K (src[1], the actual K, NOT a kbounds side-cache). + const ggml_tensor * q = op->src[0]; + const ggml_tensor * k_full = op->src[1]; + if (op->type != GGML_TYPE_I32) return false; + if (!q || q->type != GGML_TYPE_F32) return false; + if (!k_full || k_full->type != GGML_TYPE_F16) return false; + return true; + } case GGML_OP_CROSS_ENTROPY_LOSS: case GGML_OP_CROSS_ENTROPY_LOSS_BACK: case GGML_OP_OPT_STEP_ADAMW: diff --git a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu index bd4f5a0..ccedfde 100644 --- a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu +++ b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu @@ -738,6 +738,10 @@ static void set_rows_cuda_turbo3_tcq(ggml_backend_cuda_context & ctx, const ggml // opencoti-hook: asymmetric TCQ (#513) — also override the per-TU d_turbo3 copies in the // two asymmetric instance TUs so the runtime codebook reaches asym decode (not just symmetric). tcq3_set_decode_codebook_t3t2(cb); tcq3_set_decode_codebook_t2t3(cb); + // opencoti-hook: TCQ A′ — MMA-lift materialize (#553/bug-710). Also override the + // turbo-cpy.cu d_turbo3_tcq_cb_cpy copy so prefill/spec-verify (TCQ→f16 lift) decodes + // with the same runtime codebook as FA-VEC decode (else draft⇄verify KV split). + tcq3_set_cpy_codebook(cb); fprintf(stderr, "TCQ encode+decode: loaded codebook from %s (device %d)\n", cb_path, ctx.device); } else { if (f) fclose(f); fprintf(stderr, "TCQ encode: FAILED to load codebook from %s\n", cb_path); } } @@ -801,6 +805,8 @@ static void set_rows_cuda_turbo2_tcq(ggml_backend_cuda_context & ctx, const ggml // opencoti-hook: asymmetric TCQ (#513) — also override the per-TU d_turbo2 copies in the // two asymmetric instance TUs so the runtime 2-bit codebook reaches asym decode too. tcq2_set_decode_codebook_t3t2(cb); tcq2_set_decode_codebook_t2t3(cb); + // opencoti-hook: TCQ A′ — MMA-lift materialize (#553/bug-710). Mirror for the 2-bit cpy copy. + tcq2_set_cpy_codebook(cb); fprintf(stderr, "TCQ2 encode+decode: loaded 2-bit codebook from %s (device %d)\n", cb_path, ctx.device); } else { if (f) fclose(f); fprintf(stderr, "TCQ2 encode: FAILED to load codebook from %s\n", cb_path); } } diff --git a/bk-t1/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cu b/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cu new file mode 100644 index 0000000..0f8a70f --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cu @@ -0,0 +1,537 @@ +// opencoti-hook: sparse-attn — #551 Quest block-selector, B1 (the per-block kmin/kmax fill). +// See docs/features/sparse_attn.md (S2/S3 design). +// +// Reduces the freshly-written cache K into per-(KV-head·channel, block) min/max, stored in the +// kbounds side-cache [n_embd_k, n_block, 2] ([..,0]=min, [..,1]=max). Recompute-from-cache-K +// (not incremental-merge): each affected block is recomputed in full over its written positions, +// so the append-only stream self-corrects with NO ±inf init and NO merge. At decode only the +// current (last) block is affected; at prefill all blocks are computed once. + +#include "sparse-attn.cuh" + +// opencoti #551 bug-744: load one K channel as f32 regardless of K's storage type. The fill now reads +// k_cur (the NEW tokens, pre-quantization) — always f16 or f32 — so the kmin/kmax bound composes with +// EVERY KV cache type (q8_0/q4_0/q6_0/turbo*/tcq*): the cache is quantized AFTER this, the bound is the +// exact pre-quant envelope (Quest's upper bound is conservative either way). No dequant in the fill. +template static __device__ __forceinline__ float sparse_ld(const kt * p, int d); +template <> __device__ __forceinline__ float sparse_ld(const half * p, int d) { return __half2float(p[d]); } +template <> __device__ __forceinline__ float sparse_ld(const float * p, int d) { return p[d]; } + +template +static __global__ void sparse_attn_fill_kernel( + const char * __restrict__ k_base, // k_cur (NEW tokens), f16/f32; column (p-p0) at k_base+(p-p0)*nb1 + const idx_t * __restrict__ k_idxs, // newly-written cache positions (I32 or I64); k_idxs[0] = p0 + half * __restrict__ kbounds, // [n_embd, n_block, 2] contiguous (persistent across steps) + const int n_embd, + const int n_new, // # of new-token columns present in k_base + const int n_block, + const int B, + const int64_t nb1) { // byte stride between k_cur columns (= k_cur->nb[2]) + const int b = blockIdx.x; + if (b >= n_block) { + return; + } + const int p0 = (int) k_idxs[0]; // first newly-written cache position + const int n_kv = p0 + n_new; // total cache length after this write + const int b_first = p0 / B; + // Append-only: blocks strictly below the first touched block are unchanged from a prior step. + if (b < b_first) { + return; + } + const int kv_lo = b * B; + int kv_hi = (b + 1) * B; + if (kv_hi > n_kv) { + kv_hi = n_kv; + } + if (kv_lo >= kv_hi) { + return; // block beyond the current cache length + } + // Only the NEW positions [max(kv_lo,p0), kv_hi) live in k_base; a block that already held tokens + // (kv_lo < p0, i.e. the partial last block at decode) MERGES into its prior kmin/kmax, a fully-new + // block OVERWRITES. min/max is exact so this reproduces the dense recompute byte-for-byte (f16 path). + const int new_lo = kv_lo < p0 ? p0 : kv_lo; + const bool straddle = kv_lo < p0; + + for (int d = threadIdx.x; d < n_embd; d += blockDim.x) { + float vmin = INFINITY; + float vmax = -INFINITY; + for (int p = new_lo; p < kv_hi; ++p) { + const kt * col = (const kt *) (k_base + (int64_t) (p - p0) * nb1); + const float v = sparse_ld(col, d); + vmin = fminf(vmin, v); + vmax = fmaxf(vmax, v); + } + const int64_t imin = (int64_t) b * n_embd + d; // w=0 (min) + const int64_t imax = ((int64_t) n_block + b) * n_embd + d; // w=1 (max) + if (straddle) { + kbounds[imin] = __float2half(fminf(__half2float(kbounds[imin]), vmin)); + kbounds[imax] = __float2half(fmaxf(__half2float(kbounds[imax]), vmax)); + } else { + kbounds[imin] = __float2half(vmin); + kbounds[imax] = __float2half(vmax); + } + } +} + +void ggml_cuda_sparse_attn_fill(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * k = dst->src[1]; // opencoti bug-744: k_cur (new tokens), NOT the cache view + const ggml_tensor * k_idxs = dst->src[2]; + + GGML_ASSERT(dst->type == GGML_TYPE_F16); + // opencoti bug-744: k_cur is the pre-quant K projection — f16 or f32 — so sparse composes with ALL + // cache KV types (the cache quantization happens downstream, in cpy_k). + GGML_ASSERT(k->type == GGML_TYPE_F16 || k->type == GGML_TYPE_F32); + GGML_ASSERT(k_idxs->type == GGML_TYPE_I64 || k_idxs->type == GGML_TYPE_I32); + GGML_ASSERT(dst->ne[2] == 2); + + int B = 0; + memcpy(&B, dst->op_params, sizeof(int)); + GGML_ASSERT(B > 0); + + // opencoti #551 PERF-DIAG: LLAMA_SPARSE_NO_FILL skips the kmin/kmax update (kbounds goes stale — + // perf-bisection ONLY, not correctness). Isolates the fill cost from select on the decode path. + static const bool no_fill = getenv("LLAMA_SPARSE_NO_FILL") != nullptr; + if (no_fill) { return; } + + const int n_embd = dst->ne[0]; + const int n_block = dst->ne[1]; + // k_cur is [head_dim, n_head_kv, n_new, 1]; the new-token count is the position axis ne[2], and the + // per-token column stride is nb[2]. n_embd (= head_dim·n_head_kv) is the contiguous channel span. + const int n_new = k->ne[2]; + const int64_t nb1 = k->nb[2]; + + half * kb = (half *) dst->data; // dst aliases the kbounds buffer (cpy idiom) + + if (n_new <= 0 || n_block <= 0) { + return; + } + + const int nthreads = n_embd < 256 ? (n_embd < 32 ? 32 : n_embd) : 256; + const bool i64 = (k_idxs->type == GGML_TYPE_I64); + const bool kf16 = (k->type == GGML_TYPE_F16); + if (i64 && kf16) { + sparse_attn_fill_kernel<<>>( + (const char *) k->data, (const int64_t *) k_idxs->data, kb, n_embd, n_new, n_block, B, nb1); + } else if (i64) { + sparse_attn_fill_kernel<<>>( + (const char *) k->data, (const int64_t *) k_idxs->data, kb, n_embd, n_new, n_block, B, nb1); + } else if (kf16) { + sparse_attn_fill_kernel<<>>( + (const char *) k->data, (const int32_t *) k_idxs->data, kb, n_embd, n_new, n_block, B, nb1); + } else { + sparse_attn_fill_kernel<<>>( + (const char *) k->data, (const int32_t *) k_idxs->data, kb, n_embd, n_new, n_block, B, nb1); + } +} + +// --------------------------------------------------------------------------------------------- +// GGML_OP_SPARSE_ATTN_SELECT — Quest block selector (B2). Two passes: +// (1) score: per (block b, sequence s), score = max over query-heads h and rows of the Quest +// upper bound Σ_d max(q_{h,d}·kmin_{d,b}, q_{h,d}·kmax_{d,b}) → global scratch. +// (2) select: per sequence, top-K_SEL blocks by score (threshold bisection) OR the recent/sink +// window → a SORTED, −1-padded list of selected block indices dst[s*n_block + i] (#551 +// compacted list; the FA-VEC decode kernel walks it and stops at the first −1). +// Union-over-heads/rows is exact for decode (one query row) and a safe superset for prefill. + +#define SPARSE_SEL_NT 256 + +static __global__ void sparse_attn_score_kernel( + const char * __restrict__ q_base, // FA query (f32), [head_dim, n_head, n_q, n_seq] + const half * __restrict__ kbounds, // [n_embd_k, n_block, 2] contiguous + float * __restrict__ scores, // [n_block, n_seq] out + const int * __restrict__ k_idxs, // #551 PERF: KV write indices (decode position), or nullptr + const int head_dim, + const int n_head, + const int n_q, + const int n_embd_k, + const int n_block, + const int n_sub, // #551 bug-743b: kbounds ROW count (= fill's ne[1]); the max + // plane lives at physical rows [n_sub, 2*n_sub). MUST match the + // fill's base, NOT the dynamic walk n_block (= ceil(n_kv/B)). + const int sub_per_block, // #551: kbounds sub-blocks per walk-block; score = max over them + const int refresh, // re-score every N decode tokens; skip the work otherwise + const int64_t nb1, const int64_t nb2, const int64_t nb3) { // q strides (bytes) + // #551 PERF (amortized): skip the (expensive) score computation on non-refresh decode steps — the select + // kernel reuses the cached list, so these scores are never read. Gate on the DEVICE position (k_idxs[0]) + // so it stays correct under CUDA-graph replay. Uniform return across the block. + if (k_idxs && refresh > 1) { + const int pos = k_idxs[0]; + if (pos % refresh != 0) { + return; + } + } + const int b = blockIdx.x; + const int s = blockIdx.y; + if (b >= n_block) { + return; + } + const int tid = threadIdx.x; + const int n_head_kv = n_embd_k / head_dim; + const int gqa = n_head / n_head_kv; + + __shared__ float sbuf[SPARSE_SEL_NT]; + __shared__ float s_score; + if (tid == 0) { + s_score = -INFINITY; + } + __syncthreads(); + + // #551 sub-block Quest bound: walk-block b owns kbounds sub-rows [b*SUB, b*SUB+SUB). The block score is + // the max over (head, query-row, sub-block) of the per-sub Quest upper bound. Finer sub-blocks tighten + // the kmin/kmax envelope so a needle's local sub-block scores high even when the 128-wide block is loose. + const int sub0 = b * sub_per_block; + const int sub1 = min(sub0 + sub_per_block, n_sub); + for (int h = 0; h < n_head; ++h) { + const int kvh = (gqa > 0) ? (h / gqa) : 0; + for (int qr = 0; qr < n_q; ++qr) { + const char * q_col = q_base + (int64_t) s*nb3 + (int64_t) qr*nb2 + (int64_t) h*nb1; + for (int sb = sub0; sb < sub1; ++sb) { + float partial = 0.0f; + for (int j = tid; j < head_dim; j += blockDim.x) { + const float qv = *(const float *) (q_col + (int64_t) j*4); + const int ch = kvh*head_dim + j; + const float fmin = __half2float(kbounds[(int64_t) sb *n_embd_k + ch]); + const float fmax = __half2float(kbounds[(int64_t)(n_sub + sb)*n_embd_k + ch]); + partial += fmaxf(qv*fmin, qv*fmax); + } + sbuf[tid] = partial; + __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { + if (tid < r) { sbuf[tid] += sbuf[tid + r]; } + __syncthreads(); + } + if (tid == 0) { s_score = fmaxf(s_score, sbuf[0]); } + __syncthreads(); + } + } + } + if (tid == 0) { + scores[(int64_t) s*n_block + b] = s_score; + } +} + +static __global__ void sparse_attn_select_kernel( + const float * __restrict__ scores, // [n_block, n_seq] + int * __restrict__ dst, // [n_list, n_seq] I32 sorted −1-padded selected-block list (#551) + const int * __restrict__ k_idxs, // #551 PERF: KV write indices (decode position), or nullptr + const int n_block, + const int n_list, // dst row stride/capacity (>= n_block; persistent buffer adds slack) + const int k_sel, + const int recent_blocks, + const int sink_blocks, + const int refresh, // re-select every N decode tokens; reuse the cached list otherwise + const int dbg) { // opencoti #551 bug-743 DIAG: print cnt/T/score-range when >0 + const int s = blockIdx.x; + const int tid = threadIdx.x; + + // #551 PERF (amortized selection): with a persistent buffer (k_idxs != nullptr) re-select only every + // `refresh` decode tokens — read the position from k_idxs[0] on the DEVICE so the decision stays correct + // under CUDA-graph replay (the launch args are captured, but device memory is re-read each replay). On a + // skipped step every thread returns uniformly, leaving the previously written list intact for FA-VEC. + if (k_idxs && refresh > 1) { + const int pos = k_idxs[0]; + if (pos % refresh != 0) { + return; + } + } + + const float * sc = scores + (int64_t) s*n_block; + int * out = dst + (int64_t) s*n_list; + + // init the whole list row to −1 (padding / FA-VEC walk terminator). Clearing the full capacity (n_list, + // not just n_block) guarantees stale entries past the emitted count are −1 even when the persistent + // buffer was sized for a larger ctx than the current n_block. + for (int b = tid; b < n_list; b += blockDim.x) { out[b] = -1; } + __syncthreads(); + + __shared__ float sbuf[SPARSE_SEL_NT]; + __shared__ float s_lo, s_hi, s_T; + __shared__ float s_scmin, s_scmax; // opencoti #551 bug-743 DIAG: pre-bisection score range + __shared__ int sibuf[SPARSE_SEL_NT]; + + // block min/max of scores + float lmin = INFINITY, lmax = -INFINITY; + for (int b = tid; b < n_block; b += blockDim.x) { + lmin = fminf(lmin, sc[b]); lmax = fmaxf(lmax, sc[b]); + } + sbuf[tid] = lmin; __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { if (tid < r) sbuf[tid] = fminf(sbuf[tid], sbuf[tid+r]); __syncthreads(); } + if (tid == 0) { s_lo = sbuf[0]; s_scmin = sbuf[0]; } + __syncthreads(); + sbuf[tid] = lmax; __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { if (tid < r) sbuf[tid] = fmaxf(sbuf[tid], sbuf[tid+r]); __syncthreads(); } + if (tid == 0) { s_hi = sbuf[0]; s_scmax = sbuf[0]; } + __syncthreads(); + + if (k_sel <= 0 || k_sel >= n_block) { + if (tid == 0) s_T = -INFINITY; // select-all anchor (topk==0 ⇒ inert, byte-identical) + } else { + // bisection: find threshold T so that #{b: score[b] >= T} ≈ k_sel + for (int it = 0; it < 30; ++it) { + __syncthreads(); + const float mid = 0.5f*(s_lo + s_hi); + int local = 0; + for (int b = tid; b < n_block; b += blockDim.x) { local += (sc[b] >= mid); } + sibuf[tid] = local; __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { if (tid < r) sibuf[tid] += sibuf[tid+r]; __syncthreads(); } + if (tid == 0) { if (sibuf[0] > k_sel) s_lo = mid; else s_hi = mid; } + } + if (tid == 0) s_T = s_lo; + } + __syncthreads(); + const float T = s_T; + + // emit a SORTED compacted list of kept block indices (single ordered pass; n_block is a few k and the + // select kernel is ~0-cost at decode). Ascending order ⇒ select-all reproduces the dense tile order, so + // topk=0 stays byte-identical with the dense FA-VEC path; the trailing −1 entries terminate the walk. + if (tid == 0) { + int cnt = 0; + for (int b = 0; b < n_block; ++b) { + const bool keep = (sc[b] >= T) || (b < sink_blocks) || (b >= n_block - recent_blocks); + if (keep) { out[cnt++] = b; } + } + // opencoti #551 bug-743 DIAG: emit the actual kept-count + threshold + score range for s==0. + if (dbg > 0 && s == 0) { + const int last = n_block > 0 ? n_block - 1 : 0; + printf("[sparse-sel s=0] n_block=%d k_sel=%d recent=%d sink=%d T=%g scmin=%g scmax=%g sc[0]=%g sc[mid]=%g sc[last]=%g => cnt=%d first=[%d %d %d %d] last_emit=%d\n", + n_block, k_sel, recent_blocks, sink_blocks, (double) T, (double) s_scmin, (double) s_scmax, + (double) sc[0], (double) sc[n_block/2], (double) sc[last], + cnt, cnt>0?out[0]:-9, cnt>1?out[1]:-9, cnt>2?out[2]:-9, cnt>3?out[3]:-9, cnt>0?out[cnt-1]:-9); + } + } +} + +// --------------------------------------------------------------------------------------------- +// GGML_OP_SPARSE_ATTN_VSLASH — content-driven block selector (MInference-style, #551). +// score(b) = max over heads h of Σ_{q in last q_win rows} Σ_{p in block b} softmax_h(q·K)[p] +// i.e. the per-block attention MASS (sum of softmax probability that lands in the block), unioned +// across query-heads by max. This is the MInference vertical-importance estimator, NOT the max +// attention logit and NOT the Quest geometric spread bound. Both of those score a block by a MAX +// over its keys — an extreme-value statistic whose distractor inflation grows with key count, so +// the needle's lone constant-signal key gets buried as context grows and the lossless coverage +// fraction RISES with ctx (the failure mode that sank Quest AND the max-logit vslash; bug-741). +// Softmax is winner-take-all per query: block mass is bounded [0,1] and concentrates on the blocks +// the query actually attends to, regardless of ctx length. One CUDA block per (query-head, seq) so +// each head owns its softmax partition Z. The select kernel (top-K + recent/sink sorted compaction) +// is reused VERBATIM. scores must be PRE-ZEROED (the launcher memsets) — heads accumulate by max. +static __global__ void vslash_score_kernel( + const char * __restrict__ q_base, // FA query (f32), [head_dim, n_head, n_q, n_seq] + const char * __restrict__ k_base, // cache K (f16), column (p, kvh) at k_base + p*knb2 + kvh*knb1 + float * __restrict__ scores, // [n_block, n_seq] out, PRE-ZEROED; atomic-max accumulated + const int * __restrict__ k_idxs, // #551 PERF: KV write index (decode position), or nullptr + const int head_dim, + const int n_head, + const int n_q, + const int n_kv, + const int n_block, + const int B, // block size (keys per block) + const int gqa, // query heads per KV head (kvh = h / gqa) + const int q_win, // # of recent query rows that define the window + const int refresh, + const int64_t qnb1, const int64_t qnb2, const int64_t qnb3, // q strides (bytes) + const int64_t knb1, const int64_t knb2) { // k per-KV-head / per-position strides (bytes) + // #551 PERF (amortized): skip scoring on non-refresh decode steps — select reuses the cached list, + // so these scores are never read. Gate on the DEVICE position (k_idxs[0]) for CUDA-graph-replay safety. + if (k_idxs && refresh > 1) { + const int pos = k_idxs[0]; + if (pos % refresh != 0) { + return; + } + } + const int h = blockIdx.x; // one CUDA block per (query-head, sequence) + const int s = blockIdx.y; + if (h >= n_head) { + return; + } + const int tid = threadIdx.x; + const int kvh = (gqa > 0) ? (h / gqa) : 0; + + // dynamic shared: [0,n_block) = this-query block exp sums; [n_block,2*n_block) = cross-query accumulator + extern __shared__ float smem[]; + float * bsum_q = smem; + float * bsum_acc = smem + n_block; + for (int b = tid; b < n_block; b += blockDim.x) { + bsum_acc[b] = 0.0f; + } + __syncthreads(); + + __shared__ float sred[SPARSE_SEL_NT]; + __shared__ float s_M, s_Z; + + int qr0 = n_q - q_win; + if (qr0 < 0) { + qr0 = 0; + } + const char * q_head = q_base + (int64_t) s*qnb3 + (int64_t) h*qnb1; + + for (int qr = qr0; qr < n_q; ++qr) { + const char * qrow = q_head + (int64_t) qr*qnb2; + // CAUSAL bound: a prefill-window query at row qr sits at absolute position pmax = n_kv - n_q + qr + // (the chunk's queries occupy the last n_q cache positions). It may only attend keys p <= pmax — + // including future keys would pollute the estimate. At decode (n_q==1) pmax = n_kv-1 (all keys). + const int pmax = n_kv - n_q + qr; + for (int b = tid; b < n_block; b += blockDim.x) { + bsum_q[b] = 0.0f; + } + __syncthreads(); + + // Pass 1: M = max_p (q·K[p]) for softmax stability + float lmax = -INFINITY; + for (int p = tid; p <= pmax; p += blockDim.x) { + const half * kcol = (const half *) (k_base + (int64_t) p*knb2 + (int64_t) kvh*knb1); + float dot = 0.0f; + for (int d = 0; d < head_dim; ++d) { + dot += (*(const float *) (qrow + (int64_t) d*4)) * __half2float(kcol[d]); + } + lmax = fmaxf(lmax, dot); + } + sred[tid] = lmax; + __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { if (tid < r) sred[tid] = fmaxf(sred[tid], sred[tid+r]); __syncthreads(); } + if (tid == 0) { s_M = sred[0]; } + __syncthreads(); + const float M = s_M; + + // Pass 2: Z = Σ exp(dot-M); per-block bsum_q[b] += exp(dot-M) (causal: p <= pmax) + float lz = 0.0f; + for (int p = tid; p <= pmax; p += blockDim.x) { + const half * kcol = (const half *) (k_base + (int64_t) p*knb2 + (int64_t) kvh*knb1); + float dot = 0.0f; + for (int d = 0; d < head_dim; ++d) { + dot += (*(const float *) (qrow + (int64_t) d*4)) * __half2float(kcol[d]); + } + const float e = __expf(dot - M); + lz += e; + atomicAdd(&bsum_q[p / B], e); + } + sred[tid] = lz; + __syncthreads(); + for (int r = blockDim.x/2; r > 0; r >>= 1) { if (tid < r) sred[tid] += sred[tid+r]; __syncthreads(); } + if (tid == 0) { s_Z = sred[0]; } + __syncthreads(); + const float Z = s_Z > 0.0f ? s_Z : 1.0f; + + // normalize this query's block mass and fold into the cross-query accumulator + for (int b = tid; b < n_block; b += blockDim.x) { + bsum_acc[b] += bsum_q[b] / Z; + } + __syncthreads(); + } + + // union across heads: atomic-max this head's block mass into scores. Mass >= 0, so the IEEE + // int-reinterpretation is monotone and atomicMax on the int view is a valid float max. + for (int b = tid; b < n_block; b += blockDim.x) { + atomicMax((int *) &scores[(int64_t) s*n_block + b], __float_as_int(bsum_acc[b])); + } +} + +void ggml_cuda_sparse_attn_vslash(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k_full = dst->src[1]; + + GGML_ASSERT(dst->type == GGML_TYPE_I32); + GGML_ASSERT(q->type == GGML_TYPE_F32); + GGML_ASSERT(k_full->type == GGML_TYPE_F16); + + int32_t p[7] = {0,0,0,0,0,1,1}; + memcpy(p, dst->op_params, sizeof(p)); + const int B = p[0]; + const int n_block = p[1]; + const int k_sel = p[2]; + const int recent_blocks = p[3]; + const int sink_blocks = p[4]; + const int refresh = p[5] < 1 ? 1 : p[5]; + const int q_win = p[6] < 1 ? 1 : p[6]; + + const int head_dim = q->ne[0]; + const int n_head = q->ne[1]; + const int n_q = q->ne[2]; + const int n_seq = q->ne[3]; + // get_k() returns K as [head_dim, n_head_kv, n_kv, 1]; the position axis is ne[2]/nb[2] (bug-743). + const int n_kv = (int) k_full->ne[2]; + const int n_head_kv = (int) k_full->ne[1]; + const int gqa = (n_head_kv > 0) ? (n_head / n_head_kv) : 1; + const int64_t knb1 = k_full->nb[1]; // per-KV-head stride + const int64_t knb2 = k_full->nb[2]; // per-position stride + + const int n_list = dst->ne[0]; + GGML_ASSERT(n_list >= n_block); + + const ggml_tensor * k_idxs = dst->src[2]; + const int * k_idxs_ptr = k_idxs ? (const int *) k_idxs->data : nullptr; + + if (n_block <= 0 || n_seq <= 0 || n_kv <= 0) { + return; + } + + ggml_cuda_pool_alloc scores(ctx.pool(), (size_t) n_block * n_seq); + // softmax block-mass accumulates across heads by atomic-max ⇒ scores must start at 0. + CUDA_CHECK(cudaMemsetAsync(scores.ptr, 0, (size_t) n_block * n_seq * sizeof(float), ctx.stream())); + + // one CUDA block per (query-head, seq); dynamic shared = 2·n_block floats (per-query + cross-query sums). + const dim3 grid_score(n_head, n_seq, 1); + const size_t smem = (size_t) 2 * n_block * sizeof(float); + vslash_score_kernel<<>>( + (const char *) q->data, (const char *) k_full->data, scores.ptr, k_idxs_ptr, + head_dim, n_head, n_q, n_kv, n_block, B, gqa, q_win, refresh, + q->nb[1], q->nb[2], q->nb[3], knb1, knb2); + + static const int sparse_dbg = getenv("OPENCOTI_SPARSE_DBG") ? atoi(getenv("OPENCOTI_SPARSE_DBG")) : 0; + sparse_attn_select_kernel<<>>( + scores.ptr, (int *) dst->data, k_idxs_ptr, n_block, n_list, k_sel, recent_blocks, sink_blocks, refresh, sparse_dbg); +} + +void ggml_cuda_sparse_attn_select(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * kbounds = dst->src[1]; + + GGML_ASSERT(dst->type == GGML_TYPE_I32); + GGML_ASSERT(q->type == GGML_TYPE_F32); + GGML_ASSERT(kbounds->type == GGML_TYPE_F16); + GGML_ASSERT(kbounds->ne[2] == 2); + + int32_t p[7] = {0,0,0,0,0,1,1}; + memcpy(p, dst->op_params, sizeof(p)); + const int n_block = p[1]; + const int k_sel = p[2]; + const int recent_blocks = p[3]; + const int sink_blocks = p[4]; + const int refresh = p[5] < 1 ? 1 : p[5]; + const int sub_per_block = p[6] < 1 ? 1 : p[6]; // #551 sub-block Quest bound (1 = plain B-block) + + const int head_dim = q->ne[0]; + const int n_head = q->ne[1]; + const int n_q = q->ne[2]; + const int n_seq = q->ne[3]; + const int n_embd_k = kbounds->ne[0]; + // #551: list row stride. Transient dst is [n_block, n_seq]; the persistent amortized buffer is + // [n_block_cap+1, n_seq] (capacity >= current n_block), so n_list >= n_block in general. + const int n_list = dst->ne[0]; + GGML_ASSERT(n_list >= n_block); + + // #551 PERF: persistent-buffer path attaches the KV write indices as src[2] (decode position source for + // the in-kernel refresh gate). nullptr ⇒ transient per-step selection (refresh forced to 1 in ggml.c). + const ggml_tensor * k_idxs = dst->src[2]; + const int * k_idxs_ptr = k_idxs ? (const int *) k_idxs->data : nullptr; + + if (n_block <= 0 || n_seq <= 0) { + return; + } + + ggml_cuda_pool_alloc scores(ctx.pool(), (size_t) n_block * n_seq); + + // #551 bug-743b: the max plane base is the kbounds ROW count (ne[1] = ceil(n_ctx/B_alloc)), which fill + // uses; the score previously read it at the dynamic walk n_block (= ceil(n_kv/B)) → max read landed in the + // MIN plane / garbage whenever n_kv < n_ctx (always), corrupting the Quest bound → near-random selection. + const int n_sub = (int) kbounds->ne[1]; + const dim3 grid_score(n_block, n_seq, 1); + sparse_attn_score_kernel<<>>( + (const char *) q->data, (const half *) kbounds->data, scores.ptr, k_idxs_ptr, + head_dim, n_head, n_q, n_embd_k, n_block, n_sub, sub_per_block, refresh, + q->nb[1], q->nb[2], q->nb[3]); + + // opencoti #551 bug-743 DIAG: OPENCOTI_SPARSE_DBG>0 prints the actual kept-count from the select kernel. + static const int sparse_dbg = getenv("OPENCOTI_SPARSE_DBG") ? atoi(getenv("OPENCOTI_SPARSE_DBG")) : 0; + sparse_attn_select_kernel<<>>( + scores.ptr, (int *) dst->data, k_idxs_ptr, n_block, n_list, k_sel, recent_blocks, sink_blocks, refresh, sparse_dbg); +} diff --git a/bk-t1/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cuh b/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cuh new file mode 100644 index 0000000..5cb7b0e --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/sparse-attn.cuh @@ -0,0 +1,19 @@ +#pragma once + +// opencoti-hook: sparse-attn — #551 Quest block-selector. See docs/features/sparse_attn.md. +#include "common.cuh" + +// GGML_OP_SPARSE_ATTN_FILL forward: reduce the (already-written) cache K into per-block +// channel-wise min/max, stored in the kbounds side-cache [n_embd_k, n_block, 2]. +void ggml_cuda_sparse_attn_fill(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + +// GGML_OP_SPARSE_ATTN_SELECT forward: per block, score = max over query-heads/rows of the +// Quest upper bound Σ_d max(q_d·kmin_d, q_d·kmax_d); then top-K_SEL blocks per sequence +// (union over heads) OR recent/sink window → I32-packed bitmask dst[n_words, n_seq]. +void ggml_cuda_sparse_attn_select(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + +// opencoti-hook: sparse-attn-vslash — #551 content-driven block selector. GGML_OP_SPARSE_ATTN_VSLASH +// forward: per block, score = max over (last q_win query rows, key in block, head) of the actual logit +// q·key (using cache K src[1], not a kbounds bound); then reuses the select kernel's top-K + recent/sink +// sorted-compaction emit. See sparse-attn.cu. +void ggml_cuda_sparse_attn_vslash(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh b/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh index 5819745..b9cc5c0 100644 --- a/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh @@ -22,3 +22,12 @@ void tcq3_set_decode_codebook_t3t2(const float * cb); // turbo3_tcq-K / turbo2_ void tcq2_set_decode_codebook_t3t2(const float * cb); // turbo3_tcq-K / turbo2_tcq-V TU: its d_turbo2 (256) void tcq3_set_decode_codebook_t2t3(const float * cb); // turbo2_tcq-K / turbo3_tcq-V TU: its d_turbo3 (512) void tcq2_set_decode_codebook_t2t3(const float * cb); // turbo2_tcq-K / turbo3_tcq-V TU: its d_turbo2 (256) + +// opencoti-hook: TCQ A′ — MMA-lift materialize codebook (#553/bug-710). The dequant-on-lift kernels +// cpy_turbo{3,2}_tcq_f16 (turbo-cpy.cu) decode TCQ→f16 for the batched/MMA attention path (prefill + +// spec-verify) using d_turbo{3,2}_tcq_cb_cpy — a THIRD per-TU static __constant__ copy that the decode +// setters above DON'T reach. Defined in turbo-cpy.cu; set-rows.cu's TURBO_TCQ_CB / CB2 loader must call +// these too, else the lift keeps the baked codebook while FA-VEC decode uses the override → draft(VEC) +// vs verify(MMA-lift) decode the same KV bytes differently (acceptance collapse / KLD split). +void tcq3_set_cpy_codebook(const float * cb); // 512 floats → turbo-cpy.cu d_turbo3_tcq_cb_cpy +void tcq2_set_cpy_codebook(const float * cb); // 256 floats → turbo-cpy.cu d_turbo2_tcq_cb_cpy diff --git a/bk-t1/llama.cpp/ggml/src/ggml-cuda/turbo-cpy-tcq-cb.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy-tcq-cb.cuh new file mode 100644 index 0000000..15f3eba --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy-tcq-cb.cuh @@ -0,0 +1,110 @@ +#pragma once +// opencoti-hook: TCQ->f16 dequant-on-lift codebooks (#533/#505) — see docs/features/poly_kv.md. +// VERBATIM byte-identical copy of the decode codebooks in fattn-tcq.cuh (extracted, not edited), +// so the prefill-MMA materialize path (turbo-cpy.cu cpy_turbo{3,2}_tcq_f16) reconstructs the SAME +// true-space values the fused FA-VEC decode path reads. Renamed *_cb_cpy to keep a private per-TU +// __constant__ (the TURBO_TCQ_CB A'-override targets fattn-tcq.cuh's symbol only and does NOT reach +// this TU; the shipped/default codebook is baked here — A'-sweep + prefill-MMA is a documented gap). +// Regenerate (if codebook ever changes): .opencoti/m533-gen-tcq-cb.sh + +static __constant__ float d_turbo3_tcq_cb_cpy[512] = { + -0.14559399f, -0.09062801f, -0.054925077f, -0.03699251f, -0.006363985f, +0.026264573f, +0.067378916f, +0.121981815f, + -0.18648055f, -0.106522456f, -0.052047577f, -0.011695214f, +0.021953275f, +0.059698727f, +0.09831437f, +0.16083933f, + -0.16390342f, -0.12639847f, -0.09513180f, -0.05938352f, -0.028396897f, +0.005973862f, +0.049104784f, +0.11334257f, + -0.25952467f, -0.079778515f, -0.036024813f, +0.0003641268f, +0.031858794f, +0.073280424f, +0.11835553f, +0.19738495f, + -0.14218009f, -0.10224814f, -0.062498566f, -0.027066832f, +0.00393002f, +0.04069300f, +0.08257346f, +0.14548601f, + -0.18673635f, -0.13438253f, -0.088401966f, -0.05205436f, -0.02032501f, +0.012399545f, +0.05127183f, +0.10316186f, + -0.10807011f, -0.065903045f, -0.032206114f, -0.0062006037f, +0.020679146f, +0.04422085f, +0.08313074f, +0.16821936f, + -0.22979105f, -0.14431947f, -0.07689272f, -0.02755307f, +0.009225173f, +0.046684854f, +0.08834142f, +0.13766693f, + -0.22114082f, -0.12612148f, -0.06890522f, -0.016128855f, +0.03691900f, +0.08474852f, +0.14940020f, +0.23229980f, + -0.14933491f, -0.099693604f, -0.06738499f, -0.037100967f, -0.009332986f, +0.023535024f, +0.060272533f, +0.109464675f, + -0.20200425f, -0.07398328f, -0.038700905f, -0.01714807f, +0.011161969f, +0.04528101f, +0.08902637f, +0.19573534f, + -0.16645233f, -0.124482535f, -0.089342155f, -0.04427387f, -0.007353691f, +0.028033108f, +0.066108435f, +0.15552913f, + -0.22295763f, -0.059887577f, -0.018804537f, +0.020141022f, +0.059682943f, +0.097920544f, +0.14080113f, +0.25698325f, + -0.14248224f, -0.089685425f, -0.050101686f, -0.017257255f, +0.011412255f, +0.040830314f, +0.07400172f, +0.11997315f, + -0.18649384f, -0.113997504f, -0.067775466f, -0.033394672f, +0.006586988f, +0.05312057f, +0.10433043f, +0.22344802f, + -0.16138338f, -0.108194515f, -0.07600300f, -0.05135381f, -0.023365447f, +0.0087320795f, +0.045431953f, +0.09113002f, + -0.12630440f, -0.07225349f, -0.032280035f, +0.0029231994f, +0.019239848f, +0.05081419f, +0.077840395f, +0.121695265f, + -0.08928155f, -0.044983763f, -0.009889568f, +0.020831043f, +0.05684458f, +0.09409702f, +0.13867535f, +0.19084482f, + -0.14182915f, -0.11380146f, -0.06904074f, -0.002002765f, +0.034864165f, +0.070399575f, +0.11403063f, +0.15394832f, + -0.10876417f, -0.056122433f, -0.02267638f, +0.011113975f, +0.039639056f, +0.074084364f, +0.10155376f, +0.12540291f, + -0.17693359f, -0.13940524f, -0.10049578f, -0.06796275f, -0.036915872f, +0.00062823476f, +0.042142134f, +0.17906062f, + -0.09253492f, -0.04290128f, -0.006311852f, +0.023908244f, +0.049849935f, +0.078770354f, +0.10818172f, +0.15166481f, + -0.12429565f, -0.07392063f, -0.029114135f, +0.0059440783f, +0.042675965f, +0.08425635f, +0.13836108f, +0.18634140f, + -0.11795639f, -0.07033707f, -0.034163877f, -0.0008773357f, +0.03334606f, +0.07188203f, +0.12216825f, +0.17097956f, + -0.18718453f, -0.14090346f, -0.097799584f, -0.059522875f, -0.019208657f, +0.03079176f, +0.09334672f, +0.15811224f, + -0.27198875f, -0.16546582f, -0.11433405f, -0.06933013f, -0.04026183f, -0.0061146915f, +0.029263576f, +0.07322499f, + -0.18471734f, -0.102074504f, -0.06492570f, -0.034418534f, -0.009636157f, +0.023043344f, +0.05751496f, +0.09905984f, + -0.22826399f, -0.15946552f, -0.09913176f, -0.06585259f, -0.03252090f, +0.001313243f, +0.03556729f, +0.21612854f, + -0.13243781f, -0.087299444f, -0.049820945f, -0.016216082f, +0.01799807f, +0.057916876f, +0.09001349f, +0.13221787f, + -0.19516511f, -0.120894566f, -0.076130204f, -0.051442243f, -0.029535033f, -0.0020043184f, +0.029452588f, +0.075566076f, + -0.27272871f, -0.15841717f, -0.105432935f, -0.06792948f, -0.024532158f, +0.014960791f, +0.054415092f, +0.101517834f, + -0.21153601f, -0.15015371f, -0.08676790f, -0.04414934f, -0.0042129597f, +0.033762872f, +0.07589151f, +0.12768789f, + -0.090428725f, -0.037582967f, +0.0013173596f, +0.03900247f, +0.06840049f, +0.116906695f, +0.16584939f, +0.25382105f, + -0.13446195f, -0.07865091f, -0.039625354f, -0.0028398742f, +0.03019514f, +0.06799379f, +0.11850997f, +0.17521496f, + -0.11350345f, -0.058599845f, -0.017512511f, +0.019431496f, +0.055897832f, +0.093173414f, +0.14820710f, +0.22092152f, + -0.15165758f, -0.08869354f, -0.04974287f, -0.01705474f, +0.013134752f, +0.04367713f, +0.07733791f, +0.12430801f, + -0.09329869f, -0.04673005f, -0.00045857552f, +0.042781368f, +0.07802363f, +0.11887439f, +0.16250038f, +0.28612965f, + -0.12571070f, -0.07786012f, -0.03843933f, -0.0075433915f, +0.025822964f, +0.066053316f, +0.12021536f, +0.18341768f, + -0.16079275f, -0.04921760f, -0.006114644f, +0.026215268f, +0.05699377f, +0.09813471f, +0.16080129f, +0.23786584f, + -0.09980837f, -0.048535258f, -0.0096120685f, +0.025387142f, +0.05979822f, +0.09875251f, +0.14474337f, +0.20324114f, + -0.15846540f, -0.09938028f, -0.061492465f, -0.03523542f, -0.0061364113f, +0.024916094f, +0.06037314f, +0.106796466f, + -0.20557843f, -0.123237535f, -0.07734871f, -0.044549115f, -0.017114898f, +0.01616654f, +0.049574375f, +0.092319444f, + -0.19221115f, -0.14642999f, -0.091701314f, -0.055265956f, -0.021026207f, +0.017720066f, +0.05786183f, +0.110154524f, + -0.09956386f, -0.03870283f, +0.003052007f, +0.034851722f, +0.06256365f, +0.09628840f, +0.13979156f, +0.16582295f, + -0.18026546f, -0.12448310f, -0.07424377f, -0.03954519f, -0.01221123f, +0.028641058f, +0.100819774f, +0.18240699f, + -0.21520759f, -0.15573645f, -0.09820838f, -0.051450998f, -0.012993679f, +0.021135861f, +0.058727216f, +0.105848536f, + -0.11207385f, -0.08335689f, -0.048542723f, -0.023198519f, +0.0039304253f, +0.037778318f, +0.07813917f, +0.13106476f, + -0.17849164f, -0.120988995f, -0.078016765f, -0.043093704f, -0.016565649f, +0.015182641f, +0.050754096f, +0.09595712f, + -0.22132620f, -0.13407415f, -0.065785654f, -0.013291034f, +0.032098345f, +0.07478225f, +0.12431934f, +0.19174045f, + -0.095454164f, -0.051898945f, -0.015116375f, -0.012596778f, +0.018636847f, +0.05006925f, +0.087654814f, +0.13754296f, + -0.15254061f, -0.09576059f, -0.052086458f, -0.01596074f, +0.017607626f, +0.04778498f, +0.08950204f, +0.14901252f, + -0.26057002f, -0.12472382f, -0.074396215f, -0.03764066f, +0.0011168446f, +0.061569117f, +0.10793752f, +0.19771695f, + -0.08661132f, -0.045195263f, -0.016098704f, +0.012780116f, +0.040476497f, +0.074102715f, +0.074102715f, +0.12635531f, + -0.14047913f, -0.059587404f, -0.016261123f, +0.019801628f, +0.053541403f, +0.096650146f, +0.15005490f, +0.21051759f, + -0.22986396f, -0.11964334f, -0.07266585f, -0.026522418f, +0.018169926f, +0.058630653f, +0.100647695f, +0.15919648f, + -0.13251697f, -0.077567816f, -0.042766172f, -0.011389967f, +0.01831755f, +0.05304656f, +0.09620367f, +0.15567583f, + -0.119819686f, -0.06772876f, -0.028123451f, +0.00876240f, +0.014405836f, +0.048829112f, +0.08422175f, +0.13823749f, + -0.16379014f, -0.08956941f, -0.041652776f, +0.008921398f, +0.05473602f, +0.10037984f, +0.16022855f, +0.23457925f, + -0.115844205f, -0.05939626f, -0.020390417f, +0.01374377f, +0.044976473f, +0.07873563f, +0.12207942f, +0.18412720f, + -0.19048831f, -0.07587487f, -0.03220580f, -0.00011795067f, +0.02721784f, +0.04380719f, +0.07886723f, +0.13193911f, + -0.13935551f, -0.092902906f, -0.052706074f, -0.017797327f, +0.015312965f, +0.056098964f, +0.11203423f, +0.24448302f, + -0.17986591f, -0.10738580f, -0.06376371f, -0.026595421f, +0.00842492f, +0.04272362f, +0.08608052f, +0.15240218f, + -0.10953678f, -0.057022586f, -0.012483291f, +0.024463262f, +0.06076792f, +0.09776234f, +0.12983681f, +0.18648379f, + -0.16471463f, -0.089491285f, -0.037574016f, +0.004444791f, +0.039293647f, +0.07845859f, +0.12893885f, +0.23508036f +}; + +static __constant__ float d_turbo2_tcq_cb_cpy[256] = { + -0.18030643f, -0.11009848f, -0.04742626f, +0.02894132f, -0.10523465f, -0.031312924f, +0.031491395f, +0.12263535f, + -0.15660362f, -0.055477407f, +0.0046675834f, +0.06166081f, -0.07506216f, -0.016963918f, +0.043737844f, +0.116496615f, + -0.08632783f, -0.022493735f, +0.041032985f, +0.10660284f, -0.06274858f, -0.0036939639f, +0.02095157f, +0.07539709f, + -0.09802641f, -0.008419088f, +0.059072323f, +0.17311879f, -0.093109086f, -0.02654333f, +0.014827672f, +0.07793592f, + -0.031235758f, +0.01271591f, +0.08752262f, +0.17246453f, -0.14595252f, -0.07227624f, +0.013628688f, +0.08131674f, + -0.036909282f, +0.0018896917f, +0.05209119f, +0.12407892f, -0.13689458f, -0.06054520f, +0.0064648795f, +0.07551241f, + -0.18980840f, -0.110128626f, -0.046503957f, +0.026387159f, -0.034967307f, +0.04810357f, +0.072072044f, +0.14355458f, + -0.10182410f, -0.02907887f, +0.014033012f, +0.083419636f, -0.056140676f, +0.008405868f, +0.066070884f, +0.14037225f, + -0.117427245f, -0.047159385f, +0.016928354f, +0.08142885f, -0.029359628f, +0.045608785f, +0.10559447f, +0.20061271f, + -0.040425077f, +0.029068163f, +0.08408973f, +0.13628258f, -0.16633821f, -0.10711727f, -0.04196669f, +0.027895834f, + -0.0054065837f, +0.058898676f, +0.12688550f, +0.18268861f, -0.16287325f, -0.11218357f, -0.07165227f, -0.009524379f, + -0.24026902f, -0.073219374f, -0.0005165726f, +0.05959821f, -0.05532953f, +0.027044486f, +0.09425678f, +0.15356481f, + -0.14381111f, -0.10563502f, -0.037867088f, +0.023611993f, -0.03624307f, +0.049588434f, +0.12192037f, +0.23462485f, + -0.14990251f, -0.09659304f, -0.05886742f, +0.014878461f, -0.009889551f, +0.06910514f, +0.12120181f, +0.22596690f, + -0.08290075f, -0.009009629f, +0.066151775f, +0.12188313f, -0.11591514f, -0.06952189f, -0.031633306f, +0.023740824f, + -0.20510401f, -0.103369795f, +0.09148037f, +0.17268716f, -0.16597997f, -0.09207068f, -0.032810967f, +0.024847647f, + -0.02487482f, +0.049298953f, +0.09624215f, +0.14217524f, -0.18418685f, -0.10147012f, -0.05841265f, +0.008057022f, + -0.14269894f, -0.092456274f, -0.026881337f, +0.049792137f, -0.019881032f, +0.030333601f, +0.09736802f, +0.17764080f, + -0.19579841f, -0.114739306f, -0.026823774f, +0.07466014f, -0.09001050f, -0.041468445f, +0.028473806f, +0.08870695f, + -0.019396419f, +0.042828932f, +0.10885327f, +0.13335012f, -0.15005013f, -0.074581385f, -0.028608415f, +0.03848942f, + -0.09687270f, -0.057059396f, +0.0077843578f, +0.06302297f, -0.23247094f, -0.14509225f, -0.032651436f, +0.027010715f, + -0.047595482f, +0.06280303f, +0.114691675f, +0.17124057f, -0.21092793f, -0.13704823f, -0.07340412f, +0.0039013291f, + -0.062834196f, +0.012601906f, +0.012601906f, +0.08721347f, -0.13256435f, -0.024173854f, +0.07723171f, +0.14801070f, + -0.06471605f, -0.0017903054f, -0.0017903054f, +0.058302354f, -0.09731802f, -0.03400696f, +0.02762442f, +0.08986137f, + -0.08288722f, -0.019051429f, +0.045709886f, +0.15211061f, -0.09507891f, -0.015612489f, +0.025347246f, +0.087257534f, + -0.066236064f, -0.0047936034f, +0.06386274f, +0.15401669f, -0.105809286f, -0.051802177f, +0.01073050f, +0.08292137f, + -0.11884470f, -0.04404144f, +0.02550729f, +0.02550729f, -0.01731189f, +0.062161792f, +0.12127554f, +0.21981733f, + -0.17066145f, -0.11660990f, -0.049425896f, +0.021293938f, -0.04711412f, +0.026577346f, +0.055197213f, +0.12541275f, + -0.028268812f, +0.015206398f, +0.09002519f, +0.12699963f, -0.10059831f, -0.026676945f, +0.059903253f, +0.13054545f, + -0.09582803f, -0.033371232f, +0.010346129f, +0.066766635f, -0.09964944f, -0.028686784f, +0.021184925f, +0.09120017f, + -0.16957201f, -0.07594450f, +0.04172865f, +0.18313301f, -0.051526368f, +0.011877304f, +0.011877304f, +0.07956263f, + -0.13432936f, -0.05269006f, +0.03536416f, +0.117640756f, -0.022776067f, +0.042032316f, +0.10472976f, +0.18042557f +}; diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu index cf7a5dc..41611a4 100644 --- a/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu @@ -10,6 +10,20 @@ #include "common.cuh" #include "turbo-dequant.cuh" // TURBO{3}_WHT_SIGNS1/2 + TURBO{2,3,4}_CENTROIDS tables (this TU's copy) #include "turbo-cpy.cuh" +#include "turbo-cpy-tcq-cb.cuh" // opencoti-hook #533: TCQ decode codebooks (this TU's copy, == fattn-tcq.cuh) +#include "tcq-decode-cb.cuh" // opencoti-hook #553: declares tcq{3,2}_set_cpy_codebook (defined below) + +// opencoti-hook: TCQ A′ — MMA-lift codebook override (#553/bug-710). d_turbo{3,2}_tcq_cb_cpy are static +// __constant__ in THIS TU (turbo-cpy-tcq-cb.cuh); cudaMemcpyToSymbol on a static symbol only reaches the +// in-TU copy, so the TURBO_TCQ_CB / CB2 runtime override (set-rows.cu) must call these to apply to the +// dequant-on-lift path, mirroring the FA-VEC tcq{3,2}_set_decode_codebook* setters. Without it the lift +// (prefill + spec-verify) keeps the baked codebook while FA-VEC decode honors the override → mismatch. +void tcq3_set_cpy_codebook(const float * cb) { + CUDA_CHECK(cudaMemcpyToSymbol(d_turbo3_tcq_cb_cpy, cb, 512 * sizeof(float))); +} +void tcq2_set_cpy_codebook(const float * cb) { + CUDA_CHECK(cudaMemcpyToSymbol(d_turbo2_tcq_cb_cpy, cb, 256 * sizeof(float))); +} // Cooperative inverse WHT over shared buf[128] (already holding centroid×norm in rotated space), // then InnerQ ×scale_inv and a strided f16 write of element e0+t. All 128 threads of the block @@ -108,6 +122,40 @@ static __global__ void cpy_turbo8_f16(TURBO_CPY_PARAMS) { turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, scale_inv, head_dim); } +// opencoti-hook: TCQ→f16 dequant-on-lift (#533/#505) — prefill-MMA fix + DCA-compose. +// TCQ was fused-only (no f16 materialize), forcing prefill through FA-VEC (no tensor cores) +// and aborting under DCA's f16-lift. These two kernels mirror cpy_turbo3_f16 exactly but decode +// the trellis bitstream via the 512/256-entry codebook (byte-identical to fattn-tcq.cuh's reader), +// then the SAME cooperative inverse-FWHT recovers TRUE-space f16. TCQ carries NO InnerQ, so +// scale_inv=nullptr. norm is used raw — the per-side α (K×1.0/V×1.04) is already baked into norm at +// encode (turbo-tcq-cuda.cuh:359) and the decode-time α defaults to 1.0 (no double-apply), so this +// matches the FA-VEC reader at the shipped codebook/α. Bit extract == fattn-tcq.cuh:150-152 / 205-210. +static __global__ void cpy_turbo3_tcq_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo3_tcq); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + const int bp = t * 3; + const uint16_t raw = (uint16_t) blk->qs[bp / 8] | ((uint16_t) blk->qs[bp / 8 + 1] << 8); + buf[t] = d_turbo3_tcq_cb_cpy[(raw >> (bp % 8)) & 0x1FF] * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, nullptr, head_dim); +} + +static __global__ void cpy_turbo2_tcq_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo2_tcq); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + const int bp = t * 2; + const uint16_t raw = (uint16_t) blk->qs[bp / 8] | ((uint16_t) blk->qs[bp / 8 + 1] << 8); + buf[t] = d_turbo2_tcq_cb_cpy[(raw >> (bp % 8)) & 0xFF] * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, nullptr, head_dim); +} + void ggml_cpy_turbo_f16_cuda( enum ggml_type src_type, const char * cx, char * cdst, int64_t ne, int64_t ne00, int64_t ne01, int64_t ne02, @@ -127,6 +175,8 @@ void ggml_cpy_turbo_f16_cuda( case GGML_TYPE_TURBO4_0: LAUNCH(cpy_turbo4_f16); break; case GGML_TYPE_TURBO2_0: LAUNCH(cpy_turbo2_f16); break; case GGML_TYPE_TURBO8_0: LAUNCH(cpy_turbo8_f16); break; + case GGML_TYPE_TURBO3_TCQ: LAUNCH(cpy_turbo3_tcq_f16); break; // opencoti-hook #533 + case GGML_TYPE_TURBO2_TCQ: LAUNCH(cpy_turbo2_tcq_f16); break; // opencoti-hook #533 default: GGML_ABORT("ggml_cpy_turbo_f16_cuda: non-turbo src_type %d\n", (int) src_type); } #undef LAUNCH diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh index 887aac2..3f38789 100644 --- a/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh @@ -202,8 +202,14 @@ static __global__ void __launch_bounds__(TCQ3_WPB*32, 3) k_set_rows_turbo3_tcq( // each group's 8 predecessors (g*8..g*8+7) fall entirely within lane g/2, so the predecessor-min is // thread-local (no shuffle). A 64-entry per-warp smem broadcast under __syncwarp replaces the block barrier. acc_t c[16]; + // opencoti-hook: L3 (#541) decode-tps — hoist the loop-invariant TCQ codebook reads out of the 128-step + // Viterbi DP. d_turbo3_tcq_codebook[16*lane+i] depends only on (lane,i), never on t, yet the warp-sync + // design re-read it from __constant__ every timestep with 32-way per-warp serialization (the #534 + // anti-pattern, here in the encode). Loading each lane's 16 entries into registers ONCE is byte-identical + // (same loads, just hoisted) and removes 127×16 serialized constant reads/group. + acc_t cb[16]; #pragma unroll - for (int i = 0; i < 16; i++) c[i] = 0.0f; + for (int i = 0; i < 16; i++) { c[i] = 0.0f; cb[i] = d_turbo3_tcq_codebook[16*lane + i]; } for (int t = 0; t < 128; t++) { acc_t pm0 = c[0]; int pp0 = 0; #pragma unroll @@ -216,7 +222,7 @@ static __global__ void __launch_bounds__(TCQ3_WPB*32, 3) k_set_rows_turbo3_tcq( __syncwarp(); acc_t xt = x[t]; #pragma unroll - for (int i = 0; i < 16; i++) { int s = 16*lane + i; acc_t d = xt - d_turbo3_tcq_codebook[s]; c[i] = pred_min_cost[s & 0x3F] + d * d; } + for (int i = 0; i < 16; i++) { int s = 16*lane + i; acc_t d = xt - cb[i]; c[i] = pred_min_cost[s & 0x3F] + d * d; } __syncwarp(); } // --- argmin over 512 states (per-lane min over 16, then warp-shuffle) --- @@ -427,8 +433,9 @@ static __global__ void __launch_bounds__(TCQ2_WPB*32, 3) k_set_rows_turbo2_tcq( // lane L owns states {8L..8L+7} = trellis-groups 2L (states 8L..8L+3) and 2L+1 (states 8L+4..8L+7); each // group's 4 predecessors (g*4..g*4+3) fall within lane g/2, so predecessor-min is thread-local (no shuffle). acc_t c[8]; + acc_t cb[8]; // opencoti-hook: L3 (#541) — hoist loop-invariant codebook out of the Viterbi DP (see turbo3 note above) #pragma unroll - for (int i = 0; i < 8; i++) c[i] = 0.0f; + for (int i = 0; i < 8; i++) { c[i] = 0.0f; cb[i] = d_turbo2_tcq_codebook[8*lane + i]; } for (int t = 0; t < 128; t++) { acc_t pm0 = c[0]; int pp0 = 0; #pragma unroll @@ -441,7 +448,7 @@ static __global__ void __launch_bounds__(TCQ2_WPB*32, 3) k_set_rows_turbo2_tcq( __syncwarp(); acc_t xt = x[t]; #pragma unroll - for (int i = 0; i < 8; i++) { int s = 8*lane + i; acc_t d = xt - d_turbo2_tcq_codebook[s]; c[i] = pred_min_cost[s & 0x3F] + d * d; } + for (int i = 0; i < 8; i++) { int s = 8*lane + i; acc_t d = xt - cb[i]; c[i] = pred_min_cost[s & 0x3F] + d * d; } __syncwarp(); } // --- argmin over 256 states (per-lane min over 8, then warp-shuffle) --- diff --git a/llama.cpp/ggml/src/ggml.c b/llama.cpp/ggml/src/ggml.c index 2defaae..261d69e 100644 --- a/llama.cpp/ggml/src/ggml.c +++ b/llama.cpp/ggml/src/ggml.c @@ -1158,9 +1158,12 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "FLASH_ATTN_TAIL_PARTIAL", "TURBO_WHT", "FLASH_ATTN_EXT_LSE", + "SPARSE_ATTN_FILL", + "SPARSE_ATTN_SELECT", + "SPARSE_ATTN_VSLASH", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 104, "GGML_OP_COUNT != 104"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1273,9 +1276,12 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "flash_attn_tail_partial(q,k,v)", "turbo_wht(x)", "flash_attn_ext_lse(q,k,v)", + "sparse_attn_fill(kb,k)", + "sparse_attn_select(q,kb)", + "sparse_attn_vslash(q,k)", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 104, "GGML_OP_COUNT != 104"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -5821,6 +5827,136 @@ struct ggml_tensor * ggml_turbo_wht( return result; } +// opencoti-hook: sparse-attn — #551 Quest per-block kmin/kmax fill. See docs/features/sparse_attn.md. +struct ggml_tensor * ggml_sparse_attn_fill( + struct ggml_context * ctx, + struct ggml_tensor * kbounds, + struct ggml_tensor * k_cur, + struct ggml_tensor * k_idxs, + int block_size) { + GGML_ASSERT(kbounds->type == GGML_TYPE_F16); + GGML_ASSERT(kbounds->ne[2] == 2); // [n_embd_k, n_block, 2] : [..,0]=min, [..,1]=max + GGML_ASSERT(block_size > 0); + + // Write in-place into the kbounds side-cache (ggml_cpy idiom): result aliases kbounds' + // buffer so the SPARSE_ATTN_FILL node carries the write edge that orders the selector + // after the fill. src[0]=kbounds doubles as the existing values to merge (incremental). + struct ggml_tensor * result = ggml_view_tensor(ctx, kbounds); + + result->op = GGML_OP_SPARSE_ATTN_FILL; + result->src[0] = kbounds; + result->src[1] = k_cur; + result->src[2] = k_idxs; + + memcpy(result->op_params, &block_size, sizeof(int)); + + return result; +} + +struct ggml_tensor * ggml_sparse_attn_select( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * kbounds, + struct ggml_tensor * block_sel, + struct ggml_tensor * k_idxs, + int block_size, + int n_block, + int k_sel, + int recent_blocks, + int sink_blocks, + int refresh, + int sub_per_block) { + GGML_ASSERT(kbounds->type == GGML_TYPE_F16); + GGML_ASSERT(kbounds->ne[2] == 2); + GGML_ASSERT(block_size > 0); + GGML_ASSERT(n_block > 0); + GGML_ASSERT(sub_per_block >= 1); + + // dst: #551 compacted SORTED, −1-padded list of selected block indices per sequence (q->ne[3]); union + // over KV-heads and query rows. FA-VEC iterates the list (block_sel[seq*stride + i]) and stops at the + // first −1 — O(selected), not O(capacity). + const int64_t n_seq = q->ne[3]; + struct ggml_tensor * result; + if (block_sel) { + // #551 PERF (amortized): write IN-PLACE into the persistent per-stream buffer (cpy idiom, like + // SPARSE_ATTN_FILL). block_sel is [n_block+1, n_seq] I32 (capacity n_block + a reserved scratch + // slot). The kernel re-selects only every `refresh` decode tokens (position read from k_idxs[0]), + // reusing the cached list otherwise — keeps CUDA-graph topology stable (no add/remove of nodes). + GGML_ASSERT(block_sel->type == GGML_TYPE_I32); + GGML_ASSERT(block_sel->ne[0] >= n_block); + GGML_ASSERT(k_idxs); + GGML_ASSERT(refresh >= 1); + result = ggml_view_tensor(ctx, block_sel); + result->src[2] = k_idxs; + } else { + // legacy transient per-step list [n_block, n_seq] (multi-stream fallback; selects every step). + result = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_block, n_seq); + refresh = 1; + } + + result->op = GGML_OP_SPARSE_ATTN_SELECT; + result->src[0] = q; + result->src[1] = kbounds; + + int32_t params[7] = { block_size, n_block, k_sel, recent_blocks, sink_blocks, refresh, sub_per_block }; + memcpy(result->op_params, params, sizeof(params)); + + return result; +} + +// opencoti-hook: sparse-attn-vslash — #551 content-driven block selector. Unlike the Quest selector +// (geometric per-dim kmin/kmax spread bound, which over-ranks high-spread filler and DEGRADES with +// context), vslash scores each block by the EXACT per-block max attention logit over a window of +// recent queries: score(b) = max over (q in last q_win rows, key p in block b, head h) of (q_h · key_p). +// This is the "vertical" importance of MInference (arXiv:2407.02490): the question's own attention +// picks the needle's key column. Reuses the select kernel's top-K + recent/sink compaction verbatim; +// only the score source differs (actual cache K, src[1], instead of kbounds). op_params = +// {block_size, n_block, k_sel, recent_blocks, sink_blocks, refresh, q_win}. +struct ggml_tensor * ggml_sparse_attn_vslash( + struct ggml_context * ctx, + struct ggml_tensor * q, + struct ggml_tensor * k_full, + struct ggml_tensor * block_sel, + struct ggml_tensor * k_idxs, + int block_size, + int n_block, + int k_sel, + int recent_blocks, + int sink_blocks, + int refresh, + int q_win) { + GGML_ASSERT(k_full->type == GGML_TYPE_F16); + GGML_ASSERT(block_size > 0); + GGML_ASSERT(n_block > 0); + GGML_ASSERT(q_win >= 1); + + const int64_t n_seq = q->ne[3]; + struct ggml_tensor * result; + if (block_sel) { + // #551 PERF (amortized): write IN-PLACE into the persistent per-stream buffer (cpy idiom, + // mirror of SPARSE_ATTN_SELECT). The kernel re-selects only every `refresh` decode tokens + // (position read from k_idxs[0] on the device), reusing the cached list otherwise. + GGML_ASSERT(block_sel->type == GGML_TYPE_I32); + GGML_ASSERT(block_sel->ne[0] >= n_block); + GGML_ASSERT(k_idxs); + GGML_ASSERT(refresh >= 1); + result = ggml_view_tensor(ctx, block_sel); + result->src[2] = k_idxs; + } else { + result = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_block, n_seq); + refresh = 1; + } + + result->op = GGML_OP_SPARSE_ATTN_VSLASH; + result->src[0] = q; + result->src[1] = k_full; + + int32_t params[7] = { block_size, n_block, k_sel, recent_blocks, sink_blocks, refresh, q_win }; + memcpy(result->op_params, params, sizeof(params)); + + return result; +} + void ggml_flash_attn_ext_set_prec( struct ggml_tensor * a, enum ggml_prec prec) { diff --git a/llama.cpp/include/llama.h b/llama.cpp/include/llama.h index b5be5f5..a0a435c 100644 --- a/llama.cpp/include/llama.h +++ b/llama.cpp/include/llama.h @@ -381,6 +381,17 @@ extern "C" { bool dca_enabled; uint32_t dca_chunk_size; float dca_yarn_factor; + // opencoti #551 sparse-attn — Quest-style block-selector sparse attention (docs/features/sparse_attn.md) + // sparse_attn_enabled: false = off (default). block_size: KV-block granularity (B_SEL, 64). + // topk: # blocks kept per query (0 = all/off, 0xFFFFFFFF = auto ¼ floor 64). recent: + // forced-recent window tokens (0 = none). sink: leading sink blocks always kept (default 1). + bool sparse_attn_enabled; + uint32_t sparse_attn_block_size; + uint32_t sparse_attn_topk; + uint32_t sparse_attn_recent; + uint32_t sparse_attn_sink; + uint32_t sparse_attn_refresh; // opencoti #551 PERF — re-select sparse blocks every N decode tokens (amortize selection machinery) + uint32_t sparse_attn_mode; // opencoti-hook: sparse-attn-vslash — #551 0 = Quest bound (default), 1 = vertical-slash content-driven 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/llama-context.cpp b/llama.cpp/src/llama-context.cpp index 2526f84..2d87414 100644 --- a/llama.cpp/src/llama-context.cpp +++ b/llama.cpp/src/llama-context.cpp @@ -114,6 +114,14 @@ llama_context::llama_context( cparams.dca_enabled = params.dca_enabled; cparams.dca_chunk_size = params.dca_chunk_size; cparams.dca_yarn_factor = params.dca_yarn_factor; + // opencoti #551 sparse-attn — see docs/features/sparse_attn.md + cparams.sparse_attn_enabled = params.sparse_attn_enabled; + cparams.sparse_attn_block_size = params.sparse_attn_block_size; + cparams.sparse_attn_topk = params.sparse_attn_topk; + cparams.sparse_attn_recent = params.sparse_attn_recent; + cparams.sparse_attn_sink = params.sparse_attn_sink; + cparams.sparse_attn_refresh = params.sparse_attn_refresh < 1 ? 1 : params.sparse_attn_refresh; + cparams.sparse_attn_mode = params.sparse_attn_mode; // opencoti-hook: sparse-attn-vslash — #551 cparams.n_threads = params.n_threads; cparams.n_threads_batch = params.n_threads_batch; @@ -4229,6 +4237,13 @@ llama_context_params llama_context_default_params() { /*.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 + /*.sparse_attn_enabled =*/ false, // opencoti #551 sparse-attn — off by default + /*.sparse_attn_block_size =*/ 64, // opencoti #551 sparse-attn — KV-block granularity B_SEL + /*.sparse_attn_topk =*/ 0, // opencoti #551 sparse-attn — 0 = all blocks (off) + /*.sparse_attn_recent =*/ 0, // opencoti #551 sparse-attn — 0 = no forced-recent window + /*.sparse_attn_sink =*/ 1, // opencoti #551 sparse-attn — leading sink blocks always kept + /*.sparse_attn_refresh =*/ 8, // opencoti #551 PERF — re-select every N decode tokens (amortize) + /*.sparse_attn_mode =*/ 0, // opencoti-hook: sparse-attn-vslash — #551 0 = Quest (default) /*.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 index 1aa52a5..41608b8 100644 --- a/llama.cpp/src/llama-cparams.h +++ b/llama.cpp/src/llama-cparams.h @@ -56,6 +56,16 @@ struct llama_cparams { bool dca_enabled; uint32_t dca_chunk_size; float dca_yarn_factor; + // opencoti #551 sparse-attn — Quest-style block-selector (docs/features/sparse_attn.md). + // sparse_attn_enabled: per-query top-K KV-block selection before QK (false = off, default). + // sparse_attn_block_size: KV-block granularity B_SEL. topk: # blocks kept (0 = all). recent: forced window. + bool sparse_attn_enabled; + uint32_t sparse_attn_block_size; + uint32_t sparse_attn_topk; + uint32_t sparse_attn_recent; + uint32_t sparse_attn_sink; + uint32_t sparse_attn_refresh; // #551 PERF: re-run the selector every N decode tokens (amortize; 1 = every step) + uint32_t sparse_attn_mode; // opencoti-hook: sparse-attn-vslash — #551 0 = Quest bound, 1 = vertical-slash content-driven 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 index 2fe40de..e0d5e2a 100644 --- a/llama.cpp/src/llama-graph.cpp +++ b/llama.cpp/src/llama-graph.cpp @@ -22,6 +22,12 @@ #include #include +// opencoti-hook: sparse-attn #551 — carries the Quest block-selector bitmask from build_attn's +// store block (where the fill result is available, so the read edge orders select-after-fill) +// into build_attn_mha, which attaches it as the flash_attn_ext op's src[7]. Set tightly around +// the standard build_attn_mha call and cleared after, so it never leaks into other FA paths. +static thread_local ggml_tensor * opencoti_sparse_block_sel = nullptr; + // dedup helpers static ggml_tensor * build_attn_inp_kq_mask( @@ -2094,24 +2100,26 @@ ggml_tensor * llm_graph_context::build_attn_mha( const bool fused_turbo_512 = k->ne[0] == 512 && (k->type == GGML_TYPE_TURBO2_0 || k->type == GGML_TYPE_TURBO3_0 || k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ); // opencoti TCQ #447/#500 (Gemma-4 GLOBAL) - // opencoti-hook: TCQ FA-VEC decode (#447/#500). TCQ has NO materialize-lift (turbo-cpy.cu has no - // TCQ→f16 inverse-FWHT cast yet), so it must ALWAYS fuse — prefill (large n_q) rides the VEC-tiled - // path (correct, slower than MMA; a prefill-MMA dequant is a deferred perf follow-on, like buun's - // k_turbo3_tcq_dequant_f16). Without this, a prefill ubatch would fall to the else-branch with no - // cast and emit raw TCQ to FA with an UN-rotated Q → wrong logits. - const bool is_tcq_kv = (k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ); // opencoti-hook: TCQ FA-VEC (#447/#500) — no materialize-lift, always fuse - // opencoti-hook: asymmetric TCQ (#513). TCQ has no f16 materialize-lift, so the only correct path is - // the in-register fused FA-VEC read. Allow K!=V when BOTH sides are TCQ (e.g. turbo3_tcq K / turbo2_tcq - // V): K and V are FWHT-rotated identically, so the same Q-rotation + output inverse-WHT stay logit-exact - // regardless of per-side trellis bit-width. Backed by the dedicated asymmetric FA-VEC instances - // (fattn-vec-instance-turbo3_tcq-turbo2_tcq + reverse); the CUDA selector K!=V gate is relaxed to match. + // opencoti-hook #533: TCQ now HAS a materialize-lift (turbo-cpy.cu cpy_turbo{3,2}_tcq_f16: trellis + // codebook decode + inverse-FWHT → TRUE-space f16). So TCQ follows the SAME n_q threshold as plain + // turbo — fuse the in-register FA-VEC at decode (small n_q), de-fuse to the else-branch (ggml_cast + // TCQ→f16 → FA-MMA tensor cores) at prefill (large n_q). This recovers prefill throughput (the forced + // FA-VEC collapse, ~376→MMA-class) and lets TCQ compose with DCA's #444 all-KV f16-lift. The + // else-branch cast de-rotates K,V fully to TRUE space, so the UN-rotated Q there is logit-exact + // (Parseval). Previously (#447/#500) "|| is_tcq_kv" force-fused TCQ everywhere because no lift existed. + // opencoti-hook: asymmetric TCQ (#513). Allow K!=V when BOTH sides are TCQ (e.g. turbo3_tcq K / + // turbo2_tcq V). At DECODE this fuses the in-register FA-VEC read: K and V are FWHT-rotated identically, + // so the same Q-rotation + output inverse-WHT stay logit-exact regardless of per-side trellis bit-width + // (backed by the dedicated asymmetric FA-VEC instances fattn-vec-instance-turbo3_tcq-turbo2_tcq + + // reverse; the CUDA selector K!=V gate is relaxed to match). At PREFILL each side de-fuses and materializes + // independently to TRUE-space f16 via its own cpy_turbo{3,2}_tcq_f16 kernel (#533), then plain MMA. const bool both_tcq = (k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ) && (v->type == GGML_TYPE_TURBO3_TCQ || v->type == GGML_TYPE_TURBO2_TCQ); const bool fused_turbo = cparams.flash_attn && kq_b == nullptr && is_turbo_vec(k->type) && is_turbo_vec(v->type) && (k->type == v->type || both_tcq) && (k->ne[0] == 128 || k->ne[0] == 256 || fused_turbo_512) && - (tbv_nq_per_strm <= tbv_fuse_max_nq || is_tcq_kv); + tbv_nq_per_strm <= tbv_fuse_max_nq; // opencoti #533: TCQ de-fuses at prefill (now has a materialize-lift) if (fused_turbo) { // forward-rotate Q (×scale_inv → signs₁ → WHT → 1/√G·signs₂); group auto-picks 128 since @@ -2121,7 +2129,8 @@ ggml_tensor * llm_graph_context::build_attn_mha( q = ggml_turbo_wht(ctx0, q, /*direction=*/0, /*group_size=*/0, /*scale=*/nullptr); } else { const auto is_turbo = [](enum ggml_type t) { - return t == GGML_TYPE_TURBO2_0 || t == GGML_TYPE_TURBO3_0 || t == GGML_TYPE_TURBO4_0 || t == GGML_TYPE_TURBO8_0; + return t == GGML_TYPE_TURBO2_0 || t == GGML_TYPE_TURBO3_0 || t == GGML_TYPE_TURBO4_0 || t == GGML_TYPE_TURBO8_0 + || t == GGML_TYPE_TURBO3_TCQ || t == GGML_TYPE_TURBO2_TCQ; // opencoti #533: TCQ→f16 materialize-lift (prefill-MMA + DCA) }; if (is_turbo(k->type)) { k = ggml_cast(ctx0, k, GGML_TYPE_F16); } if (is_turbo(v->type)) { v = ggml_cast(ctx0, v, GGML_TYPE_F16); } @@ -2186,6 +2195,15 @@ ggml_tensor * llm_graph_context::build_attn_mha( ggml_flash_attn_ext_add_sinks(cur, sinks); ggml_flash_attn_ext_set_prec (cur, GGML_PREC_F32); + // opencoti-hook: sparse-attn #551 — attach the Quest block bitmask (parked by build_attn's + // store block) as the FA op's src[7]. launch_fattn reads src[7] and passes it to the FA-VEC + // kernel for per-block skipping. Consume-once (clear) so it never leaks to the next FA op. + // src[5]/src[6] are reserved for the streaming/position-window k_tail/v_tail. + if (opencoti_sparse_block_sel) { + cur->src[7] = opencoti_sparse_block_sel; + opencoti_sparse_block_sel = nullptr; + } + if (fused_turbo) { // opencoti-hook: turboquant perf-lift (M6-S2 Level B). Inverse-WHT + InnerQ un-equalize // on the SMALL FA output (signs₂ → WHT → 1/√G·signs₁ → ×scale_inv), recovering true-space @@ -2767,6 +2785,100 @@ ggml_tensor * llm_graph_context::build_attn( ggml_build_forward_expand(gf, mctx_cur->cpy_k(ctx0, k_cur, k_idxs, il)); ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, v_idxs, il)); + + // opencoti-hook: sparse-attn — #551 B1/B2. Reduce the just-written cache K into the per-block + // kmin/kmax side-cache (fill), then run the Quest selector (Q·bounds upper bound → top-K + + // recent/sink → I32 block bitmask). The selector reads the fill RESULT so the scheduler orders + // select-after-fill; the bitmask is parked in opencoti_sparse_block_sel and attached as the + // standard FA op's src[7] in build_attn_mha. Off ⇒ get_kbounds nullptr ⇒ no op ⇒ byte-identical. + // opencoti-hook: sparse-attn-vslash — #551. mode==0 (default) = Quest bound (this block, unchanged); + // mode==1 = vertical-slash content-driven selector (the separate block immediately below). + if (cparams.sparse_attn_enabled && cparams.sparse_attn_mode == 0) { + ggml_tensor * kbounds = mctx_cur->get_kbounds(ctx0, il); + if (kbounds) { + ggml_tensor * k_full = mctx_cur->get_k(ctx0, il); + // opencoti #551: block size must be a multiple of the FA-VEC tile width (nthreads=128) so the + // compacted-list consumer's tiles align to blocks (sparse_m = B/128 ≥ 1). Clamp up. + const int B_raw = (int) cparams.sparse_attn_block_size; + const int B = B_raw < 128 ? 128 : (B_raw / 128) * 128; + // opencoti #551 sub-block Quest bound: kbounds is filled at the finer raw granularity (subB, + // = kbounds alloc granularity), and the score coalesces SUB sub-blocks per walk-block B via a + // max. Tightens the kmin/kmax envelope (B=128 is ~8x Quest's page → loose). subB must divide + // the walk B and be ≤ B; otherwise fall back to subB=B (SUB=1 ⇒ identical to plain B-block). + const int subB = (B_raw > 0 && B_raw <= B && B % B_raw == 0) ? B_raw : B; + const int SUB = B / subB; + // opencoti bug-744: the fill reads k_cur (the NEW tokens, pre-quantization f16/f32) instead + // of the quantized cache view, so the Quest kmin/kmax bound composes with EVERY cache KV + // type (q8_0/q4_0/turbo*/tcq*). k_full is kept for the n_kv/n_block sizing below. + ggml_tensor * k_fill = ggml_is_contiguous(k_cur) ? k_cur : ggml_cont(ctx0, k_cur); + ggml_tensor * fill = ggml_sparse_attn_fill(ctx0, kbounds, k_fill, k_idxs, subB); + ggml_build_forward_expand(gf, fill); + // opencoti #551 PERF: the FA-VEC DECODE kernel is the only consumer of block_sel; the + // MMA prefill kernel ignores it. Building the O(n_block·n_q) score+select at prefill + // (n_q>1) is pure overhead for zero gain (measured 5-8× prefill collapse). Run the + // selector ONLY at decode (n_q==1); fill still runs every step so kbounds stays current. + if (q_cur->ne[2] == 1) { + // opencoti #551 bug-743 FIX: get_k() returns [head_dim, n_head_kv, n_kv, 1] — the KV + // length is ne[2], NOT ne[1] (=n_head_kv). Reading ne[1] gave n_kv≈8 ⇒ n_block=1 ⇒ the + // selector saw one block ⇒ FA-VEC walked only the first ~128 keys of the cache (garbage + // past 128 tokens; invisible at short ctx where 1 block covers all). Use the position axis. + const int n_kv = (int) k_full->ne[2]; + const int n_block = (n_kv + B - 1) / B; + // topk==0 ⇒ select-all (inert B2 byte-identical anchor). --sparse-attn-recent is in + // TOKENS → round up to blocks; sink=1 forces block 0. select's src[1]=fill orders + // select-after-fill. topk==0xFFFFFFFF ⇒ auto: ¼ of blocks, floor 64. + const int recent_blk = ((int) cparams.sparse_attn_recent + B - 1) / B; + const int sink_blk = (int) cparams.sparse_attn_sink; + const int k_sel = (cparams.sparse_attn_topk == 0xFFFFFFFFu) + ? (n_block / 4 > 64 ? n_block / 4 : 64) + : (int) cparams.sparse_attn_topk; + // #551 PERF (amortized selection): when the KV cache exposes a persistent per-stream + // block_sel buffer (single-stream), the SELECT op writes into it in-place and re-selects + // only every cparams.sparse_attn_refresh decode tokens (in-kernel gate on the position + // from k_idxs), reusing the cached list in between. nullptr ⇒ transient per-step select. + ggml_tensor * block_sel_buf = mctx_cur->get_block_sel(ctx0, il); + opencoti_sparse_block_sel = ggml_sparse_attn_select( + ctx0, q_cur, fill, block_sel_buf, block_sel_buf ? k_idxs : nullptr, B, n_block, + k_sel, recent_blk, sink_blk, (int) cparams.sparse_attn_refresh, SUB); + ggml_build_forward_expand(gf, opencoti_sparse_block_sel); + } + } + } + // opencoti-hook: sparse-attn-vslash — #551 content-driven selector (mode==1), MInference prefill-window. + // The selector is estimated ONCE at PREFILL (n_q>1) over the last q_win prompt rows — which, for a + // retrieval prompt, ARE the question and genuinely attend the needle — by softmax block-mass (vertical + // importance), written into the persistent per-layer block_sel buffer. At DECODE (n_q==1) we do NOT + // recompute (the prompt queries are gone and a single decode query is unreliable at long ctx — the + // q_win=1 decode estimate failed the 40k gate AND cost a full attention pass per step, bug-741); we + // simply re-attach the cached block_sel as FA src[7]. The MMA prefill path is dense and ignores src[7], + // so the prefill emit only side-writes the buffer (opencoti_sparse_block_sel stays null at prefill). + if (cparams.sparse_attn_enabled && cparams.sparse_attn_mode == 1) { + ggml_tensor * k_full = mctx_cur->get_k(ctx0, il); + ggml_tensor * block_sel_buf = k_full ? mctx_cur->get_block_sel(ctx0, il) : nullptr; + if (block_sel_buf) { + const int B_raw = (int) cparams.sparse_attn_block_size; + const int B = B_raw < 128 ? 128 : (B_raw / 128) * 128; + const int n_kv = (int) k_full->ne[2]; // position axis (bug-743) + const int n_block = (n_kv + B - 1) / B; + const int recent_blk = ((int) cparams.sparse_attn_recent + B - 1) / B; + const int sink_blk = (int) cparams.sparse_attn_sink; + const int k_sel = (cparams.sparse_attn_topk == 0xFFFFFFFFu) + ? (n_block / 4 > 64 ? n_block / 4 : 64) + : (int) cparams.sparse_attn_topk; + if (q_cur->ne[2] > 1) { + // PREFILL: estimate over the last q_win rows of this chunk; refresh=1 (always compute, never + // amortize at prefill). The final prompt chunk's write persists into decode. Side-write only. + const int q_win = q_cur->ne[2] < 8 ? (int) q_cur->ne[2] : 8; // #551 fast correctness read (was 64; prefill cost ∝ q_win) + ggml_tensor * sel = ggml_sparse_attn_vslash( + ctx0, q_cur, k_full, block_sel_buf, k_idxs, B, n_block, + k_sel, recent_blk, sink_blk, /*refresh=*/1, q_win); + ggml_build_forward_expand(gf, sel); + } else { + // DECODE: reuse the cached prefill selection — attach as FA src[7], no recompute. + opencoti_sparse_block_sel = block_sel_buf; + } + } + } } const auto & kq_mask = inp->get_kq_mask(); @@ -3070,6 +3182,82 @@ ggml_tensor * llm_graph_context::build_attn( ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, v_idxs, il)); } + // opencoti-hook: sparse-attn #551 — Quest block selector on the iSWA path, GLOBAL (!is_swa) + // layers ONLY. SWA layers are already windowed (block-selection is moot); the global + // full-attention layers are exactly the long-ctx KV-bandwidth hogs sparse targets. Mirrors the + // unified overload's store block (build_attn @ ~2794). Clear unconditionally first so a global + // layer's bitmask never bleeds into a following SWA layer's build_attn_mha consume (the + // window_layer branch below does not consume it). Engages only on build_attn_mha (resident KV); + // the POSITION_WINDOW overflow tactic falls through without sparse (first-cut limitation). + opencoti_sparse_block_sel = nullptr; + // opencoti-hook: sparse-attn-vslash — #551. mode==0 (default) = Quest bound (this block); mode==1 = + // vertical-slash content-driven selector (the separate block below). iSWA global (resident-KV) path. + if (!is_swa && cparams.sparse_attn_enabled && cparams.sparse_attn_mode == 0 && k_cur) { + ggml_tensor * kbounds = mctx_cur->get_kbounds(ctx0, il); + if (kbounds) { + ggml_tensor * k_full = mctx_cur->get_k(ctx0, il); + // opencoti #551: block size must be a multiple of the FA-VEC tile width (nthreads=128). + const int B_raw = (int) cparams.sparse_attn_block_size; + const int B = B_raw < 128 ? 128 : (B_raw / 128) * 128; + // opencoti #551 sub-block Quest bound (iSWA global path) — see unified path for rationale. + const int subB = (B_raw > 0 && B_raw <= B && B % B_raw == 0) ? B_raw : B; + const int SUB = B / subB; + // opencoti bug-744: fill from k_cur (new tokens, pre-quant) so sparse composes with all KV types + // (iSWA global path). k_full kept for n_kv/n_block sizing below. + ggml_tensor * k_fill = ggml_is_contiguous(k_cur) ? k_cur : ggml_cont(ctx0, k_cur); + ggml_tensor * fill = ggml_sparse_attn_fill(ctx0, kbounds, k_fill, inp->get_k_idxs(), subB); + ggml_build_forward_expand(gf, fill); + // opencoti #551 PERF: selector (score+select) only at decode (n_q==1) — MMA prefill ignores + // block_sel, so building it at prefill is pure overhead. Fill runs every step (kbounds current). + if (q_cur->ne[2] == 1) { + // opencoti #551 bug-743 FIX: KV length is k_full->ne[2] (position axis), NOT ne[1] + // (=n_head_kv). ne[1] gave n_block=1 ⇒ only the first ~128 keys attended. (iSWA global path.) + const int n_kv = (int) k_full->ne[2]; + const int n_block = (n_kv + B - 1) / B; + const int recent_blk = ((int) cparams.sparse_attn_recent + B - 1) / B; + const int sink_blk = (int) cparams.sparse_attn_sink; + const int k_sel = (cparams.sparse_attn_topk == 0xFFFFFFFFu) + ? (n_block / 4 > 64 ? n_block / 4 : 64) + : (int) cparams.sparse_attn_topk; + // #551 PERF (amortized): persistent per-stream block_sel buffer (single-stream) ⇒ in-place + // write + re-select every cparams.sparse_attn_refresh decode tokens; nullptr ⇒ transient. + ggml_tensor * block_sel_buf = mctx_cur->get_block_sel(ctx0, il); + opencoti_sparse_block_sel = ggml_sparse_attn_select( + ctx0, q_cur, fill, block_sel_buf, block_sel_buf ? inp->get_k_idxs() : nullptr, + B, n_block, k_sel, recent_blk, sink_blk, (int) cparams.sparse_attn_refresh, SUB); + ggml_build_forward_expand(gf, opencoti_sparse_block_sel); + } + } + } + // opencoti-hook: sparse-attn-vslash — #551 content-driven selector (mode==1), iSWA global path. + // MInference prefill-window: estimate ONCE at prefill (n_q>1) over the last q_win prompt rows by softmax + // block-mass into the persistent block_sel buffer; at decode reuse the cached selection (no recompute). + // Mirrors the unified-site placement above (see rationale + bug-741 there). + if (!is_swa && cparams.sparse_attn_enabled && cparams.sparse_attn_mode == 1 && k_cur) { + ggml_tensor * k_full = mctx_cur->get_k(ctx0, il); + ggml_tensor * block_sel_buf = k_full ? mctx_cur->get_block_sel(ctx0, il) : nullptr; + if (block_sel_buf) { + const int B_raw = (int) cparams.sparse_attn_block_size; + const int B = B_raw < 128 ? 128 : (B_raw / 128) * 128; + const int n_kv = (int) k_full->ne[2]; // position axis (bug-743) + const int n_block = (n_kv + B - 1) / B; + const int recent_blk = ((int) cparams.sparse_attn_recent + B - 1) / B; + const int sink_blk = (int) cparams.sparse_attn_sink; + const int k_sel = (cparams.sparse_attn_topk == 0xFFFFFFFFu) + ? (n_block / 4 > 64 ? n_block / 4 : 64) + : (int) cparams.sparse_attn_topk; + if (q_cur->ne[2] > 1) { + const int q_win = q_cur->ne[2] < 8 ? (int) q_cur->ne[2] : 8; // #551 fast correctness read (was 64; prefill cost ∝ q_win) + ggml_tensor * sel = ggml_sparse_attn_vslash( + ctx0, q_cur, k_full, block_sel_buf, inp->get_k_idxs(), B, n_block, + k_sel, recent_blk, sink_blk, /*refresh=*/1, q_win); + ggml_build_forward_expand(gf, sel); + } else { + opencoti_sparse_block_sel = block_sel_buf; + } + } + } + const auto & kq_mask = is_swa ? inp->get_kq_mask_swa() : inp->get_kq_mask(); ggml_tensor * q = q_cur; diff --git a/llama.cpp/src/llama-kv-cache-iswa.cpp b/llama.cpp/src/llama-kv-cache-iswa.cpp index 79fefd0..ba644cc 100644 --- a/llama.cpp/src/llama-kv-cache-iswa.cpp +++ b/llama.cpp/src/llama-kv-cache-iswa.cpp @@ -30,7 +30,8 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( uint32_t n_ubatch, uint32_t n_pad, const layer_filter_cb & filter, - const layer_reuse_cb & reuse) : hparams(model.hparams), unified(unified) { + const layer_reuse_cb & reuse, + uint32_t sparse_attn_block_size) : hparams(model.hparams), unified(unified) { // chain filters const layer_filter_cb filter_base = [&](int32_t il) { @@ -69,7 +70,7 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( model, type_k, type_v, v_trans, offload, unified, size_base, kv_size_initial, slot_shrink_idle_ms, headinfer_gpu_heads_frac, vram_target_mib, pcie_bw_gbps, kv_residency_mode, - n_seq_max, n_pad, 0, LLAMA_SWA_TYPE_NONE, filter_base, reuse); + n_seq_max, n_pad, 0, LLAMA_SWA_TYPE_NONE, filter_base, reuse, sparse_attn_block_size); LLAMA_LOG_INFO("%s: creating SWA KV cache, size = %u cells\n", __func__, size_swa); @@ -77,7 +78,7 @@ llama_kv_cache_iswa::llama_kv_cache_iswa( model, type_k, type_v, v_trans, offload, unified, size_swa, kv_size_initial, slot_shrink_idle_ms, headinfer_gpu_heads_frac, vram_target_mib, pcie_bw_gbps, kv_residency_mode, - n_seq_max, n_pad, hparams.n_swa, hparams.swa_type, filter_swa, reuse); + n_seq_max, n_pad, hparams.n_swa, hparams.swa_type, filter_swa, reuse, sparse_attn_block_size); } void llama_kv_cache_iswa::clear(bool data) { diff --git a/llama.cpp/src/llama-kv-cache-iswa.h b/llama.cpp/src/llama-kv-cache-iswa.h index 94feb62..d274b0a 100644 --- a/llama.cpp/src/llama-kv-cache-iswa.h +++ b/llama.cpp/src/llama-kv-cache-iswa.h @@ -45,7 +45,11 @@ public: uint32_t n_ubatch, uint32_t n_pad, const layer_filter_cb & filter, - const layer_reuse_cb & reuse); + const layer_reuse_cb & reuse, + // opencoti #551 sparse-attn — forwarded to both kv_base and + // kv_swa. Quest block-selector B_SEL; 0 = off (no side-cache). + // See docs/features/sparse_attn.md. + uint32_t sparse_attn_block_size = 0); ~llama_kv_cache_iswa() = default; diff --git a/llama.cpp/src/llama-kv-cache.cpp b/llama.cpp/src/llama-kv-cache.cpp index 4bab0d1..3687725 100644 --- a/llama.cpp/src/llama-kv-cache.cpp +++ b/llama.cpp/src/llama-kv-cache.cpp @@ -332,12 +332,17 @@ llama_kv_cache::llama_kv_cache( uint32_t n_swa, llama_swa_type swa_type, const layer_filter_cb & filter, - const layer_reuse_cb & reuse) : + const layer_reuse_cb & reuse, + uint32_t sparse_attn_block_size) : model(model), hparams(model.hparams), v_trans(v_trans), n_seq_max(n_seq_max), n_stream(unified ? 1 : n_seq_max), n_pad(n_pad), n_swa(n_swa), swa_type(swa_type) { GGML_ASSERT(kv_size % n_pad == 0); + // opencoti #551 sparse-attn — stash B_SEL (0 = off) for the side-cache alloc + // below + the selector (S2). See docs/features/sparse_attn.md. + sparse_attn_block_size_ = sparse_attn_block_size; + // opencoti F4 M3 Phase 2 — see docs/decisions/0001-lazy-slot-context.md // Initialize the hard cap and the soft cap. kv_size_initial==0 (the // legacy / Phase-1 default) means "no soft cap" so the whole kv_size @@ -424,6 +429,37 @@ llama_kv_cache::llama_kv_cache( window_cells, kv_size, kv_size - window_cells); } + // opencoti F5 rolling-KV S0 — per-layer residency budget (task #585). + // OPENCOTI_KV_LOCAL_LAYERS="0,1,2,4,11" marks provably-local layers (from the HAL + // attention-locality profile — layers whose retrieval mass stays inside the recent + // window) as FULLY RESIDENT: window_cells==kv_size ⇒ no host tail for that layer. + // Full-attention models (Qwen) then stream a tail on only the RETRIEVING layers, + // the way Gemma iSWA already does for free. Manual env gate for the live niah gate; + // S4 auto-derives the set from an online prefill locality estimate. Empty/unset ⇒ + // uniform behavior (byte-inert). Only meaningful in window_mode. GATE per-layer niah, + // never widen past retrieval survival. See docs/features/rolling_kv_step_prefetch.md. + std::vector kv_local_layer(hparams.n_layer, false); + if (window_mode) { + const char * ll = getenv("OPENCOTI_KV_LOCAL_LAYERS"); + if (ll && *ll) { + std::string s(ll); + size_t pos = 0; + while (pos < s.size()) { + const size_t comma = s.find(',', pos); + const std::string tok = s.substr(pos, comma == std::string::npos ? std::string::npos : comma - pos); + if (!tok.empty()) { + const long v = strtol(tok.c_str(), nullptr, 10); + if (v >= 0 && (uint32_t) v < hparams.n_layer) kv_local_layer[v] = true; + } + if (comma == std::string::npos) break; + pos = comma + 1; + } + uint32_t nloc = 0; for (const bool b : kv_local_layer) nloc += b ? 1u : 0u; + LLAMA_LOG_INFO("%s: rolling-kv per-layer residency — %u/%u layer(s) forced RESIDENT " + "(no host tail) via OPENCOTI_KV_LOCAL_LAYERS\n", __func__, nloc, hparams.n_layer); + } + } + // define a comparator for the buft -> ctx map to ensure that the order is well-defined: struct ggml_backend_buft_comparator { bool operator()(const ggml_backend_buffer_type_t & lhs, const ggml_backend_buffer_type_t & rhs) const { @@ -457,8 +493,14 @@ llama_kv_cache::llama_kv_cache( auto & ctx_map = ctx_map_per_stream[s]; auto it = ctx_map.find(buft); if (it == ctx_map.end()) { + // opencoti #551 sparse-attn — the per-(buft,stream) metadata pool budgets + // 4 tensors/layer (k, v, k_cpu, v_cpu worst case). When sparse-attn is on, + // each device layer also allocates kmin + kmax into this ctx, so bump the + // budget to 8/layer (covers k,v,kmin,kmax + the cpu-coresident worst case + // with margin). no_alloc ⇒ metadata only, so the KV data buffer is + // unaffected; off (block_size==0) keeps the original 4u (byte-identical). ggml_init_params params = { - /*.mem_size =*/ size_t(4u*n_layer_kv*ggml_tensor_overhead()), + /*.mem_size =*/ size_t((sparse_attn_block_size > 0 ? 8u : 4u)*n_layer_kv*ggml_tensor_overhead()), /*.mem_buffer =*/ NULL, /*.no_alloc =*/ true, }; @@ -607,8 +649,16 @@ llama_kv_cache::llama_kv_cache( // host tail tensor spans [wc, kv_size) (tail_c cells); both carry ALL // heads. Outside window mode wc == kv_size and tail_c == 0, so the // tensor shapes below are byte-identical to the pre-3c layout. - const uint32_t wc = window_mode ? window_cells : kv_size; - const uint32_t tail_c = window_mode ? (kv_size - window_cells) : 0; + // opencoti F5 rolling-KV S0 (#585): a layer marked local (OPENCOTI_KV_LOCAL_LAYERS) + // is FULLY RESIDENT — wc==kv_size, no host tail — so only retrieving layers stream. + const bool layer_local = window_mode && kv_local_layer[il]; + // A local layer behaves EXACTLY like a non-window layer: fully resident + // (wc == kv_size), NO host tail, and window_cells sentinel 0 → the accessors + // and graph take the plain-FA resident path (never touch a 0-row tail). Gate + // ALL per-layer window/tail machinery on layer_window, not the global window_mode. + const bool layer_window = window_mode && !layer_local; + const uint32_t wc = layer_window ? window_cells : kv_size; + const uint32_t tail_c = layer_window ? (kv_size - window_cells) : 0; // opencoti F4 M3 Phase 4 — per-stream tensor split. // Each stream gets its own OWNING 3-D tensor of shape @@ -624,6 +674,12 @@ llama_kv_cache::llama_kv_cache( std::vector v_cpu_per_stream; std::vector k_stream; std::vector v_stream; + // opencoti #551 sparse-attn — per-stream kbounds side-cache. Stay EMPTY + // when sparse-attn off (sparse_attn_block_size == 0) ⇒ byte-identical. + std::vector kbounds_per_stream; + // opencoti #551 sparse-attn PERF — per-stream persistent selected-block list (amortized + // selection). EMPTY when sparse off OR n_stream > 1 (per-step transient fallback). + std::vector block_sel_per_stream; for (uint32_t s = 0; s < n_stream; ++s) { ggml_context * ctx_s = ctx_for_buft(buft, s); @@ -673,15 +729,15 @@ llama_kv_cache::llama_kv_cache( // mode forces gpu_heads == 0), so exactly one branch sizes the // tensor. The host buffer is the pinned cpu_buft when CUDA is the // layer dev — full PCIe DMA for the stream-back. - if (gpu_heads > 0 || window_mode) { + if (gpu_heads > 0 || layer_window) { ggml_context * ctx_cpu_s = ctx_for_buft(cpu_buft, s); if (!ctx_cpu_s) { throw std::runtime_error("failed to create CPU ggml context for headinfer kv cache"); } // window mode: ALL heads, tail_c cells; M2: CPU head subset, kv_size cells. - const uint32_t host_k_embd = window_mode ? n_embd_k_gqa : n_embd_k_cpu; - const uint32_t host_v_embd = window_mode ? n_embd_v_gqa : n_embd_v_cpu; - const uint32_t host_cells = window_mode ? tail_c : kv_size; + const uint32_t host_k_embd = layer_window ? n_embd_k_gqa : n_embd_k_cpu; + const uint32_t host_v_embd = layer_window ? n_embd_v_gqa : n_embd_v_cpu; + const uint32_t host_cells = layer_window ? tail_c : kv_size; ggml_tensor * k_cpu_s = has_k ? ggml_new_tensor_3d(ctx_cpu_s, type_k, host_k_embd, host_cells, 1) : nullptr; @@ -690,8 +746,8 @@ llama_kv_cache::llama_kv_cache( : nullptr; // Name distinguishes the two layouts for state-IO grep: // "cache_k_tail" (position tail) vs "cache_k_cpu" (head half). - const char * k_base = window_mode ? "cache_k_tail" : "cache_k_cpu"; - const char * v_base = window_mode ? "cache_v_tail" : "cache_v_cpu"; + const char * k_base = layer_window ? "cache_k_tail" : "cache_k_cpu"; + const char * v_base = layer_window ? "cache_v_tail" : "cache_v_cpu"; if (k_cpu_s) { if (n_stream == 1) { ggml_format_name(k_cpu_s, "%s_l%d", k_base, il); @@ -731,6 +787,33 @@ llama_kv_cache::llama_kv_cache( v_stream.push_back(v_s ? ggml_view_2d(ctx_s, v_s, n_embd_v_gpu, wc, v_s->nb[1], 0) : nullptr); + + // opencoti #551 sparse-attn — per-(KV-head, block) min/max bounds, + // allocated ONLY when enabled. n_block_sel = ceil(wc / B_SEL). Lives in + // the device ctx_s alongside k_s (same n_embd_k_gpu dim0 as the resident + // window), f16 regardless of KV-quant (bounds come from dequant K at the + // S2 write). Off ⇒ vectors stay empty ⇒ no tensor, byte-identical. + if (sparse_attn_block_size > 0 && k_s) { + const uint32_t n_block_sel = (wc + sparse_attn_block_size - 1) / sparse_attn_block_size; + // [n_embd_k_gpu, n_block_sel, 2] f16 : [..,0]=min, [..,1]=max + ggml_tensor * kbounds_s = ggml_new_tensor_3d(ctx_s, GGML_TYPE_F16, n_embd_k_gpu, n_block_sel, 2); + if (n_stream == 1) { + ggml_format_name(kbounds_s, "cache_kbounds_l%d", il); + } else { + ggml_format_name(kbounds_s, "cache_kbounds_l%d_s%u", il, s); + } + kbounds_per_stream.push_back(kbounds_s); + + // opencoti #551 sparse-attn PERF (amortized selection) — persistent selected-block + // list, single-stream only (n_stream > 1 falls back to per-step transient selection). + // [n_block_sel + 1, 1] I32: [0,n_block_sel) = sorted −1-padded compacted list, + // [n_block_sel] = reserved scratch. Init to [0,−1,…] post-alloc (see below). + if (n_stream == 1) { + ggml_tensor * block_sel_s = ggml_new_tensor_2d(ctx_s, GGML_TYPE_I32, n_block_sel + 1, 1); + ggml_format_name(block_sel_s, "cache_block_sel_l%d", il); + block_sel_per_stream.push_back(block_sel_s); + } + } } map_layer_ids[il] = layers.size(); @@ -741,7 +824,9 @@ llama_kv_cache::llama_kv_cache( k_stream, v_stream, k_cpu_per_stream, v_cpu_per_stream, gpu_heads, - window_mode ? window_cells : 0u, + layer_window ? window_cells : 0u, + kbounds_per_stream, + block_sel_per_stream, }); } @@ -826,6 +911,23 @@ llama_kv_cache::llama_kv_cache( } } + // opencoti #551 sparse-attn PERF — initialize each persistent block_sel list to [0, −1, −1, …] + // (sink block 0 selected, rest empty). ensure_cleared() only zeroes the K/V tensors, never + // block_sel, so without this the buffer is uninitialized garbage on the first decode steps before + // the selector's first refresh. [0,−1,…] makes a never-refreshed buffer attend ONLY the sink + // block — safe (the FA-VEC consumer breaks at the first −1; never empty ⇒ no NaN, never OOB). + if (!model.hparams.no_alloc) { + for (const auto & layer : layers) { + for (ggml_tensor * bs : layer.block_sel_per_stream) { + if (!bs) { continue; } + const int64_t n = bs->ne[0]; + std::vector init((size_t) n, -1); + init[0] = 0; + ggml_backend_tensor_set(bs, init.data(), 0, (size_t) n * sizeof(int32_t)); + } + } + } + if (model.hparams.no_alloc) { // opencoti F4 M3 Phase 1+2 — dummy buffers are size 0; treat all // cells as "already cleared" so ensure_cleared is a no-op. @@ -2831,6 +2933,34 @@ ggml_tensor * llama_kv_cache::get_k_gpu(ggml_context * ctx, int32_t il, uint32_t return result; } +// opencoti-hook: sparse-attn — #551. Returns the per-(layer,stream) kbounds side-cache +// tensor [n_embd_k_gpu, n_block_sel, 2] f16, or nullptr when sparse-attn is off (the +// side-cache was never allocated). Referenced raw as a graph src (like k_l); the fill op +// views it in-place. n_kv is irrelevant — kbounds spans all selectable blocks. +ggml_tensor * llama_kv_cache::get_kbounds(ggml_context * ctx, int32_t il, const slot_info & sinfo) const { + GGML_UNUSED(ctx); + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + if (layer.kbounds_per_stream.empty()) { + return nullptr; + } + return layer.kbounds_per_stream[sinfo.s0]; +} + +// opencoti-hook: sparse-attn — #551 PERF. Returns the per-(layer,stream) persistent selected-block +// list tensor [n_block_sel + 1, 1] I32, or nullptr when sparse-attn is off OR multi-stream (the +// vector is only populated for n_stream == 1). The SELECT op views it in-place (cpy idiom) and +// re-writes it only every cparams.sparse_attn_refresh decode tokens; FA-VEC reads it every step. +ggml_tensor * llama_kv_cache::get_block_sel(ggml_context * ctx, int32_t il, const slot_info & sinfo) const { + GGML_UNUSED(ctx); + const int32_t ikv = map_layer_ids.at(il); + const auto & layer = layers[ikv]; + if (layer.block_sel_per_stream.empty()) { + return nullptr; + } + return layer.block_sel_per_stream[sinfo.s0]; +} + ggml_tensor * llama_kv_cache::get_k_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const { const int32_t ikv = map_layer_ids.at(il); const auto & layer = layers[ikv]; @@ -5146,6 +5276,16 @@ ggml_tensor * llama_kv_cache_context::get_v(ggml_context * ctx, int32_t il) cons return kv->get_v(ctx, il, n_kv, sinfos[i_cur]); } +// opencoti-hook: sparse-attn — #551. Per-batch wrapper around the kbounds accessor. +ggml_tensor * llama_kv_cache_context::get_kbounds(ggml_context * ctx, int32_t il) const { + return kv->get_kbounds(ctx, il, sinfos[i_cur]); +} + +// opencoti-hook: sparse-attn — #551 PERF. Per-batch wrapper around the persistent block_sel accessor. +ggml_tensor * llama_kv_cache_context::get_block_sel(ggml_context * ctx, int32_t il) const { + return kv->get_block_sel(ctx, il, sinfos[i_cur]); +} + // opencoti F5 M3 neo — per-batch wrappers around the head-split accessors. ggml_tensor * llama_kv_cache_context::get_k_gpu(ggml_context * ctx, int32_t il) const { return kv->get_k_gpu(ctx, il, n_kv, sinfos[i_cur]); diff --git a/llama.cpp/src/llama-kv-cache.h b/llama.cpp/src/llama-kv-cache.h index 5e8ea9b..c6027f1 100644 --- a/llama.cpp/src/llama-kv-cache.h +++ b/llama.cpp/src/llama-kv-cache.h @@ -136,7 +136,11 @@ public: uint32_t n_swa, llama_swa_type swa_type, const layer_filter_cb & filter, - const layer_reuse_cb & reuse); + const layer_reuse_cb & reuse, + // opencoti #551 sparse-attn — Quest block-selector KV-block + // size B_SEL (docs/features/sparse_attn.md). 0 = sparse-attn + // OFF: no kmin/kmax side-cache allocated (byte-identical). + uint32_t sparse_attn_block_size = 0); ~llama_kv_cache() = default; @@ -251,6 +255,11 @@ public: ggml_tensor * get_v_gpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; ggml_tensor * get_k_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; ggml_tensor * get_v_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const; + // opencoti-hook: sparse-attn — #551. Per-(layer,stream) kbounds side-cache, or nullptr when off. + ggml_tensor * get_kbounds(ggml_context * ctx, int32_t il, const slot_info & sinfo) const; + // opencoti-hook: sparse-attn — #551 PERF. Per-(layer,stream) persistent selected-block list, or + // nullptr when off / multi-stream (then the selector emits a transient per-step list instead). + ggml_tensor * get_block_sel(ggml_context * ctx, int32_t il, const slot_info & sinfo) const; bool headinfer_split_active(int32_t il) const; uint32_t headinfer_gpu_heads (int32_t il) const; @@ -423,6 +432,27 @@ private: // [window_cells, kv_size). 0 ⇒ no position window (vanilla / M2 layout). uint32_t window_cells = 0; + // opencoti #551 sparse-attn — per-(KV-head·channel, block) min/max key + // bounds side-cache for the Quest block-selector. EMPTY when sparse-attn + // is off (sparse_attn_block_size_ == 0) ⇒ no tensor allocated, + // byte-identical. Shape per stream: [n_embd_k_gpu, n_block_sel, 2] f16 + // ([..,0]=min, [..,1]=max), with n_block_sel = ceil(window_cells / B_SEL). + // Filled by GGML_OP_SPARSE_ATTN_FILL right after the cpy_k write (B1); + // read by the selector (B2) + FA-VEC block-skip (B2/S3). + // See docs/features/sparse_attn.md. + std::vector kbounds_per_stream; + + // opencoti #551 sparse-attn PERF (amortized selection) — per-stream persistent + // I32 selected-block list [n_block_sel + 1, 1]: slots [0,n_block_sel) hold the + // SORTED −1-padded compacted list (FA-VEC consumer), slot [n_block_sel] is a + // scratch pad reserved for future markers. Persists across decode steps so the + // GGML_OP_SPARSE_ATTN_SELECT kernel can re-run only every cparams.sparse_attn_refresh + // tokens (reading the decode position from k_idxs) and reuse the cached selection in + // between — amortizing the per-step selector cost. EMPTY when sparse off OR n_stream>1 + // (multi-stream falls back to per-step transient selection). Init to [0,−1,−1,…] so a + // never-refreshed buffer attends only the sink block (safe, never OOB / never empty). + std::vector block_sel_per_stream; + // Convenience accessors for unified-mode (n_stream == 1) call // sites. Returns nullptr in non-unified mode — that's the // signal that the caller must use per-stream access instead. @@ -471,6 +501,10 @@ private: // this is the SWA type of the cache - not to be confused with the model SWA type const llama_swa_type swa_type = LLAMA_SWA_TYPE_NONE; + // opencoti #551 sparse-attn — Quest block-selector KV-block size B_SEL. + // 0 = sparse-attn off (no kmin/kmax side-cache). Set from the ctor arg. + uint32_t sparse_attn_block_size_ = 0; + // ggml contexts for the KV cache along with the allocated backend buffers: std::vector> ctxs_bufs; @@ -622,6 +656,11 @@ public: ggml_tensor * get_k(ggml_context * ctx, int32_t il) const; ggml_tensor * get_v(ggml_context * ctx, int32_t il) const; + // opencoti-hook: sparse-attn — #551. Per-(layer,stream) kbounds side-cache, or nullptr when off. + ggml_tensor * get_kbounds(ggml_context * ctx, int32_t il) const; + // opencoti-hook: sparse-attn — #551 PERF. Per-batch wrapper around the persistent block_sel accessor. + ggml_tensor * get_block_sel(ggml_context * ctx, int32_t il) const; + // opencoti F5 M3 neo — head-split-aware K/V accessors. Per-batch wrappers // that delegate to the underlying llama_kv_cache with the current n_kv / // slot info. See llama-kv-cache.h for semantics. diff --git a/llama.cpp/src/llama-model.cpp b/llama.cpp/src/llama-model.cpp index b18a498..017b755 100644 --- a/llama.cpp/src/llama-model.cpp +++ b/llama.cpp/src/llama-model.cpp @@ -2104,7 +2104,9 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, cparams.n_ubatch, 1, filter, - reuse); + reuse, + // opencoti #551 sparse-attn — B_SEL when enabled, else 0 (off) + cparams.sparse_attn_enabled ? cparams.sparse_attn_block_size : 0); } else { GGML_ASSERT(!hparams.is_swa_any()); @@ -2133,7 +2135,9 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, hparams.n_swa, hparams.swa_type, filter, - nullptr); + nullptr, + // opencoti #551 sparse-attn — B_SEL when enabled, else 0 (off) + cparams.sparse_attn_enabled ? cparams.sparse_attn_block_size : 0); } } }