Repository navigation
Conversation
…lots (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 738022b; likely what cb6024e (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 (03da020). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01F6cBr5HtsKMSNV8KsK8FTd
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
Fixes an intermittent race condition in "fattn_gemm_softmax" that can cause NaNs with f16/q8_0 KV caches or inconsistent results with fp32 precision.
The kernel reuses shared memory for its sum reduction before all warps have finished reading the previous maximum. Adding a single "__syncthreads()" prevents the race.
Additional information
This addresses a separate race from the earlier "cb6024e6b" fix.
Requirements