From 8d601afcfc2056d7e8411201fc4b87edc2155d88 Mon Sep 17 00:00:00 2001 From: Tyler Date: Thu, 8 Oct 2026 03:42:54 +0000 Subject: [PATCH] cuda: barrier before fattn_gemm_softmax reuses its shared reduction slots (f16/q8_0 KV long prompts: intermittent NaN -> none) fattn_gemm_softmax reads the block's row max back from __shared__ red[0] and then reuses red[] for the row-sum reduction with no barrier in between: warp 0 can store its partial sum into red[0] before a slower warp has read the max, and that warp then uses the sum as the max in exp(v - m). With GGML_CUDA_FA_GEMM_PREC=32 that is run-to-run drift; with the default fp16 path a probability can overflow half to inf and the attention output becomes NaN. Fix: one __syncthreads() after vmax = red[0]. Affected: the generic branch of ggml_cuda_flash_attn_ext_gemm (pre-Volta, nkv >= 4096, >= 128 query rows, KV type other than q4_0, e.g. f16/q8_0). q4_0 caches go to the fold kernels, which have barriers between every red write and read; decode/verify never enter this path. Present since 738022bda; likely what cb6024e6b (out-of-place S/P) was chasing. Repro (2x P100, -sm tensor, Qwen3.8-27B Q6_K, 13 prompts of 5.8k-7.9k tokens x3 passes, b/ub 512, 8k ctx): f16 KV: NaN 10/39 -> 0/39; prompts whose logits differ between identical passes 8/13 -> 0/13 q8_0 KV: NaN 8/39 -> 0/39; differing prompts -> 0/13 Before the fix: GGML_CUDA_FA_GEMM=0 (tile) 0/39 and deterministic; GGML_CUDA_FA_GEMM_PREC=32 0/39 but 11/13 nondeterministic -> a race in this kernel, not overflow on particular data. No cost to the default q4_0 path: fixed vs pre-fix libggml-cuda, ABAB at 150 W, pp2048@8k 408.96/406.82 vs 407.63/406.64, tg256@8k 26.36/26.37 vs 26.33/26.35. tools/gate.sh --full: PPL 2.6074 unchanged; one unrelated MUL_MAT q5_1 tolerance flake that passes 5/5 on both libraries. Found and fixed with Claude Code (AI-assisted); first landed on Keoian/dual-p100-qwen-3-8-27b tyler-port (03da0202b). Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01F6cBr5HtsKMSNV8KsK8FTd --- ggml/src/ggml-cuda/fattn-gemm.cu | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/ggml/src/ggml-cuda/fattn-gemm.cu b/ggml/src/ggml-cuda/fattn-gemm.cu index 1855d8a5b1a7..2f2a0d84a263 100755 --- a/ggml/src/ggml-cuda/fattn-gemm.cu +++ b/ggml/src/ggml-cuda/fattn-gemm.cu @@ -168,6 +168,11 @@ static __global__ void fattn_gemm_softmax( } __syncthreads(); vmax = red[0]; + // every thread must have read the max out of red[0] before pass 2 reuses red for the row sum: + // without this barrier warp 0 can store its partial sum into red[0] first, a late warp then + // takes that sum as the row max, and exp(v - m) is wrong for its keys -- run-to-run drift in + // fp32, and in fp16 a probability that overflows to inf and NaNs the output + __syncthreads(); } const float m_old = m_state[h*nt + t];