* CUDA: use the mask tensor stride (nb31) in the tile FA kernels, not ne11
The tile flash-attention kernels (no-tensor-core GPUs, e.g. Pascal/sm_60) indexed
the KQ mask using ne11 (= K->ne[1]) as the row stride. The SWA windowing in
ggml_cuda_flash_attn_ext (the n_swa branch) re-points K/V/mask to the last nton
tokens, setting ne11 = nton while the mask keeps its original row stride nb31.
Indexing by ne11 then reads across mask rows and yields NaN once the window slice
engages (context past nton). Use the mask's own stride nb31/sizeof(half); in the
non-sliced case this equals ne11, so it is a no-op there. Matches how the vec
kernel already computes its mask offset, and how mainline llama.cpp fixed the same
latent bug in its unified tile kernel.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* CUDA: apply attention sinks in the tile FA kernels (gpt-oss)
The tile flash-attention kernels took a sinks argument but never applied it, so
gpt-oss (which has a per-head sink logit) got a wrong softmax denominator with
-fa 1 while the -fa 0 path (ggml_soft_max_add_sinks) matched CPU. Apply the sink
after the KV loop, mirroring the vec kernel: the sink joins the running max and
adds exp(sink - max) to the denominator once, rescaling kqsum and VKQ. Only ip==0
adds it so it is counted once across the parallel_blocks KV split (the epilogue
writes the sink-inclusive max/denominator into dst_meta before the combine). The
per-head index is blockIdx.y (the same index the slope uses). Guarded by non-null
sinks, so non-sink models are unaffected. This is the tile-kernel equivalent of
mainline llama.cpp PR #15178; the wmma kernel is left untouched (no Turing/Volta
hardware to validate here, but the same defect likely applies).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
On GPUs without FP16 tensor cores (Pascal / sm_60, e.g. Tesla P100) MLA
flash-attention decode falls back to the CPU. The !fp16_mma_available path
routes decode to the f16 vector kernel, whose is_supported check requires
K == V head sizes; MLA's absorbed head sizes are 576/512 (asymmetric), so it is
rejected and attention runs on the CPU. With --cpu-moe that recomputes the full
MLA attention on the CPU every decoded token, which dominates decode at long
context.
Route Pascal MLA decode (Q->ne[1] <= 8 && K == 576 && V == 512) to the f32
vector kernel and enable that kernel for the 576/512 case, including Q8_0 KV.
Scope: decode only (batch <= 8). Prefill (batch > 8) and -fa 0 are untouched;
tensor-core GPUs never reach this branch. Aligned head sizes are byte-identical
(the asymmetric/Q8_0 work folds to a no-op at compile time), so no other model
or configuration is affected.
- fattn.cu: route 576/512 decode to vec_f32 in the !fp16_mma dispatch and its
is_supported mirror.
- fattn-vec-f32.cu/.cuh: accept + instantiate 576/512 (F16 and Q8_0); fix latent
issues exposed by the first asymmetric/large-head use (KQ-row granularity uses
FATTN_KQ_STRIDE not Dv; guard the dst store to tid < Dv; guard the softmax exp
on the KV tail; size Q_i32 by ceil; only convert K/V to F16 when the type is
F16). All are no-ops for the previously-exercised symmetric cases.
- fattn-vec-f32.cuh / fattn-vec-common.cuh: guard the Q8_0 ragged tail (Dk=576 is
144 int32 lanes = 4.5 warps) with the ragged-dim idiom; compile-time-constant
for aligned head dims, so it folds away.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* GLM-DSA: improve TG performance even more
* Fix crash when using MTP
* GLM-DSA: much better PP performance (CPU-only)
* GLM-DSA: allow for quantized indexer cache
Commit 75a5f6d0 (Per model CUDA contexts) added a third parameter
"model" to ggml_backend_cuda_init, but rpc-server was missed.
The RPC server creates backends at startup before any model is
loaded, so passing nullptr for the model parameter is correct: the
RPC server acts as a pure compute proxy and does not host models
locally. The nullptr serves as the context grouping key in
all_ctx, which is sufficient since all RPC backends in the same
process share the same group for multi-GPU reduce operations.
Fixes: cmake build failure with GGML_CUDA=ON + GGML_RPC=ON
The scalar fast-path in ggml_cuda_op_mul routes to ggml_cuda_op_scale_tensor,
which reads src0 and writes dst as flat contiguous buffers. With a non-contiguous
src0 (for example a row-gapped view) this ignored the per-row strides: only the
first row was correct and later rows read the wrong memory. The CPU backend
respects the strides, so the two backends diverged.
Guard the fast-path on contiguous src0 and dst; non-contiguous inputs now fall
through to the general bin_bcast path, which honours the strides.
Add a non-contiguous test_bin_bcast variant covering both the scalar (scale) and
vector (general) paths.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
The batch >8 indexer path in build_deepseek2_dsa_indexer accumulates the
per-head scores with ggml_add_inplace into an accumulator that is seeded
from a view of KQ_mask. On -fa 0, KQ_mask is the raw F32 input tensor, so
the in-place writes land in the shared KQ_mask buffer and corrupt the causal
mask that build_deepseek2_dsa_sparse_mask and the later softmax layers read
back, which gives garbage perplexity.
-fa 1 is unaffected (its F16 mask is cast to a private F32 buffer), and the
small-batch path added in #2067 is unaffected (it uses a non-inplace add).
Take a private copy of the seed in the batch >8 -fa 0 path (raw F32 mask)
before the accumulation, matching what those two paths already do.
4K -fa 0 --dsa PPL goes from thousands to 2.7134 (dense 2.6972, -fa 1 --dsa
2.7111). -fa 1 and non-DSA builds are byte-identical.
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* 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>
DFlashDraftModel.modify_tensors rewrites flat-named norm.weight and layers.N.* to their model.* form but not embed_tokens.weight, so a draft that carries its own (self-contained) token embeddings with flat naming fails with 'Can not map tensor embed_tokens.weight'. Handle it the same way as norm.weight so self-contained drafts convert.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
The openai gpt-oss models, and the z-lab gpt-oss DFlash drafts that reuse their
tokenizer, ship the o200k_harmony BPE tokenizer, which shares the GPT-4o
pre-tokenizer regex already handled by LLAMA_VOCAB_PRE_TYPE_GPT4O. Its checksum
was unrecognized, so get_vocab_base_pre() fell through to the "BPE pre-tokenizer
was not recognized" warning and conversion failed.
Register the gpt-oss-20b tokenizer hash against the existing "gpt-4o" pre-type
in convert_hf_to_gguf_update.py and add the matching entry in
get_vocab_base_pre(), following the dual-location pattern already used for the
z-lab Qwen3.5 DFlash drafts.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
Adds support for the z-lab gpt-oss DFlash drafts (e.g. z-lab/gpt-oss-20b-DFlash),
which use a Qwen3 backbone trained with attention_bias=true.
- dflash-draft loader: create optional attention bias tensors bq/bk/bv/bo
(TENSOR_NOT_REQUIRED); the Qwen3.5/MiMo drafts have none and are unaffected.
- build_dflash: add those biases at the q/k/v/o projections (cross-context K/V
and the noise block), each guarded so bias-free drafts are unchanged.
- Narrow the DFlash graph contract validator to reject only fused qkv biases
(bqkv/bqk/bkv), which the graph still does not implement, and remove a dead
duplicate of the validator in llama.cpp (the live copy is in llama-dflash.cpp).
The gpt-oss tokenizer recognition needed to convert the draft is a separate,
general gpt-oss conversion fix submitted independently.
Verified on gpt-oss-20b: coherent output and 42% draft acceptance
(~2.95 accepted tokens/cycle) at cross_ctx=128.
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* Split mode graph for GLM4MoE MTP in split_mode_tensor_parallel
Analogous to what was done for Qwen35 dense MTP in PR 2027:
- ctx_for_layer_split(): return buft_matrix for GLM4_MOE MTP tail
layers instead of buft, so split tensor preparation uses the
correct split context.
- create_glm4_moe_tensors(): remove the ctx_split = ctx_layer
override for MTP tail layers. ctx_for_layer_split(i) now
returns buft_matrix for GLM4_MOE MTP tails, so regular layer
tensors (attn, ffn) are created in the split context. NextN
tensors (eh_proj, enorm, hnorm, shared_head_head,
shared_head_norm) stay monolithic via ctx_for_layer().
- create_tensors(): add LLM_ARCH_GLM4_MOE to the MTP tail
layer splitting exclusion so split processing visits them.
- build_glm4_moe_mtp(): pass inp_out_ids to build_std_attention
instead of post-processing with ggml_get_rows. Use build_output
for the output projection to properly handle split mode.
- build_output(): add LLM_ARCH_GLM4_MOE to the is_qwen_mtp
check to ensure MTP output is properly materialized.
* Skip loading shared_head_head for GLM4MoE MTP
Add TENSOR_SKIP flag to shared_head_head so it is never loaded, even
when present in the GGUF file. The graph code already falls back to
model.output when shared_head_head is nullptr (line 375-376), which
frees ~306 MiB on CUDA0 for models that include this tensor (e.g.,
GLM-Steam-106B). The 355B GLM-4.6 model does not have this tensor and
already uses the same fallback path.
* cuda-graph: GLM4_MOE - Don't load layer.nextn.embed_tokens
---------
Co-authored-by: Nexesenex <124105151+Nexesenex@users.noreply.github.com>
* Support MiMo DFlash draft conversion
* Fix MiMo2 DFlash capture row pruning
* Fix MiMo DFlash draft RoPE and value scale
* Honor partial_rotary_factor in DFlash draft RoPE dim count
The draft set rope.dimension_count to the full head_dim (128), ignoring the
MiMo DFlash draft's partial_rotary_factor=0.5. The correct count is
head_dim*partial_rotary_factor=64; the remaining dims are NoPE. With the full
head_dim the upper half of each head receives position rotation it was never
trained for, which roughly halves draft acceptance on code (~26% -> ~60% once
corrected). RoPE base (5e6) and value scale (0.612) were already correct.
* Filter weight-map shard discovery to files that exist
get_model_part_names_from_weight_map() returned shard names straight from the
index weight_map without checking they exist. A model dir with a stale
model.safetensors.index.json but no safetensors shards would then set
is_safetensors=True and skip the pytorch_model*.bin fallback, failing later when
opening the missing files. Filter to shards present on disk so a stale index
falls through to the other weight formats.
* DFlash: store backbone_rotary_base in dedicated GGUF key
backbone_rotary_base (the target model's RoPE theta used when encoding
context K/V) was written to rope.freq_base, clobbering the draft
model's own rope_theta. For MiMo this swapped 10000 → 5000000 in the
draft attention path.
Fix: write backbone_rotary_base to a dedicated dflash.backbone_rotary_base
GGUF key and read it into hparams.dflash_backbone_rotary_base. In
build_dflash_kv_cache, use target_freq_base (the new hparam when set,
falling back to freq_base) for the context-K RoPE call. The draft model's
own rope.freq_base is now set correctly from rope_theta.
Existing MiMo DFlash GGUFs must be reconverted.
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
examples/llava has been replaced by mtmd since late 2025, and has
been out-of-build in ik_llama.cpp since examples/CMakeLists.txt removed it
in #798.
Repointed descriptions from llava to mtmd where they remained.
Loading a model with no GPU layers on a binary built with CUDA crashes in
`llm_load_tensors`. The GPU-fit block is guarded by `if (device_count > 0)`, but
`device_count` comes from `model.splits`, which always has at least one entry
(`{1.0f}`). The memory array it indexes, `device_mem`, is sized by `model.devices`,
which is empty when no GPU is present or when the model is loaded with `-ngl 0`. So the
block runs with `device_count >= 1` and reads `device_mem[0]` out of bounds.
Repro: build with `-DGGML_CUDA=ON` on a host that has no usable GPU, or hide the GPUs
with `CUDA_VISIBLE_DEVICES=""`, then load any model. The load segfaults inside the fit
loop (confirmed with DeepSeek-V2-Lite-Q4_K_M). With a real GPU present `model.devices`
is non-empty even at `-ngl 0`, so the crash needs the empty-device case.
The fix is to also require `!model.devices.empty()` before entering the GPU-fit block.
CPU-only placement is already handled earlier, all layers go to the CPU when there are
no GPU layers, so skipping this block on a CPU-only load is correct.
GPU loads still take the block since `model.devices` is non-empty. CPU-only loads on a
CUDA build now finish and decode normally instead of crashing.
Co-authored-by: local-llm <local-llm@local-llm-R740.cruvis.org>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
The CPU `set_rows` kernel for F32 sources fetches `type_traits[dst->type].from_float`
and calls it for every scattered row. F32 has no `from_float` entry, it is NULL in
`type_traits`, so any `set_rows` into an F32 destination calls a NULL function pointer
and segfaults. Other destination types work because they all have a real `from_float`.
Repro (CPU backend, standalone ggml graph):
dst = new_tensor_2d(F32, 8, 6) // F32 destination
src = new_tensor_2d(F32, 8, 4)
idx = new_tensor_1d(I64, 4) // {0,2,4,5}
out = ggml_set_rows(dst, src, idx)
// ggml_backend_graph_compute(cpu, ...) -> SIGSEGV on current main
When the destination is F32, copy the row with `memcpy` instead of going through
`from_float`. The I32 and I64 index branches both get the same treatment. An assert
guards the remaining case, non-F32 dst with a NULL `from_float`, so a future
unsupported type fails loudly instead of crashing.
I ran a normal model after this and it still decodes fine (DeepSeek-V2-Lite-Q4_K_M,
CPU, coherent output), and the non-F32 path is untouched. On the F32 path you pay one
`memcpy` per row in place of the indirect call.
Co-authored-by: local-llm <local-llm@local-llm-R740.cruvis.org>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* WIP: Split mode graph for Gemma4 assistant
Something is not right - acceptance drops to nearly zero.
* Per model CUDA contexts
Still not working!?
* This works
The issue was that I was not correctly calculating the number
of KV heads for the split KV cache.
* Compiler warnings
* It is better to use llama_context pointers as keys
* Split mode graph for dense Qwen35 MTP
Data server_tokens.map_idx_to_media.tokens_image.batch_f32 is read exactly once,
by mtmd_encode, however it was retained as long as the input image was present
in the sequence. Add a manual free function to clear out this data after encoding.
Solves:
* Memory wasted in struct server_tokens
* The same wasted memory in the ram cache
* Long copy durations cloning this data to/from ram cache
* Accounting failures in ram cache (`batch_f32` can be larger than a sequence's entire KV)
* The above accounting failures leading to terminal memory leaks in pathological cases
* Remove JSON serialization for `batch_32` which was unused, and had no foreseeable usecase
* DFlash: bound intra-block draft tokens to the SWA window
The SWA mask builder applied the sliding-window distance check only to the
cross-context section; the intra-block draft-token loop masked causal-only,
so a draft token could attend to earlier block tokens beyond n_swa. Apply
the same window bound ((j - block_k) < swa_window) in both the F16 and F32
paths so it matches the cross-context section.
Behavior-neutral for dense models: the SWA mask tensor is only allocated
when the model has SWA layers (build_dflash.cpp needs_swa_mask gate), so
for dense targets the changed block is unreachable.
* DFlash: enable sliding-window attention for draft models
DFlash drafts can be trained with sliding-window attention for long context,
but the runtime ignored it: the draft loader never read the window keys and the
converter never emitted them, so SWA-trained drafts always ran full-attention.
Enable it end to end and fix the dormant SWA graph path it exposes:
- convert_hf_to_gguf.py (DFlashDraftModel): emit attention.sliding_window + an
all-layers sliding_window_pattern when the source config sets use_sliding_window.
- llama-hparams.cpp (LLM_ARCH_DFLASH_DRAFT): read sliding_window + pattern into
n_swa / swa_layers.
- build_dflash.cpp + llama-dflash.cpp: the SWA mask path had never run; an all-SWA
draft turned the full kq_mask into a dead graph node the scheduler never backs
with a buffer, then the input-set wrote it unconditionally (GGML_ASSERT buf!=NULL).
Create + set each mask only when a layer uses it; derive mask dims from whichever
mask is live. Dense/mixed drafts are byte-identical.
Validated on gemma-4-26B-A4B at long context (cross_ctx 8176 > window 2048): no
crash, no short-context regression, SWA-on recovers long-context draft acceptance.
* DFlash: derive draft SWA pattern from layer_types
The converter emitted an all-layers SWA pattern ([True]*n_layers). The z-lab
DFlash drafts are sliding-window on every layer except a final full-attention
(global) layer, so this ran that global layer as sliding-window and clipped its
long-context view. Read layer_types and emit the matching per-layer pattern
(sliding_attention -> True), falling back to all-SWA only when layer_types is
absent.
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
The flash-attn vec kernels walk the KV cache in blocks of Dk rows for the
score loop but accumulate V in blocks of Dv. With Dk == Dv that is the same
thing, so normal attention shapes are fine. For absorbed MLA shapes where the
K and V head sizes differ (Dk=576/Dv=512 and Dk=192/Dv=128) the two loops step
a different number of KV rows, so K and V drift out of sync after the first
block and the V pointer reads the wrong cache rows.
This only shows up at decode (batch=1) on cards that fall back to the vec
kernel for MLA, which on NVIDIA is pre-Volta. There deepseek2/GLM MLA models
with -mla 1 -fa 1 or -mla 3 -fa 1 decode coherently for short prompts but
collapse into garbage once n_kv passes the first KV block (Dk=576). Prefill/PPL
is unaffected because prefill takes the tile kernel, not the vec kernel.
Fix: the score loop already covers Dk KV rows, so the V loop and the V pointer
step Dk rows too. For asymmetric Dk>Dv the V row is only Dv wide, so threads
with tid >= Dv have no V element (their VKQ lane is discarded at the output
store anyway) and read 0 instead of stepping past the row.
The change keys off the compile-time Dk != Dv, so every symmetric instantiation
compiles to byte-identical code and modern GPUs (which never take this vec path
for MLA) are unaffected.
Validated on a Tesla P100 (sm_60) with DeepSeek-V2-Lite Q4_K_M: decode coherence
restored for -mla 1/3 -fa 1, KLD vs the -fa 0 soft_max path drops from 4.79 to
1.4e-4 (same top token 27% -> 100%) at c1024, and TG is unchanged (82.8 t/s).
Co-authored-by: mb8565 <244351746+mb8565@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* WIP: Split mode graph for Gemma4 assistant
Something is not right - acceptance drops to nearly zero.
* Per model CUDA contexts
Still not working!?
* This works
The issue was that I was not correctly calculating the number
of KV heads for the split KV cache.
* Compiler warnings
* It is better to use llama_context pointers as keys
* fix: wrong stride in batched quantized add1 (nb0 -> nb3)
ggml_compute_forward_add1_q_f32 used i3*nb0 (element stride) instead of
i3*nb3 (batch stride) for the destination row pointer. This causes all
add1 operations with quantized types and batch > 1 to write to wrong
memory locations. The src0 pointer on the line above correctly uses nb03.
* fix: wrong dimension limits in dup_f16 non-contiguous path
The destination index wrapping in ggml_compute_forward_dup_f16 used
source dimensions (ne00/ne01/ne02/ne03) instead of destination dimensions
(ne0/ne1/ne2/ne3). While source and destination shapes are currently
identical for dup, using the wrong variables is incorrect by design.
* fix: wrong dimension limits in dup_bf16 non-contiguous path
Same fix as the dup_f16 path: destination index wrapping used source
dimensions (ne00/ne01/ne02/ne03) instead of destination dimensions
(ne0/ne1/ne2/ne3). Copy-paste error from the contiguous path.
* fix: ACC work size uses src[1] instead of src[0]
The dequantization work buffer for quantized ACC was sized using
src[1]->ne[0] instead of src[0]->ne[0]. Since src[0] is the tensor
being dequantized, its dimensions should determine the buffer size.
* fix: missing work size for SOFT_CAP_MAX and ROPE_BACK
Both ops dereference params->wdata in their forward functions but had
no work size allocation (cur = 0), causing a NULL pointer dereference
when any thread attempted to use wdata.
* fix: wrong dim in sum_rows_f32 dimension decomposition
Line 14404 used ne01*ne0 (= ne01*1) instead of ne01*ne02 for the
i3 term in the flat row index formula. When ne02 > 1 (batched 2D
inputs), this causes wrong memory access and corrupted results.
* Handle Cohere2MoE unopened thinking before tools
* Cohere2MoE: route unopened thinking to reasoning_content; test in active target
Follow-up to #1968. Gate extract_reasoning on reasoning_format only (drop the
"&& enable_thinking" addition) so the unopened-thinking handling does not also
change where an opened thinking block is routed. Under --reasoning off
(enable_thinking=false, reasoning_format defaults to DEEPSEEK) an orphaned
thinking block is now quarantined in reasoning_content with clean content and a
native tool call, instead of leaking the thinking prose into the user-facing
answer.
Move the Cohere2MoE end-to-end parser cases into tests/test-chat-auto-parser.cpp,
which CMake actually builds. tests/test-chat.cpp has been disabled in
tests/CMakeLists.txt since #723, so cohere coverage added there never ran in CI;
revert the local band-aids to that file.
* Cohere2MoE: harden parser from NMC eval findings
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
`{% set %}` of a non-loop variable inside a `{% for %}` body leaks
across iterations when the assignment is conditionally skipped. Each
iteration should start with a clean scope, matching standard Jinja2
semantics.
This fixes the issue with GLM-5.2 chat template when:
* turn 1 is a tool call with reasoning
* turn 2 is a tool call without reasoning
In this case, the reasoning content for turn 1 would be wrongly
duplicated to turn 2, resulting in degraded model performance.
ggml_print_tensor() accumulated `sum` inside the truncated print loop, which
skips the middle of each row when ne[0] > 2*n (n = 3). The printed `sum =`
therefore only covered the first n and last n elements per row, not the whole
tensor. For example a {2688, 5} tensor reported the sum of 30 of its 13440
values. That makes the value useless for numerically comparing two runs, and it
disagrees with mainline llama.cpp, whose eval-callback sums every element in a
separate pass.
This factors the per-element read into a small helper and computes the sum in
its own loop over all elements (double accumulator). The truncated print is a
separate, unchanged pass, so only the printed `sum =` value changes. The change
is confined to examples/eval-callback.
Co-authored-by: mb8565 <244351746+mb8565@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* fix: wrong tensor index in BF16 fused RMS norm add path (norm.cu:1039)
The BF16 branch of ggml_cuda_op_fused_rms_rms_add used dst->src[2]->data
for the second weight pointer, but should have used dst->src[3]->data.
This caused reading float weights from the wrong bf16 input tensor.
The F32 and F16 branches both correctly reference src[3], and the
assertions at lines 1013-1015 confirm src[3] is the F32 weight tensor.
* fix: off-by-one bounds check in 7 dmmv kernels (row > nrows -> row >= nrows)
Seven K-quant dequantize_mul_mat_vec kernels used row > nrows for bounds
checking instead of row >= nrows. Since rows are 0-indexed (0..nrows-1),
the check missed the row == nrows case, allowing a potential out-of-bounds
memory write when grid dimensions produce exactly nrows.
The templated dequantize_mul_mat_vec<type> kernel at line 667 already used
the correct row >= nrows pattern.
* fix: typo in function name iqk_mul_mat_vec_q_kerne -> iqk_mul_mat_vec_q_kernel
Truncated function name in iqk_mmvq_templates.cuh was missing trailing 'l'.
* fix: print actual split_dim value in set_tensor error message (ggml-cuda.cu)
fprintf used extra->split_dim == 0 which evaluates to boolean 0 or 1
instead of the actual split dimension value. When this fatal error is
hit for an unsupported split_dim, the user could not diagnose which
value caused the problem.
* fix: wrong src index in gate bias stride for fused up-gate MoE path
ggml_cuda_add_id for the gate bias used dst->src[4]->nb[1] as the stride
argument instead of dst->src[5]->nb[1]. This was a copy-paste error from
the up-bias code (lines 3220-3224) where src[4] is correct. If src[4]
and src[5] have different strides, the bias addition produces incorrect
results.
* fix: wrong row count for gate projection MMQ in fused up-gate MoE path
ggml_cuda_op_mul_mat_q for the gate projection (src0_2) used
src0_1->ne[1] as row_high instead of src0_2->ne[1]. This copy-paste
error causes processing the wrong number of rows if the up and gate
projections have different row counts. The gemv path (line ~3563)
correctly used src0_2->ne[1].
* CUDA : typo
* CUDA: Add missing GGML_CALL to function definition
* CUDA: only log GGML_CUDA_FORCE_MMQ/CUBLAS when enabled
* CUDA: Fix softcap bug in flash_attn_tile_ext_f16
The else branch (softcap != 0) incorrectly called launch_fattn_tile_f16_64_128
with use_softcap=false instead of true, causing logit softcap to be silently
ignored for the col_per_block=32, parallel_blocks=1 path.
