Another DSA attempt by ikawrakow · Pull Request #2063 · ikawrakow/ik_llama.cpp

ikawrakow · 2026-07-01T13:43:44Z

I haven't abandoned your PR #2045, but have been trying various things starting from your branch.

This PR completely resolves the excessive memory usage problem, but something is not quite right. The model starts getting slightly confused after 3k tokens or so, and completely loses coherence after 4k tokens or so.

In summary, the changes added compared to #2045

Loop over attention heads
Removed the llama.cpp-style Hadamard tranformation via a matrix multiplication and replaced with the builtin ggml_hadamard (one of my comments in your PR)
Added a new ggml_blend op, which I thought makes it easier to construct the sparse attention mask
Added sorting of more than 1024 elements on CUDA so the sorting does not need to be send back to the CPU. That part was simply copied from mainline.
When FA is ON, the sparse mask is constructed differently using newly ggml_blend
Rebased on the main branch and resolved the merge conflict that exist in deepseek2 : GLM-DSA sparse attention (lightning indexer), --dsa off by default #2045

It runs quite a bit faster compared to #2045. I don't have a good hypothesis of what could be wrong with the implementation, so I thought I publish the way it is hoping that you or someone else will find the bug(s).

…seq prefill) Implements the sparse top-k "lightning indexer" attention for LLM_ARCH_GLM_DSA in build_deepseek2_layer_attention (ik's deepseek2 graph). What it does (per layer, gated on model.arch==GLM_DSA && indexer_attn_q_b): - indexer_q = indexer_attn_q_b(q_lora latent), split rope(64)/nope(64), NEOX-rope the pe part, concat. indexer_k = indexer_attn_k(attn_norm out), LayerNorm w/ bias, same rope/concat (single key head, MQA). - scores = relu(indexer_k . indexer_q), scaled per-head weights (indexer_proj), summed over heads, + base causal mask, then ggml_top_k(min(top_k, n_tokens)). - sparse mask: ggml_fill(-inf) -> ggml_set_rows(0) at top_k positions -> + causal, used in the soft_max_ext attention path (-mla 1 -fa 0) instead of KQ_mask. Simplifications (intentional, proven sound): - Batch-local: no indexer KV-cache. Indexer keys are the current batch tokens. - Walsh-Hadamard transform omitted: orthonormal rotation, (Hq).(Hk)==q.k, no score change. Validation (GLM-5.2-UD-IQ2_M, 3x P100, -mla 1 -fa 0): - Compiles clean (CUDA sm_60); loads and runs. - c512 -b512 (n_seq=1) PPL = 2.7760, byte-identical to dense baseline (indexer disabled) = 2.7760, all 8 chunks match -> indexer is an exact no-op when top_k>=n_tokens. Proves correctness-preservation. - 3105-token prompt completion (top_k=2048 < 3105 -> indexer ACTIVELY masks): prompt-eval produces coherent, accurate continuation, identical to dense for the prompt+early-gen tokens. No NaN/crash. Confirms the masking path works in prefill. Known limitations (documented follow-ups, NOT handled): - Single-sequence prefill only. Multi-sequence batches (n_seq>1, e.g. perplexity default n_batch>n_ctx) and kv_head>0 (decode) break the batch-local key->slot mapping. n_seq>1 -> NaN (use n_batch==n_ctx). Decode (kv_head>0): each generated token sees only itself as an indexer key, so generation degenerates into repetition after the prompt (dense A/B stays coherent) -- this is the decode-cache stub, the documented next step. - Flash-attn path (-fa 1, F16 mask) still uses dense KQ_mask (soft_max path only). - Decode indexer KV-cache + Hadamard cached-K storage not implemented. Runtime gate: DSA_INDEXER_DISABLE=1 falls back to dense attention (for A/B). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Make the lightning-indexer correct for DECODE (not just prefill). Previously the indexer was batch-local, so a generated token only scored against itself and generation degenerated. Now the indexer keys are cached across the full context. Changes - llama_kv_cache: add per-layer indexer-key cache `kr_l` [indexer_head_size, kv_size] (F16, MQA single head), allocated alongside the MLA latent cache for GLM_DSA. - build_deepseek2_dsa_indexer: write the batch's (Hadamard-rotated) indexer keys to kr_l at kv_head, read back the full [128, n_kv] cached keys, and score the indexer queries against ALL past keys. Returns the full descending argsort of the scores. - Walsh-Hadamard rotation of indexer q/k (cparams.dsa_indexer_hadamard, default on; filled in llama_set_inputs). Score-preserving; improves cached-K F16 precision. - build_deepseek2_dsa_sparse_mask: rank-based full-coverage scatter (write a 0/-BIG penalty into EVERY key slot keyed by rank) instead of partial set_rows into a -inf fill — the CUDA in-place set_rows does not preserve an un-written base, which had corrupted decode when n_kv > top_k. - Attention-sink force-inclusion (DSA_SINK, default 1): boost the first key(s) so the sink always survives top-k. The IQ2_M-quantized indexer under-ranks the sink, and masking it collapsed decode; with the boost, top_k=2048 over n_kv>2048 stays coherent. ggml backend fixes (needed by the indexer) - CUDA argsort: report unsupported when padded ncols > 1024 (one-thread-per-column bitonic launch limit) so the scheduler falls back to the CPU argsort. Fixes "invalid configuration argument" for top_k over a large n_kv. - CUDA cpy/dup: support I32 -> I32 (top_k index copies / cross-backend moves). Validation (GLM-5.2-UD-IQ2_M, 3xP100 + --cpu-moe, -mla 1 -fa 0) - c512 PPL = 2.0743, byte-identical to dense (all 8 chunks): no-op path exact. - Short-context decode (300 tok): coherent, identical to dense. - Long-context decode (2521-tok prompt, n_kv>top_k, real masking of ~474 keys, 120+ tok generated): coherent with the sink boost; dense A/B also coherent. Gated behind arch==GLM_DSA + indexer tensors + kr_l cache; DSA_INDEXER_DISABLE=1 forces dense. Remaining: FA path still uses the dense KQ_mask; multi-sequence (n_seq>1) batches; deepseek32 arch wiring. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…-fa 1) The DSA sparse top-k mask is now applied on the -fa 1 path (our serving config), not just -fa 0 soft_max. c512 PPL on -fa 1 = 2.0743, byte-identical to dense (no regression, indexer no-op exact at n_kv <= top_k). Gated arch==GLM_DSA with DSA_INDEXER_DISABLE escape; -fa 0 path unchanged. Long-context -fa 1 decode coherence (n_kv > top_k, mask actually biting) validation is still running at commit time; the FA mask reuses the same full-coverage scatter proven coherent on the -fa 0 decode path, so it should hold, but confirm before relying on long-context -fa 1. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

… multi-seq characterized Document the re-validation after cherry-picking the MLA-FA vec-decode fix (5f18dcc0): - FA path is ALIVE. Long-ctx -fa 1 decode (2521-tok prompt > top_k, mask actively biting) is now COHERENT at -mla 1 and -mla 3, vs the pre-fix degeneration into "0.0.0.0..." repetition. Matches dense (DSA_INDEXER_DISABLE) and -fa 0 controls. - c512 -fa 1 PPL: indexer-ON == dense == 2.0854, byte-identical all 8 chunks (exact no-op when n_kv <= top_k; no regression). The 2.0743->2.0854 shift is the MLA-FA fix changing V accumulation, not an indexer artifact (ON==dense proves it). - Indexer is feature-complete + validated for single-seq prefill+decode on both -fa 0 and -fa 1, at -mla 1 and -mla 3 (the R740 serving target). Remaining PR gaps, characterized honestly: - Multi-seq (n_seq>1) with active mask is BROKEN (n_seq=2 c4096 PPL 62.6 vs dense multi-seq 2.54 and single-seq indexer 3.05). No NaN/crash anymore. Root cause: the indexer uses a single scalar kv_head/n_kv for the whole ubatch; multi-seq needs per-sequence cache writes + per-sequence top-k. Fix deferred (structural). - deepseek32 arch: N/A in this fork. DSA lives entirely under LLM_ARCH_GLM_DSA; there is no LLM_ARCH_DEEPSEEK32 enum. Documented the steps to add one if a real deepseek32 GGUF is ever served. Also commit DSA_REFERENCE.md (verbatim mainline deepseek32/glm-dsa source, the port reference), trimmed of a stray agent-handoff footer. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…eq>1) UPDATE 5. The DSA lightning indexer was numerically broken for multi-sequence batches once the top-k mask bites (n_kv > top_k): c4096 n_seq=2 PPL 62.6 vs dense 2.54, while single-seq was fine. Root cause: the attention-sink force-include boosted the GLOBAL key range [0, n_sink) by +1e20, which only protects sequence 0's sink. With several sequences packed contiguously into one ubatch (seq 0 at cells [0,n0), seq 1 at [n0,n1), ...), every non-first sequence's sink lives at cell n0.. (not cell 0), got no boost, and was dropped from top-k once the mask bites — collapsing that sequence (chunk[2]=61.2 while chunk[1]=2.33). The cache write and score/argsort were already per-sequence correct: tokens are placed contiguously like the main K cache, and the base KQ_mask (filled from kv_self.cells[i].has_seq_id) already drives cross-seq keys to -inf before argsort. Only the sink was anchored at the wrong (global) cell. Fix: replace the global arange sink boost with a per-graph input tensor inp_dsa_sink {n_kv, n_tokens} (F32), filled on the CPU in llama_set_inputs from kv_self.cells exactly like the KQ_mask: inp_dsa_sink[j,i] = 1e20 iff cell[i].pos in [0,n_sink) AND cell[i].has_seq_id(seq_of_query_j), else 0 so each query force-includes only its OWN sequence's sink. For a single contiguous sequence from pos 0 this is exactly the old "cell index < n_sink" set with the same magnitude, so n_seq==1 is byte-identical. Validation (3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, wikitext-2): - c4096 n_seq=2 indexer chunk[2]: 61.2 -> 3.07 (== single-seq 3.05). - c2048 topk=1024 (mask bites): n_seq=4 == n_seq=1 chunk-for-chunk (2.5005/2.6080/2.7759/3.1137 vs .../3.1138) -> multi-seq is numerically identical to processing each sequence alone. - c512 n_seq=1 indexer ON == dense, all 4 chunks byte-identical (no regression). n_seq=4 at full c4096 (n_kv=16384) OOMs the P100 compute buffer (capacity, not correctness; n_seq=4 proven correct at c2048/n_kv=8192). GLM-5.2 DSA indexer is now sequence-correct for n_seq>=1, prefill+decode, soft_max+FA, -mla 1/-mla 3. Fully general and PR-ready. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…across shift/defrag/seq-ops; per-seq sink on first-present pos) An adversarial review found the indexer was proven on the perplexity path but not the serving path: the persistent indexer-K cache kr_l was written/read but never *maintained* by the KV-cache mutators, and the attention sink anchored on absolute pos<n_sink (wrong after multi-turn seq_rm). This closes those gaps and pins down what is actually reachable on the MLA model. kr_l maintenance: - build_k_shift (llama-build-context.cpp): rotate the indexer keys by the same per-cell delta as the main K. The cached key is H*concat(RoPE(k_pe,pos),k_nope), so un-Hadamard (H sym/orthonormal => H*H=I) -> RoPE-delta the pe sub-block -> re-Hadamard. Exact because GLM-DSA has no rope-scaling metadata (ext_factor=0, attn_factor=1, freq_scale=1), so NEOX RoPE is pure/composable. Params mirror the forward indexer RoPE exactly (rope_factors=nullptr); no DEEPSEEK2 yarn-shift leak. Non-in-place (cont->rope->concat->re-Had->cpy), no aliasing. K-shift Hadamard input filled in llama_set_k_shift with the identical Sylvester construction. - build_defrag: kr_l row-move mirrors the k_l move (defrag never changes pos, so no re-RoPE). max_moves divisor 6->9 *n_layer when the indexer cache is present. - seq_rm/seq_cp/seq_keep are metadata-only (verified) so kr_l rows stay matched to cells; seq_add/seq_div set has_shift and route through K-shift. No seq-op change. Per-seq sink (llama.cpp llama_set_inputs): anchor on each sequence's FIRST PRESENT pos (min present pos over the scored n_kv span), not absolute pos<n_sink. After multi-turn seq_rm drops a sequence's early tokens its earliest survivor has pos>=n_sink; the absolute test would protect nothing. Fresh seq at pos 0 => min=0 => byte-identical to the old behaviour. Serving-shift finding (the whole point): a RoPE context-shift on this model is REFUSED BY THE ENGINE. get_can_shift() returns false for all MLA models (is_mla_model() includes GLM_DSA); llama_kv_cache_update returns 1 -> "main : failed to eval". Reproduced AND isolated with a dense control (DSA_INDEXER_DISABLE=1): dense fails identically at the same token. The failure is pre-existing MLA engine behaviour, independent of the indexer. On the MLA path the shift never happens, so the indexer's kr_l can never desync via K-shift; the build_k_shift kr_l block is correct-and-dormant (documented loudly in code). Validation (3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, GGML_CUDA_NO_PINNED=1, numactl --interleave=all, wikitext-2): - No regression: c512 n_seq=1 indexer ON == dense == 2.1957 +/- 0.12031, byte-identical all 4 chunks (2.2770/2.8741/2.3956/2.1957). - Multi-seq: c4096 n_seq=2 chunk[1]=2.33 chunk[2]=3.07 healthy (== UPDATE 5; per-seq sink change did not regress). - Serving shift: engine-refused for MLA, dense control fails identically. - Independent adversarial review: GO, no correctness defect in the diff. - Build clean (llama-cli, llama-perplexity, sm_60). Comments updated (build_deepseek2.cpp): multi-seq+FA no longer limitations; sink description matches per-seq min-pos anchoring; BIG=1e30 masks on both soft_max and FA paths. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…ission for the kr_l indexer cache update_cache_copies() re-points the K/V cache writes to the current kv_head whenever a compute graph is REUSED (can_reuse_graph reuses iff kv_self.n == prev->n_kv). The persistent indexer-key cache write (kr_l) is a separate ggml_cpy whose destination view bakes kv_head at graph-build time, and it was NEVER registered for that fixup. Under FA the cache pads to 256, so consecutive single-token decode ubatches share the same padded n_kv and the graph IS reused; without the fixup the kr_l write keeps landing in the first ubatch's slot and later ubatches never populate their own recent index-key cells (those cells stay at the alloc-zeroed 0.0). Structurally identical to the MiniMax MSA bug (fork commit 133d14c9). Fix (mirrors the K/V cache_copies fixup, same shape as MSA 133d14c9): - llama-context.h: new std::vector<CacheCopy> dsa_cache_copies. - llama.cpp ctor: resize dsa_cache_copies to n_layer (null entries -> no-op when DSA off). - build_deepseek2.cpp: register the kr_l ggml_cpy as dsa_cache_copies[il] = {kr_cpy, kr->nb[1]}. - llama.cpp update_cache_copies(): re-point each registered cpy view_offs = kv_head*step and patch src[1]->data/data, exactly like K/V, with the c.cpy->view_src == kv_self.kr_l[il] (+ null/op) guard the MSA fix omitted. soft_max / non-DSA paths byte-identical. Validation (GLM-5.2-UD-IQ2_M, 3x P100 -ngl 99 --cpu-moe -t 32, NO_PINNED, P2P-disable patch re-applied to get a working multi-GPU baseline — see UPDATE 7.3; that patch was lost in the upstream rebase and is required separately): - c512 -fa1 -mla3 indexer ON: 2.1983 (== prior baseline; build healthy). - Long-ctx FA decode, 2735-tok recall prompt, -mla3 -fa1 temp0, reuse ON (default): coherent, correct deep-context recall ("Dr. Mariana Velasquez ... Daniel Okonkwo") on BOTH the fixed and the unfixed binary. - ub128 PPL -fa1 -mla3 reuse ON, unfixed: 1.7239/1.8211/2.1888/2.4517, healthy (no inflation). Honest scope: the bug is real in code but LATENT for GLM-DSA at its configured top_k=2048 (permissive selection keeps the genuinely-attended recent blocks even when reuse leaves some recent index-key cells stale), unlike MSA's tighter top-k where it inflated PPL ~2x. The fix is correct and prevents the latent corruption from biting at any tighter top_k / longer ctx / future serving config. The pre-P2P-patch "nan" seen at ub128 was P2P corruption, not this bug. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…off by default) Implements ikawrakow's direction from discussion #2040: the DSA sparse indexer must be controllable via command-line argument (not environment variables), and must be OFF by default for now. Control surface, before -> after: DSA_INDEXER_DISABLE (env, inverted: on-by-default) -> --dsa / -dsa (cparams.dsa, default false; opt-in, dense-by-default) DSA_TOPK_OVERRIDE (env) -> --dsa-top-k N / -dsatk N (cparams.dsa_top_k, default -1 == model's configured indexer_top_k) DSA_HADAMARD_DISABLE, DSA_SINK (env) -> kept as DEBUG-ONLY env knobs (clearly commented; no CLI surface, not system on/off controls) Plumbing mirrors existing boolean/int feature flags (-mla, -khad): include/llama.h llama_context_params {bool dsa; int dsa_top_k;} src/llama.cpp default_params (false / -1); cparams assignment src/llama-cparams.h llama_cparams {bool dsa=false; int dsa_top_k=-1;} common/common.h gpt_params {bool dsa=false; int dsa_top_k=-1;} common/common.cpp arg parse + help text + cparams copy src/graphs/build_deepseek2.cpp gate now checks cparams.dsa instead of getenv; top-k override reads cparams.dsa_top_k. Stays arch-gated to LLM_ARCH_GLM_DSA. When --dsa is off (default) the indexer function is never called -> existing dense MLA path, byte-identical to no-feature. Validation (GLM-5.2-UD-IQ2_M, 3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, wikitext-2, 4 chunks @ c2560): --dsa OFF (default, dense): PPL 2.4151 (graph nodes 4166) --dsa ON, default top_k=2048: PPL 2.4697 (graph nodes 8846) --dsa ON, --dsa-top-k 1024: PPL 3.5107 Off-by-default runs the dense path; ON activates the indexer (node count jumps, PPL shifts as the top-k mask bites once n_kv > top_k). No env var is consulted for the primary on/off or the top-k knob. Graph-parallel (-sm graph) interaction (the item ikawrakow flagged): Under -sm graph the MLA layers are TP-split (wo->extra) and route to build_deepseek2_tp_attention(), which contains NO indexer code. So --dsa is silently a NO-OP under -sm graph: it does not error or crash, it runs dense. Empirically, --dsa --dsa-top-k 1024 under -sm graph gives PPL 2.4308 (chunks 1.6967/1.7906/2.1664/2.4308) -- the dense baseline (2.4151), NOT the DSA top_k=1024 numbers (3.5107). The 0.016 delta is f16 TP-reduce numerics, not DSA. Conclusion: DSA "works under deepseek2" only on the non-TP (layer) path; serving DSA with -sm graph would require wiring the indexer into the TP attention path (or a dedicated DSA arch). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…ns dense MLA) The DSA lightning indexer is built only in the layer-mode (non-TP) attention path. Under -sm graph / -sm attn the tensor-parallel attention path has no indexer, so --dsa would silently run dense MLA. Emit a clear one-time LLAMA_LOG_WARN at context creation instead of degrading silently. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

DSA_REFERENCE.md and the R740 progress note are development scratch, not part of the submission. Remove them so the PR diff is code-only. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

PR #2045 adds GLM-DSA sparse attention but was validated on CUDA (--cpu-moe). A CPU-only build (-ngl 0 --dsa) crashes in four spots where the CUDA backend tolerates something the CPU backend does not. These make GLM-5.2 --dsa run coherently on CPU; with --dsa off they are no-ops (DSA CPU path only). 1. set_rows into an F32 dest segfaults (ggml.c set_rows_f32): type_traits[F32].from_float is NULL, so the DSA sparse-mask scatter calls a NULL fn (segfault at ip=0). memcpy when the dest is F32. CUDA has a real F32 set_rows path, so this only bit the CPU build. 2. ggml_add(F32 score, F16 mask) aborts on CPU (build_deepseek2_dsa_indexer and build_deepseek2_dsa_sparse_mask): under -fa 1 the dense KQ_mask is F16 and CPU add only accepts F32+F16 when src0 is F16. Cast the causal mask view to F32. CUDA's add accepts the mixed types. 3. dsa_fa_mask dim-1 concat must be F32 on CPU (build_deepseek2_dsa_fa_mask): CPU ggml_concat only supports F16 along dim 0; do the row (dim-1) concat in F32 then cast the result to F16. CUDA supports the F16 dim-1 concat. 4. indexer k_norm epsilon is 0 -> ggml_norm aborts (llama-hparams.cpp): the lightning-indexer k_norm is a non-RMS LayerNorm using f_norm_eps, but the GLM-DSA GGUF only carries the RMS eps so f_norm_eps stays 0 (GGML_ASSERT(eps > 0)). Mirror the RMS eps. CUDA's norm doesn't assert on eps=0. Validated: GLM-5.2 UD-Q4_K_M, single-socket Xeon w7-2475X, CPU-only (-ngl 0 --dsa) - coherent at 49K+ ctx, correct 30K needle retrieval, prefill flat with length (~32 tok/s, the O(L) DSA signature) vs the dense build's O(L^2) decline. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

usrlocalben · 2026-07-01T18:56:16Z

+        __syncthreads();
+    }
+    for (int i = threadIdx.x; i < nidx; i += blockDim.x) {
+        y_row[idx[i]] = b;


s/idx/idx_row ?

usrlocalben · 2026-07-01T19:15:48Z

Check out 2da6c80

This concept seems to be missing for GLM-5.2 where not all layers are full index layers.

Also, it seems like "heavy" quantization of the index tensors is potentially harmful. The two GGUFs I have on hand have different index quants, one w/Q6 and the other w/Q3. The Q3 index tensors go off the rails w/10K prompt but the Q6 seems OK. Then, the Q6 fails at 100K. Still mostly OK output, but not coherent enough to leave <think> phase.

The OEM's FP8 safetensors distro has the index tensors in FP8.

indexer.weights_proj - BF16
indexer.wk.weight - F8_E4M3 (+scale)
indexer.wq_b.weight - F8_E4M3 (+scale)

I'm pulling the BF16 now to build new GGUFs with Q8 index.

AI disclosure: I gathered the DS32 paper, hf transformers source, ik branches etc. and worked with DSv4-flash to find these things. The writing here is my own.

usrlocalben · 2026-07-01T19:22:51Z

...and there is still a lot of memory usage, but it's usable.
w/200K + b=8192 it fits in ~75GB vram. I observe that at each 8192 block completed during long prefill that VRAM ticks up 1GB (as indicated by nvidia-smi)

Or said a different way, at runtime/post-init an addl 13GB vram is needed to reach 100K during PP

mb8565 · 2026-07-01T23:54:06Z

@usrlocalben you called it on the shared index layers. We landed on the same thing from the other side and can add a measured number for it.

GLM-5.2's indexer_types marks 21 "full" layers (0, 1, 2, then every 4th up to 74) that compute their own lightning-indexer top-k, and 57 "shared" layers that reuse the previous full layer's top-k. As you saw in transformers, the shared layers set self.indexer = None and take topk_indices = prev_topk_indices. Both #2045 and #2063 compute an independent top-k on every layer, because the GGUF ships indexer weights on every layer and nothing marks which are shared, so 57 of 78 layers select a different key set than the model was trained with. DeepSeek-V3.2 has no such split (its config has no indexer_types or per-layer index designation, so its indexer runs on every layer), which is why porting from it looked correct.

Measured it on an unsloth IQ2_M GGUF at 4K, top_k 2048, CPU: making the shared layers reuse the previous full layer's top-k drops perplexity from 3.2102 to 2.7185, against a dense baseline of 2.6972. That closes about 95.8% of the gap. Adding the k_norm eps fix below takes it to 2.7099, about 97.5%. In one long-context generation run the buggy version drifts into repetition and the fix does not, though that is a single sample, not multi-turn serving. As a control, the machinery is otherwise correct: with top_k forced at or above n_kv the sparse path matches dense perplexity exactly, so this is purely which keys get selected.

On the index-tensor quant you raised, separate from IndexShare: both unsloth GGUFs we tested (IQ2_M and Q3_K_XL) keep the index tensors at Q8_0, so they show IndexShare cleanly with no index-quant confound, which is why our Q3_K_XL still degraded before the fix. IndexShare is engine-side, not in the weights, and no current GGUF inference path (#2045, #2063, mainline) implements it, so your Q3 build almost certainly carries the same missing-reuse degradation on top of whatever the Q3 index itself costs. A Q8-index build with the reuse fix would separate the two cleanly. One caveat: we have only tested to 10K, not your 100K regime, so I would not assume 100K is solved.

Two things still open here: the indexer k_norm eps should be 1e-6 (the reference hardcodes it) where we fall back to the model's 1e-5 rms eps, and we have measured prefill perplexity plus a long-context generation check, not full multi-turn serving yet. The clean fix is to carry indexer_types (or the freq/offset pair) into the GGUF in conversion and drive the reuse from it, always on, rather than the hardcoded layer rule we used to prove it out.

GLM-5.2's indexer_types marks 21 'full' layers that compute their own lightning-indexer top-k and 57 'shared' layers that reuse the previous full layer's top-k. This port computed an independent top-k on every layer, which mis-selects keys on the 57 shared layers (the transformers reference sets indexer=None on shared layers and reuses prev_topk). Shared layers now reuse the most-recent full layer's selection. Full/ shared map derived from the config rule (full iff il<=1 or il%4==2), which reproduces indexer_types exactly; loader can later override from GGUF metadata. Built on ikawrakow#2063's tree; head-loop/ggml_hadamard/ggml_blend/ argsort/FA-mask unchanged. 4K PPL (unsloth IQ2_M, top_k 2048, CPU): DSA-on 3.1922 -> 2.7111, dense 2.6972 (~97% of the gap). top_k>=n_kv reproduces dense exactly. Single- seq and 4x8 parallel decode coherent. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

mb8565 · 2026-07-02T03:17:49Z

We tracked the coherence bug down to IndexShare, which @usrlocalben also identified. We rebased #2045 onto your #2063 branch so it carries your head-loop, ggml_hadamard, ggml_blend, argsort, and FA-mask work, with the IndexShare fix on top.

What the fix does: the 21 "full" indexer layers (from config indexer_types) compute their own top-k as before, and the 57 "shared" layers now reuse the previous full layer's top-k instead of computing their own. In our branch it is always on, since it is a correctness fix rather than a tunable, but if you would rather gate it behind a flag until it has run on the 100K regime and on GPU, that is a small change to add. The full/shared map is derived from the config rule (full iff layer <= 1 or layer % 4 == 2), which reproduces GLM-5.2's indexer_types exactly, and the loader is set up so a future GGUF metadata key can override that derivation for variants with a different pattern. It is a small change: a per-layer flag in hparams, a last-full-sorted tensor that persists and gets reset on each graph build, and the indexer call site branches on full vs shared.

Results on an unsloth IQ2_M GGUF (Q8_0 index tensors), 4K, top_k 2048, CPU:

DSA-on before: 3.1922
DSA-on with IndexShare: 2.7111
dense: 2.6972

That closes about 97% of the gap. As a control, forcing top_k >= n_kv reproduces dense exactly on this tree (2.6972, chunk for chunk), so the remainder is ordinary sparse approximation and not a path bug. Single-sequence generation at a ~9K prompt stays coherent, and a 4-client / 8-sequence parallel run stays coherent with no cache issues, so the shared-layer reuse holds up under batching, which the single perplexity run on its own doesn't prove. All of this is 4K on CPU with one IQ2_M GGUF. We have not run the 100K regime and have not run it on GPU, so those are still open on our side.

A few things to mention:

eps: we looked at the indexer k_norm eps (HF hardcodes 1e-6, we mirror the model RMS eps 1e-5). 1e-6 vs 1e-5 is within the 4-chunk PPL noise on this tree, so we left the mirror. Not part of the fix.
MTP: index_share_for_mtp_iteration is true in the config, but the MTP tail layer stays dense (unchanged from your tree). Not implemented here.
Tensor-parallel: DSA is still absent under -sm graph/attn (pre-existing).
Defrag: forcing a defrag session aborts on a ggml view-bounds assert, but this reproduces identically in dense mode with no --dsa on this tree, so it looks like a pre-existing defrag issue orthogonal to IndexShare, not something this change introduces. Defrag is off by default (-dt -1).
Index-tensor quant (the separate axis @usrlocalben raised): all our tests use Q8_0 index tensors, so this isolates IndexShare cleanly.

mb8565 · 2026-07-02T05:55:18Z

Separate from IndexShare, we ran into another issue while testing: with -fa 0 (no flash attention), the DSA path produces garbage perplexity. On the unsloth IQ2_M GGUF, 4K, top_k 2048, CPU:

-mla 3 -fa 0 --dsa: PPL ~4350
-mla 3 -fa 0 dense (no --dsa): 2.69
-mla 3 -fa 1 --dsa: 2.71

So it is specific to the -fa 0 DSA path, which routes through build_deepseek2_dsa_sparse_mask rather than the -fa 1 ggml_blend mask path, and not dense.

It is not IndexShare, it breaks the same with the shared-layer reuse off, and not the mask base-fill, byte-identical with it on or off. It also works on the prior #2045 base at -fa 0 --dsa (PPL 3.16, in line with that base's -fa 1), so it looks like the -fa 0 path regressed somewhere in #2063 rather than being a prior issue.

-fa 1 is unaffected, so this only matters for anyone running -fa 0 with --dsa. We are digging into where in #2063 it regressed, but it may be a bit before we have a fix.

@ikawrakow

…y default (#2045) * Add GLM-5.2/DeepSeek-V3.2 DSA lightning indexer (batch-local, single-seq prefill) Implements the sparse top-k "lightning indexer" attention for LLM_ARCH_GLM_DSA in build_deepseek2_layer_attention (ik's deepseek2 graph). What it does (per layer, gated on model.arch==GLM_DSA && indexer_attn_q_b): - indexer_q = indexer_attn_q_b(q_lora latent), split rope(64)/nope(64), NEOX-rope the pe part, concat. indexer_k = indexer_attn_k(attn_norm out), LayerNorm w/ bias, same rope/concat (single key head, MQA). - scores = relu(indexer_k . indexer_q), scaled per-head weights (indexer_proj), summed over heads, + base causal mask, then ggml_top_k(min(top_k, n_tokens)). - sparse mask: ggml_fill(-inf) -> ggml_set_rows(0) at top_k positions -> + causal, used in the soft_max_ext attention path (-mla 1 -fa 0) instead of KQ_mask. Simplifications (intentional, proven sound): - Batch-local: no indexer KV-cache. Indexer keys are the current batch tokens. - Walsh-Hadamard transform omitted: orthonormal rotation, (Hq).(Hk)==q.k, no score change. Validation (GLM-5.2-UD-IQ2_M, 3x P100, -mla 1 -fa 0): - Compiles clean (CUDA sm_60); loads and runs. - c512 -b512 (n_seq=1) PPL = 2.7760, byte-identical to dense baseline (indexer disabled) = 2.7760, all 8 chunks match -> indexer is an exact no-op when top_k>=n_tokens. Proves correctness-preservation. - 3105-token prompt completion (top_k=2048 < 3105 -> indexer ACTIVELY masks): prompt-eval produces coherent, accurate continuation, identical to dense for the prompt+early-gen tokens. No NaN/crash. Confirms the masking path works in prefill. Known limitations (documented follow-ups, NOT handled): - Single-sequence prefill only. Multi-sequence batches (n_seq>1, e.g. perplexity default n_batch>n_ctx) and kv_head>0 (decode) break the batch-local key->slot mapping. n_seq>1 -> NaN (use n_batch==n_ctx). Decode (kv_head>0): each generated token sees only itself as an indexer key, so generation degenerates into repetition after the prompt (dense A/B stays coherent) -- this is the decode-cache stub, the documented next step. - Flash-attn path (-fa 1, F16 mask) still uses dense KQ_mask (soft_max path only). - Decode indexer KV-cache + Hadamard cached-K storage not implemented. Runtime gate: DSA_INDEXER_DISABLE=1 falls back to dense attention (for A/B). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: decode-correct via persistent indexer-K cache Make the lightning-indexer correct for DECODE (not just prefill). Previously the indexer was batch-local, so a generated token only scored against itself and generation degenerated. Now the indexer keys are cached across the full context. Changes - llama_kv_cache: add per-layer indexer-key cache `kr_l` [indexer_head_size, kv_size] (F16, MQA single head), allocated alongside the MLA latent cache for GLM_DSA. - build_deepseek2_dsa_indexer: write the batch's (Hadamard-rotated) indexer keys to kr_l at kv_head, read back the full [128, n_kv] cached keys, and score the indexer queries against ALL past keys. Returns the full descending argsort of the scores. - Walsh-Hadamard rotation of indexer q/k (cparams.dsa_indexer_hadamard, default on; filled in llama_set_inputs). Score-preserving; improves cached-K F16 precision. - build_deepseek2_dsa_sparse_mask: rank-based full-coverage scatter (write a 0/-BIG penalty into EVERY key slot keyed by rank) instead of partial set_rows into a -inf fill — the CUDA in-place set_rows does not preserve an un-written base, which had corrupted decode when n_kv > top_k. - Attention-sink force-inclusion (DSA_SINK, default 1): boost the first key(s) so the sink always survives top-k. The IQ2_M-quantized indexer under-ranks the sink, and masking it collapsed decode; with the boost, top_k=2048 over n_kv>2048 stays coherent. ggml backend fixes (needed by the indexer) - CUDA argsort: report unsupported when padded ncols > 1024 (one-thread-per-column bitonic launch limit) so the scheduler falls back to the CPU argsort. Fixes "invalid configuration argument" for top_k over a large n_kv. - CUDA cpy/dup: support I32 -> I32 (top_k index copies / cross-backend moves). Validation (GLM-5.2-UD-IQ2_M, 3xP100 + --cpu-moe, -mla 1 -fa 0) - c512 PPL = 2.0743, byte-identical to dense (all 8 chunks): no-op path exact. - Short-context decode (300 tok): coherent, identical to dense. - Long-context decode (2521-tok prompt, n_kv>top_k, real masking of ~474 keys, 120+ tok generated): coherent with the sink boost; dense A/B also coherent. Gated behind arch==GLM_DSA + indexer tensors + kr_l cache; DSA_INDEXER_DISABLE=1 forces dense. Remaining: FA path still uses the dense KQ_mask; multi-sequence (n_seq>1) batches; deepseek32 arch wiring. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: wire sparse mask into the flash-attention path (-fa 1) The DSA sparse top-k mask is now applied on the -fa 1 path (our serving config), not just -fa 0 soft_max. c512 PPL on -fa 1 = 2.0743, byte-identical to dense (no regression, indexer no-op exact at n_kv <= top_k). Gated arch==GLM_DSA with DSA_INDEXER_DISABLE escape; -fa 0 path unchanged. Long-context -fa 1 decode coherence (n_kv > top_k, mask actually biting) validation is still running at commit time; the FA mask reuses the same full-coverage scatter proven coherent on the -fa 0 decode path, so it should hold, but confirm before relying on long-context -fa 1. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: UPDATE 4 — MLA-FA fix merged, FA path validated, multi-seq characterized Document the re-validation after cherry-picking the MLA-FA vec-decode fix (5f18dcc0): - FA path is ALIVE. Long-ctx -fa 1 decode (2521-tok prompt > top_k, mask actively biting) is now COHERENT at -mla 1 and -mla 3, vs the pre-fix degeneration into "0.0.0.0..." repetition. Matches dense (DSA_INDEXER_DISABLE) and -fa 0 controls. - c512 -fa 1 PPL: indexer-ON == dense == 2.0854, byte-identical all 8 chunks (exact no-op when n_kv <= top_k; no regression). The 2.0743->2.0854 shift is the MLA-FA fix changing V accumulation, not an indexer artifact (ON==dense proves it). - Indexer is feature-complete + validated for single-seq prefill+decode on both -fa 0 and -fa 1, at -mla 1 and -mla 3 (the R740 serving target). Remaining PR gaps, characterized honestly: - Multi-seq (n_seq>1) with active mask is BROKEN (n_seq=2 c4096 PPL 62.6 vs dense multi-seq 2.54 and single-seq indexer 3.05). No NaN/crash anymore. Root cause: the indexer uses a single scalar kv_head/n_kv for the whole ubatch; multi-seq needs per-sequence cache writes + per-sequence top-k. Fix deferred (structural). - deepseek32 arch: N/A in this fork. DSA lives entirely under LLM_ARCH_GLM_DSA; there is no LLM_ARCH_DEEPSEEK32 enum. Documented the steps to add one if a real deepseek32 GGUF is ever served. Also commit DSA_REFERENCE.md (verbatim mainline deepseek32/glm-dsa source, the port reference), trimmed of a stray agent-handoff footer. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: per-sequence attention sink — fix multi-seq (n_seq>1) UPDATE 5. The DSA lightning indexer was numerically broken for multi-sequence batches once the top-k mask bites (n_kv > top_k): c4096 n_seq=2 PPL 62.6 vs dense 2.54, while single-seq was fine. Root cause: the attention-sink force-include boosted the GLOBAL key range [0, n_sink) by +1e20, which only protects sequence 0's sink. With several sequences packed contiguously into one ubatch (seq 0 at cells [0,n0), seq 1 at [n0,n1), ...), every non-first sequence's sink lives at cell n0.. (not cell 0), got no boost, and was dropped from top-k once the mask bites — collapsing that sequence (chunk[2]=61.2 while chunk[1]=2.33). The cache write and score/argsort were already per-sequence correct: tokens are placed contiguously like the main K cache, and the base KQ_mask (filled from kv_self.cells[i].has_seq_id) already drives cross-seq keys to -inf before argsort. Only the sink was anchored at the wrong (global) cell. Fix: replace the global arange sink boost with a per-graph input tensor inp_dsa_sink {n_kv, n_tokens} (F32), filled on the CPU in llama_set_inputs from kv_self.cells exactly like the KQ_mask: inp_dsa_sink[j,i] = 1e20 iff cell[i].pos in [0,n_sink) AND cell[i].has_seq_id(seq_of_query_j), else 0 so each query force-includes only its OWN sequence's sink. For a single contiguous sequence from pos 0 this is exactly the old "cell index < n_sink" set with the same magnitude, so n_seq==1 is byte-identical. Validation (3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, wikitext-2): - c4096 n_seq=2 indexer chunk[2]: 61.2 -> 3.07 (== single-seq 3.05). - c2048 topk=1024 (mask bites): n_seq=4 == n_seq=1 chunk-for-chunk (2.5005/2.6080/2.7759/3.1137 vs .../3.1138) -> multi-seq is numerically identical to processing each sequence alone. - c512 n_seq=1 indexer ON == dense, all 4 chunks byte-identical (no regression). n_seq=4 at full c4096 (n_kv=16384) OOMs the P100 compute buffer (capacity, not correctness; n_seq=4 proven correct at c2048/n_kv=8192). GLM-5.2 DSA indexer is now sequence-correct for n_seq>=1, prefill+decode, soft_max+FA, -mla 1/-mla 3. Fully general and PR-ready. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: UPDATE 6 — serving-correctness (kr_l maintained across shift/defrag/seq-ops; per-seq sink on first-present pos) An adversarial review found the indexer was proven on the perplexity path but not the serving path: the persistent indexer-K cache kr_l was written/read but never *maintained* by the KV-cache mutators, and the attention sink anchored on absolute pos<n_sink (wrong after multi-turn seq_rm). This closes those gaps and pins down what is actually reachable on the MLA model. kr_l maintenance: - build_k_shift (llama-build-context.cpp): rotate the indexer keys by the same per-cell delta as the main K. The cached key is H*concat(RoPE(k_pe,pos),k_nope), so un-Hadamard (H sym/orthonormal => H*H=I) -> RoPE-delta the pe sub-block -> re-Hadamard. Exact because GLM-DSA has no rope-scaling metadata (ext_factor=0, attn_factor=1, freq_scale=1), so NEOX RoPE is pure/composable. Params mirror the forward indexer RoPE exactly (rope_factors=nullptr); no DEEPSEEK2 yarn-shift leak. Non-in-place (cont->rope->concat->re-Had->cpy), no aliasing. K-shift Hadamard input filled in llama_set_k_shift with the identical Sylvester construction. - build_defrag: kr_l row-move mirrors the k_l move (defrag never changes pos, so no re-RoPE). max_moves divisor 6->9 *n_layer when the indexer cache is present. - seq_rm/seq_cp/seq_keep are metadata-only (verified) so kr_l rows stay matched to cells; seq_add/seq_div set has_shift and route through K-shift. No seq-op change. Per-seq sink (llama.cpp llama_set_inputs): anchor on each sequence's FIRST PRESENT pos (min present pos over the scored n_kv span), not absolute pos<n_sink. After multi-turn seq_rm drops a sequence's early tokens its earliest survivor has pos>=n_sink; the absolute test would protect nothing. Fresh seq at pos 0 => min=0 => byte-identical to the old behaviour. Serving-shift finding (the whole point): a RoPE context-shift on this model is REFUSED BY THE ENGINE. get_can_shift() returns false for all MLA models (is_mla_model() includes GLM_DSA); llama_kv_cache_update returns 1 -> "main : failed to eval". Reproduced AND isolated with a dense control (DSA_INDEXER_DISABLE=1): dense fails identically at the same token. The failure is pre-existing MLA engine behaviour, independent of the indexer. On the MLA path the shift never happens, so the indexer's kr_l can never desync via K-shift; the build_k_shift kr_l block is correct-and-dormant (documented loudly in code). Validation (3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, GGML_CUDA_NO_PINNED=1, numactl --interleave=all, wikitext-2): - No regression: c512 n_seq=1 indexer ON == dense == 2.1957 +/- 0.12031, byte-identical all 4 chunks (2.2770/2.8741/2.3956/2.1957). - Multi-seq: c4096 n_seq=2 chunk[1]=2.33 chunk[2]=3.07 healthy (== UPDATE 5; per-seq sink change did not regress). - Serving shift: engine-refused for MLA, dense control fails identically. - Independent adversarial review: GO, no correctness defect in the diff. - Build clean (llama-cli, llama-perplexity, sm_60). Comments updated (build_deepseek2.cpp): multi-seq+FA no longer limitations; sink description matches per-seq min-pos anchoring; BIG=1e30 masks on both soft_max and FA paths. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-5.2 DSA indexer: UPDATE 7 — FIX latent graph-reuse cache-fixup omission for the kr_l indexer cache update_cache_copies() re-points the K/V cache writes to the current kv_head whenever a compute graph is REUSED (can_reuse_graph reuses iff kv_self.n == prev->n_kv). The persistent indexer-key cache write (kr_l) is a separate ggml_cpy whose destination view bakes kv_head at graph-build time, and it was NEVER registered for that fixup. Under FA the cache pads to 256, so consecutive single-token decode ubatches share the same padded n_kv and the graph IS reused; without the fixup the kr_l write keeps landing in the first ubatch's slot and later ubatches never populate their own recent index-key cells (those cells stay at the alloc-zeroed 0.0). Structurally identical to the MiniMax MSA bug (fork commit 133d14c9). Fix (mirrors the K/V cache_copies fixup, same shape as MSA 133d14c9): - llama-context.h: new std::vector<CacheCopy> dsa_cache_copies. - llama.cpp ctor: resize dsa_cache_copies to n_layer (null entries -> no-op when DSA off). - build_deepseek2.cpp: register the kr_l ggml_cpy as dsa_cache_copies[il] = {kr_cpy, kr->nb[1]}. - llama.cpp update_cache_copies(): re-point each registered cpy view_offs = kv_head*step and patch src[1]->data/data, exactly like K/V, with the c.cpy->view_src == kv_self.kr_l[il] (+ null/op) guard the MSA fix omitted. soft_max / non-DSA paths byte-identical. Validation (GLM-5.2-UD-IQ2_M, 3x P100 -ngl 99 --cpu-moe -t 32, NO_PINNED, P2P-disable patch re-applied to get a working multi-GPU baseline — see UPDATE 7.3; that patch was lost in the upstream rebase and is required separately): - c512 -fa1 -mla3 indexer ON: 2.1983 (== prior baseline; build healthy). - Long-ctx FA decode, 2735-tok recall prompt, -mla3 -fa1 temp0, reuse ON (default): coherent, correct deep-context recall ("Dr. Mariana Velasquez ... Daniel Okonkwo") on BOTH the fixed and the unfixed binary. - ub128 PPL -fa1 -mla3 reuse ON, unfixed: 1.7239/1.8211/2.1888/2.4517, healthy (no inflation). Honest scope: the bug is real in code but LATENT for GLM-DSA at its configured top_k=2048 (permissive selection keeps the genuinely-attended recent blocks even when reuse leaves some recent index-key cells stale), unlike MSA's tighter top-k where it inflated PPL ~2x. The fix is correct and prevents the latent corruption from biting at any tighter top_k / longer ctx / future serving config. The pre-P2P-patch "nan" seen at ub128 was P2P corruption, not this bug. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-DSA: convert sparse-attention control from env vars to CLI args (off by default) Implements ikawrakow's direction from discussion #2040: the DSA sparse indexer must be controllable via command-line argument (not environment variables), and must be OFF by default for now. Control surface, before -> after: DSA_INDEXER_DISABLE (env, inverted: on-by-default) -> --dsa / -dsa (cparams.dsa, default false; opt-in, dense-by-default) DSA_TOPK_OVERRIDE (env) -> --dsa-top-k N / -dsatk N (cparams.dsa_top_k, default -1 == model's configured indexer_top_k) DSA_HADAMARD_DISABLE, DSA_SINK (env) -> kept as DEBUG-ONLY env knobs (clearly commented; no CLI surface, not system on/off controls) Plumbing mirrors existing boolean/int feature flags (-mla, -khad): include/llama.h llama_context_params {bool dsa; int dsa_top_k;} src/llama.cpp default_params (false / -1); cparams assignment src/llama-cparams.h llama_cparams {bool dsa=false; int dsa_top_k=-1;} common/common.h gpt_params {bool dsa=false; int dsa_top_k=-1;} common/common.cpp arg parse + help text + cparams copy src/graphs/build_deepseek2.cpp gate now checks cparams.dsa instead of getenv; top-k override reads cparams.dsa_top_k. Stays arch-gated to LLM_ARCH_GLM_DSA. When --dsa is off (default) the indexer function is never called -> existing dense MLA path, byte-identical to no-feature. Validation (GLM-5.2-UD-IQ2_M, 3x P100, -ngl 99 --cpu-moe -mla 3 -fa 1, wikitext-2, 4 chunks @ c2560): --dsa OFF (default, dense): PPL 2.4151 (graph nodes 4166) --dsa ON, default top_k=2048: PPL 2.4697 (graph nodes 8846) --dsa ON, --dsa-top-k 1024: PPL 3.5107 Off-by-default runs the dense path; ON activates the indexer (node count jumps, PPL shifts as the top-k mask bites once n_kv > top_k). No env var is consulted for the primary on/off or the top-k knob. Graph-parallel (-sm graph) interaction (the item ikawrakow flagged): Under -sm graph the MLA layers are TP-split (wo->extra) and route to build_deepseek2_tp_attention(), which contains NO indexer code. So --dsa is silently a NO-OP under -sm graph: it does not error or crash, it runs dense. Empirically, --dsa --dsa-top-k 1024 under -sm graph gives PPL 2.4308 (chunks 1.6967/1.7906/2.1664/2.4308) -- the dense baseline (2.4151), NOT the DSA top_k=1024 numbers (3.5107). The 0.016 delta is f16 TP-reduce numerics, not DSA. Conclusion: DSA "works under deepseek2" only on the non-TP (layer) path; serving DSA with -sm graph would require wiring the indexer into the TP attention path (or a dedicated DSA arch). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-DSA: warn that --dsa is inactive under -sm graph/attn (TP path runs dense MLA) The DSA lightning indexer is built only in the layer-mode (non-TP) attention path. Under -sm graph / -sm attn the tensor-parallel attention path has no indexer, so --dsa would silently run dense MLA. Emit a clear one-time LLAMA_LOG_WARN at context creation instead of degrading silently. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-DSA: drop in-tree dev reference docs from the PR branch DSA_REFERENCE.md and the R740 progress note are development scratch, not part of the submission. Remove them so the PR diff is code-only. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * GLM-DSA: fix CPU-only crashes in the sparse-attention path PR #2045 adds GLM-DSA sparse attention but was validated on CUDA (--cpu-moe). A CPU-only build (-ngl 0 --dsa) crashes in four spots where the CUDA backend tolerates something the CPU backend does not. These make GLM-5.2 --dsa run coherently on CPU; with --dsa off they are no-ops (DSA CPU path only). 1. set_rows into an F32 dest segfaults (ggml.c set_rows_f32): type_traits[F32].from_float is NULL, so the DSA sparse-mask scatter calls a NULL fn (segfault at ip=0). memcpy when the dest is F32. CUDA has a real F32 set_rows path, so this only bit the CPU build. 2. ggml_add(F32 score, F16 mask) aborts on CPU (build_deepseek2_dsa_indexer and build_deepseek2_dsa_sparse_mask): under -fa 1 the dense KQ_mask is F16 and CPU add only accepts F32+F16 when src0 is F16. Cast the causal mask view to F32. CUDA's add accepts the mixed types. 3. dsa_fa_mask dim-1 concat must be F32 on CPU (build_deepseek2_dsa_fa_mask): CPU ggml_concat only supports F16 along dim 0; do the row (dim-1) concat in F32 then cast the result to F16. CUDA supports the F16 dim-1 concat. 4. indexer k_norm epsilon is 0 -> ggml_norm aborts (llama-hparams.cpp): the lightning-indexer k_norm is a non-RMS LayerNorm using f_norm_eps, but the GLM-DSA GGUF only carries the RMS eps so f_norm_eps stays 0 (GGML_ASSERT(eps > 0)). Mirror the RMS eps. CUDA's norm doesn't assert on eps=0. Validated: GLM-5.2 UD-Q4_K_M, single-socket Xeon w7-2475X, CPU-only (-ngl 0 --dsa) - coherent at 49K+ ctx, correct 30K needle retrieval, prefill flat with length (~32 tok/s, the O(L) DSA signature) vs the dense build's O(L^2) decline. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * DSA: loop over attention heads + use builtin Hadamard * DSA: ggml_blend * DSA: remove a bunch of unnecessary ggml_cont * DSA: fix CUDA blend - but something is still wrong * DSA: use ggml_top_k instead of ggml_argsort when FA is ON * CUDA: add CUB based argsort * DSA: avoid graph leaves * Various * GLM-5.2 DSA: IndexShare (shared layers reuse full-layer top-k) GLM-5.2's indexer_types marks 21 'full' layers that compute their own lightning-indexer top-k and 57 'shared' layers that reuse the previous full layer's top-k. This port computed an independent top-k on every layer, which mis-selects keys on the 57 shared layers (the transformers reference sets indexer=None on shared layers and reuses prev_topk). Shared layers now reuse the most-recent full layer's selection. Full/ shared map derived from the config rule (full iff il<=1 or il%4==2), which reproduces indexer_types exactly; loader can later override from GGUF metadata. Built on #2063's tree; head-loop/ggml_hadamard/ggml_blend/ argsort/FA-mask unchanged. 4K PPL (unsloth IQ2_M, top_k 2048, CPU): DSA-on 3.1922 -> 2.7111, dense 2.6972 (~97% of the gap). top_k>=n_kv reproduces dense exactly. Single- seq and 4x8 parallel decode coherent. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * Apply suggestion from @ikawrakow --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: mgkwill <168222+mgkwill@users.noreply.github.com> Co-authored-by: Kawrakow <iwankawrakow@gmail.com>

ikawrakow · 2026-07-02T07:48:03Z

The changes made here were merged in #2045 -> closing

sayap · 2026-07-02T23:43:56Z

Then, the Q6 fails at 100K. Still mostly OK output, but not coherent enough to leave phase.

I have uploaded another GGUF that has Q8 indexer, and it is only about 5MiB larger 😅

Can give it a try if the BF16 download is still ongoing.

I only have enough RAM + VRAM to fit 100k Q6 context for now, and previously with dense attention the model was coherent up until the point the 100k context got filled.

usrlocalben · 2026-07-03T00:28:54Z

Then, the Q6 fails at 100K. Still mostly OK output, but not coherent enough to leave phase.

I have uploaded another GGUF that has Q8 indexer, and it is only about 5MiB larger 😅

Can give it a try if the BF16 download is still ongoing.

I only have enough RAM + VRAM to fit 100k Q6 context for now, and previously with dense attention the model was coherent up until the point the 100k context got filled.

I built one with q4/q4/q6 _exps, bf16 index, and your q6 params for the remainder. It still falls apart with -dsa and a 100K test. There seems to be more to do here.

mb8565 and others added 19 commits June 30, 2026 18:29

GLM-DSA: drop in-tree dev reference docs from the PR branch

3b20dbc

DSA_REFERENCE.md and the R740 progress note are development scratch, not part of the submission. Remove them so the PR diff is code-only. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

DSA: loop over attention heads + use builtin Hadamard

0549b55

DSA: ggml_blend

914da63

DSA: remove a bunch of unnecessary ggml_cont

87a783b

DSA: fix CUDA blend - but something is still wrong

4f73639

DSA: use ggml_top_k instead of ggml_argsort when FA is ON

ce391f7

CUDA: add CUB based argsort

bbac552

DSA: avoid graph leaves

4df1d31

Various

06e6530

ikawrakow mentioned this pull request Jul 1, 2026

Reduce memory required for DSA #2058

Closed

usrlocalben reviewed Jul 1, 2026

View reviewed changes

joelfarthing mentioned this pull request Jul 2, 2026

model: add openPangu-2.0-Flash (92B-A6B) with MLA-latent cache, DSA/SWA, mHC, and multi-head MTP #2065

Merged

4 tasks

mb8565 mentioned this pull request Jul 2, 2026

deepseek2 : GLM-DSA sparse attention (lightning indexer), --dsa off by default #2045

Merged

ikawrakow closed this Jul 2, 2026

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Another DSA attempt#2063

Another DSA attempt#2063
ikawrakow wants to merge 19 commits into
mainfrom
ik/dsa_loop_hadamard_blend

ikawrakow commented Jul 1, 2026

Uh oh!

usrlocalben Jul 1, 2026

Uh oh!

usrlocalben commented Jul 1, 2026 •

edited

Loading

Uh oh!

usrlocalben commented Jul 1, 2026 •

edited

Loading

Uh oh!

mb8565 commented Jul 1, 2026

Uh oh!

mb8565 commented Jul 2, 2026

Uh oh!

mb8565 commented Jul 2, 2026

Uh oh!

ikawrakow commented Jul 2, 2026

Uh oh!

sayap commented Jul 2, 2026

Uh oh!

usrlocalben commented Jul 3, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

ikawrakow commented Jul 1, 2026

Uh oh!

usrlocalben Jul 1, 2026

Choose a reason for hiding this comment

Uh oh!

usrlocalben commented Jul 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

usrlocalben commented Jul 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

mb8565 commented Jul 1, 2026

Uh oh!

mb8565 commented Jul 2, 2026

Uh oh!

mb8565 commented Jul 2, 2026

Uh oh!

ikawrakow commented Jul 2, 2026

Uh oh!

sayap commented Jul 2, 2026

Uh oh!

usrlocalben commented Jul 3, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

usrlocalben commented Jul 1, 2026 •

edited

Loading

usrlocalben commented Jul 1, 2026 •

edited

Loading