* 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>
a8cf53fd69