ggml-cuda: fix divergent barrier in f16 flash attention by siavashnorouzi · Pull Request #27870 · ggml-org/llama.cpp

siavashnorouzi · 2026-08-28T09:48:05Z

Overview

The issue #27678 reported a divergent block-wide barrier in the flash_attn_ext_f16_process_tile kernel revolving around a block-wide __syncthreads() in the metadata-combine process.

I moved the block-wide synchronization out of the threadIdx.y % np == 0 branch so all threads in the np > 1 case reach the same __syncthreads(), while keeping the metadata calculation and SHMEM writeback limited to the intended warps.

Additional information

The reporter shared that compute-sanitizer reveals this issue:

compute-sanitizer --tool synccheck --show-backtrace device \
    build/bin/test-backend-ops -b CUDA0 -o FLASH_ATTN_EXT \
    -p 'hsk=192,hsv=128,nh=4,nr23=.8,1.,kv=512,nb=3,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16'

which reports thousands of divergent-barrier errors.

The issue was initially reported on RTX 4070, CUDA toolkit 13.3.73, driver 610.57.04. I was able to reproduce, fix, and test the fix on RTX 5090, CUDA 13.3, driver 610.43.02.

compute-sanitizer reports 3232 errors before and 0 after this fix.

Also, measured performance before and after this commit by running this config:

Model: Qwen3.8-27B-UD-Q4_K_XL.gguf
FlashAttention: Enabled
Context depth: 100k

and the results are:

Test	Before	After	Change
pp512 @ d100000 (t/s)	1390.27 ± 21.66	1388.10 ± 21.43	-0.15%
tg128 @ d100000 (t/s)	59.75 ± 0.06	59.73 ± 0.05	-0.03%
pp1024+tg256 @ d100000 (t/s)	255.36 ± 0.21	255.30 ± 0.19	-0.02%

Therefore, No measurable performance regression.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES, GPT-5.6 Sol was used for assistance with investigation and fix.

siavashnorouzi · 2026-09-02T05:01:57Z

The issue is still reproducible on the latest master branch (commit: b96806d) with recent changes in the same kernel and gives thousands of synccheck errors in compute-sanitizer with the same test parameters. Cherry-picking this fix applies cleanly and gives zero errors.

siavashnorouzi · 2026-09-02T21:53:06Z

@JohannesGaessler would you be able to take a look at his PR? It fixes the FlashAttention issue described above.

JohannesGaessler · 2026-09-04T08:26:30Z

+            const int jc_meta = threadIdx.y*cols_per_warp + (np*cols_per_warp < warp_size ? threadIdx.x % (np*cols_per_warp) : threadIdx.x);
+            const float2 * const meta_ptr = ((float2 *) tile_Q) + jc_meta*(tile_stride/2) + nbatch_combine/2;


Move these outside the conditional statement in order to have only one place where they are defined.

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

siavashnorouzi · 2026-09-11T01:57:49Z

Thank you @JohannesGaessler for taking the time to review this and @ggerganov for merging it. I really appreciate it and am looking forward to contributing more!

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

Port of upstream ggml-org/llama.cpp ggml-org#27870 (b74f590, PR by Siavash Norouzi), adapted to keep the beellama 0xFFFFFFFFULL shuffle masks and the KVarN dst_final_meta whole-tile publishing block. The old shape (kept since ggml-org#13469) put the meta-combine __syncthreads() inside 'if (np > 1 && threadIdx.y % np == 0)' with a separate 'else if (np > 1) __syncthreads()'. The paired barriers are formally divergent per the CUDA spec, and in practice the if-branch's post-barrier meta write-back races with the other warps re-entering the tile_Q reuse loop, which corrupts attention output data- and timing-dependently. This is the root cause of the catastrophic per-chunk KLD/PPL blowups seen on K2-Horizon-7B (GQA ratio 4, head dim 128 -> np = 4) with f16/bf16 KV caches and any exact tail: f16/f16, kvarn4 (intrinsic f16 tail) and bf16-tail runs collapse on whole chunks (chunk-9 PPL 5.11 -> 10.68 on Q5_K_M|f16) while quantized-body runs without a tail stay clean because they dispatch to fattn-vec. CPU runs of the same chunks/configs are clean, isolating the fault to this CUDA kernel. Restructure per upstream: all threads enter 'if (np > 1)', the combine reads run under 'threadIdx.y % np == 0', the __syncthreads() is unconditional, and the write-back runs after it under the same guard.

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

…barrier (upstream b74f590) process_tile's combine for np > 1 had one __syncthreads() inside threadIdx.y % np == 0 and another in the else branch: a barrier the warps reached at different instructions. Every warp now reaches one barrier, and the combine and the write-back run in the np == 0 warps around it. Ported by hand from ggml-org/llama.cpp b74f590 (ggml-org#27870). No arithmetic changes. Measured against c986f0c's build: - compute-sanitizer synccheck on the RTX 5070 Ti: upstream's repro (hsk=192 hsv=128 nh=4 [8,1] kv=512 nb=3) 3680 errors before, 0 after; GLM-5.3's cases (512/512 [64,1] 8192, sparse, nb 1 and 3) 512 before, 0 after. - test-backend-ops FLASH_ATTN_EXT 3229/3229 on both cards. - PPL bit for bit: 386825.5220 (-c 256 -ub 1), 388534.0622 (-ub 3), 387892.2627 (-sm tensor), 349039.3239 (-c 8192 -sm tensor). - Kernel time, ncu at base clocks with L2 flushed (gpu__time_duration of flash_attn_ext_f16 over GLM's grid, 3 cache lengths x 8 batches, sparse and dense, 48 cells): new/base within 1.79 %; the instrument's base/base spread 1.44 %. Co-authored-by: Siavash Norouzi <35790025+siavashnorouzi@users.noreply.github.com>

@64k

Hypothesis: every prefill FA launch for head_dim 256 was the scalar flash_attn_tile<256,256>, because fattn.cu gated the AMD WMMA (MMA_F16) path on D <= 128. FA was 59% of prefill at 32K depth and ~85% at 128K. Change (RDNA4 under HIP only; every other arch keeps the old expressions): - Port of upstream ggml-org/llama.cpp #28102 (gfx1201 FA tuning): WMMA for D=256 once Q rows * gqa_ratio_eff > 16; ncols2 chosen to divide the GQA ratio at D=256 (GQA 6 -> 2); RDNA D=256 config rows; stream-k only when tile efficiency < 75% (D=64 always). - Port of upstream #27870: divergent __syncthreads in the np>1 combine. - Ours: D=256 ncols=64 config stages Q in LDS with 32-half2 K/V chunks (Q_in_reg=false): VGPR spills 424 -> 48, FA kernel 5.44 -> 4.33 ms at 16K KV, 32.5 -> 27.7 ms at 64K. - test-backend-ops: model-shape cases (D=256, 4 KV heads, GQA 6; KV 512/4096/8192; nb 1..512 and 1024/1030 for -ub 1024; q8_0 KV). Benchmark (RX 9070 XT, Bonsai 2 27B PQ2_0, llama-bench -r 2): Before (deddbab): pp2048 @ 32K depth 518 t/s; pp512 1168; tg128 49.8 After: pp2048 @ 32K depth 906 t/s (+75%); pp512, tg128 same Earlier run (upstream config): @64k q8_0 340 -> 704, @128k 199 -> 498. Agent edit workload (ngram-mod + MTP draft 3, ~33-97-token verify batches, which now take the WMMA kernel): 307.2 -> 316.8 t/s, same acceptance. VRAM +32 MiB. Correctness: test-backend-ops FLASH_ATTN_EXT 2952/2952 (15/15 at the model shape). PPL (wikitext-2, 20 chunks) 7.8773 -> 7.8695 (-0.10%), identical at -ub 512 and -ub 1024; KLD vs previous 1.45e-4, same as the model's own ub1-vs-ub512 noise (1.57e-4): fp reduction order in attention. Decode (1 token) and MTP verify (<= 4 tokens) stay on the old kernels. Decision: KEEP Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

siavashnorouzi requested a review from a team as a code owner August 28, 2026 09:48

github-actions Bot added ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Aug 28, 2026

Midaychi mentioned this pull request Aug 28, 2026

CUDA illegal memory access in cudaStreamSynchronize (flash-attn path) with Qwen3.6-35B MoE + partial expert offload — deterministic, cross-build (b10107, b10243); disappears with -fa off #26609

Open

siavashnorouzi mentioned this pull request Sep 2, 2026

Misc. bug: CUDA FlashAttention synccheck reports divergent __syncthreads #27678

Open

JohannesGaessler reviewed Sep 4, 2026

View reviewed changes

JohannesGaessler self-assigned this Sep 4, 2026

siavashnorouzi added 2 commits September 4, 2026 18:38

ggml-cuda: fix divergent barrier in f16 flash attention

672f59d

ggml-cuda: avoid duplicate metadata pointer setup

c4b9b7e

siavashnorouzi force-pushed the fix-27678-fattn-barrier branch from ddc7fa9 to c4b9b7e Compare September 5, 2026 01:48

JohannesGaessler approved these changes Sep 5, 2026

View reviewed changes

JohannesGaessler added the merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. label Sep 5, 2026

ggerganov merged commit b74f590 into ggml-org:master Sep 7, 2026
19 of 22 checks passed

leehack mentioned this pull request Sep 7, 2026

Qualify post-v0.4.0 upstream grammar and GDN correctness fixes leehack/llamadart-native#76

Closed

x1250 pushed a commit to x1250/llama.cpp that referenced this pull request Sep 9, 2026

ggml-cuda: fix divergent barrier in f16 flash attention (ggml-org#27870)

894807f

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

zbrad pushed a commit to zbrad/llama.cpp that referenced this pull request Sep 10, 2026

ggml-cuda: fix divergent barrier in f16 flash attention (ggml-org#27870)

eabaac1

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

ggerganov mentioned this pull request Sep 14, 2026

ggml : bump version to 0.24.0 ggml-org/ggml#1627

Merged

BrewTestBot mentioned this pull request Sep 14, 2026

ggml 0.24.0 Homebrew/homebrew-core#310890

Merged

1 task done

harveyff mentioned this pull request Sep 15, 2026

[UPDATE][llamacppllmbasev3][1.2.38] Update CUDA engine to b10969 beclab/apps#3544

Merged

1 task

pl752 pushed a commit to pl752/llama.cpp that referenced this pull request Sep 15, 2026

ggml-cuda: fix divergent barrier in f16 flash attention (ggml-org#27870)

040626b

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

zsogitbe pushed a commit to zsogitbe/llama.cpp that referenced this pull request Sep 17, 2026

ggml-cuda: fix divergent barrier in f16 flash attention (ggml-org#27870)

254824b

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

Te-eMster pushed a commit to Te-eMster/mx-llama.cpp that referenced this pull request Sep 18, 2026

ggml-cuda: fix divergent barrier in f16 flash attention (ggml-org#27870)

f060a72

* ggml-cuda: fix divergent barrier in f16 flash attention * ggml-cuda: avoid duplicate metadata pointer setup

JohannesGaessler mentioned this pull request Sep 21, 2026

CUDA : fix divergent FlashAttention barrier #27955

Closed

marcospaulo mentioned this pull request Oct 1, 2026

perf: Qwen3.5 levers3; GLM-5.3 decode and MTP-verify levers (routed-expert ring, top-k at 2-4 rows, L2 issue stop, hyper-connection front) torad-labs/llama.cpp#84

Merged

Navigation Menu

ggml-cuda: fix divergent barrier in f16 flash attention - #27870

ggml-cuda: fix divergent barrier in f16 flash attention#27870
ggerganov merged 2 commits into
ggml-org:masterfrom
siavashnorouzi:fix-27678-fattn-barrier

siavashnorouzi commented Aug 28, 2026 •

edited

Loading

Uh oh!

siavashnorouzi commented Sep 2, 2026 •

edited

Loading

Uh oh!

siavashnorouzi commented Sep 2, 2026 •

edited

Loading

Uh oh!

JohannesGaessler Sep 4, 2026

Uh oh!

Uh oh!

siavashnorouzi commented Sep 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		const int jc_meta = threadIdx.ycols_per_warp + (npcols_per_warp < warp_size ? threadIdx.x % (np*cols_per_warp) : threadIdx.x);
		const float2 * const meta_ptr = ((float2 ) tile_Q) + jc_meta(tile_stride/2) + nbatch_combine/2;

Conversation

siavashnorouzi commented Aug 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Requirements

Uh oh!

siavashnorouzi commented Sep 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

siavashnorouzi commented Sep 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

JohannesGaessler Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

siavashnorouzi commented Sep 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

siavashnorouzi commented Aug 28, 2026 •

edited

Loading

siavashnorouzi commented Sep 2, 2026 •

edited

Loading

siavashnorouzi commented Sep 2, 2026 •

edited

Loading