mirror of
https://github.com/ikawrakow/ik_llama.cpp.git
synced 2026-07-21 10:15:35 +00:00
a8cf53fd69
* 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>