sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path by meatposes · Pull Request #25880 · ggml-org/llama.cpp

meatposes · 2026-07-19T02:36:08Z

Overview

Fixes the long-context garbage output in the SYCL oneDNN flash-attention path (introduced in
32b741c / #25222) at its root cause, restoring the single-device no-sync fast path.

The SDPA scale was uploaded with an async memcpy whose source is a stack local. On the in-order
queue that copy is ordered behind the K/V staging kernels; once n_kv is large enough (>= ~26k
observed on Arc Pro B70) the staging outlives the host stack frame, the copy reads recycled
stack memory, and the SDPA runs with a garbage scale. Output collapses to a single repeated
token and the poisoned KV corrupts the rest of the session. Short contexts win the race by
accident; test-backend-ops caps FLASH_ATTN_EXT at kv=1024 and syncs between ops, so CI cannot
see it.

Fix: upload the scale synchronously, once, into a device scalar cached per (device, value) --
the scale is constant per model. The single-device path then needs no per-call host sync
(every device-side hazard already serializes on the in-order queue); the conservative
multi-GPU wait is kept unchanged. Also adds an optional GGML_SYCL_FA_ONEDNN_MAX_KV n_kv
ceiling (default unlimited) and large-KV FLASH_ATTN_EXT test cases up to kv=65536.

Additional information

Related: #25741 fixes the same symptom with an unconditional wait_and_throw(); that works by
keeping the stack frame alive across the copy, but costs a host sync on every prefill FA call
and leaves the use-after-return in place. If #25741 lands first, this applies on top by
reverting the unconditional wait to the multi-GPU-only branch; the scale fix is independent.

Reproducer (unmodified build, GGML_SYCL_F16=ON + GGML_SYCL_DNN=ON, Arc Pro B70): send one
~28k-token prompt ending "Reply with exactly: HELLO WORLD" at temperature 0.
Before: GGML_SYCL_FA_ONEDNN=1 -> repeated-token garbage / HTTP 500; =0 -> "HELLO WORLD".
After: both -> "HELLO WORLD".

Testing on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072 -b 2048 -ub 2048 -fa on:

temp-0 output byte-identical to the native FA path through 32k-deep prefill
prefill depth-flat at 820-840 t/s at 8k-32k depth (native FA: ~600 at 8k, ~350 at 32k);
token generation unchanged at all depths
test-backend-ops -o FLASH_ATTN_EXT: all pre-existing cases pass. The new kv>=32k cases show
~6.5e-4..1.3e-3 NMSE vs the 5e-4 gate on the oneDNN path -- f16 accumulation drift inside
the fused kernel that predates and is unrelated to this fix.

The scale cache (one 2-byte device allocation per distinct scale value) is never freed,
matching the lifetime behavior of the existing sdpa partition cache in the same file.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES - the root-cause analysis, patch, and benchmark/testing harness were
developed with AI assistance (Claude, disclosed via the Assisted-by commit trailer). All
changes were reviewed by the submitter and validated end-to-end on real hardware (Arc Pro
B70) before submission.

…ntion path The scale was uploaded with an async memcpy sourced from a stack local. On the in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale. Output then collapses to a single repeated token and the KV cache is poisoned for the rest of the session. Short contexts win the race by accident, and test-backend-ops caps FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous device_count > 1 wait_and_throw() gate (and reverting it, PR ggml-org#25741) fixes the symptom only by keeping the frame alive across the copy at the cost of a host sync on every FA call. Fix: cache one device scalar per (device, value) -- the scale is constant per model -- and upload it synchronously once. The single-device fast path (no per-call host sync) is then safe: every device-side hazard already serializes on the in-order queue. The multi-GPU conservative wait is kept unchanged. Also: - GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that routes very long sequences to the native FA kernel. - test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind spot. Note the race itself needs a live multi-op pipeline to reproduce; single-op runs pass even on broken builds. Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output byte-identical at temp 0 to the native FA path through 32k-deep prefill, with prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth). Assisted-by: Claude Fable 5

ggml-gh-bot · 2026-07-19T02:40:14Z

Hi @meatposes, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

arthw · 2026-07-20T02:48:19Z

+    static const int64_t max_kv = [] {
+        const char * v = getenv("GGML_SYCL_FA_ONEDNN_MAX_KV");
+        return v ? atoll(v) : (long long) 0;
+    }();


Please following the existed style to handle the running env vars, like g_ggml_sycl_enable_optimize.
The benefit is easy to check the value by log (-lv 4).

Please update the explain of "GGML_SYCL_FA_ONEDNN_MAX_KV" in SYCL.md. Guide user how to set the value.

I did my best to follow your recommendations and moved the env var to follow same example as g_ggml_sycl_enable_optimize toggle

As requested I also added GGML_SYCL_FA_ONEDNN_MAX_KV to the runtime env table in SYCL.md with guidance on picking a value (0 = no cap; otherwise set it near the context depth where driver watchdog resets start appearing). While I was in there I went ahead and added GGML_SYCL_FA_ONEDNN as that toggle wasn't documented yet.

…rs and document it Review feedback on ggml-org#25880: - read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it) - document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md runtime table Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536 shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels (long-sequence fp16 accumulation drift, present before this PR) and would fail CI for an unrelated reason. Assisted-by: Claude Fable 5

arthw · 2026-07-21T02:13:50Z

 | GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. |
+| GGML_SYCL_FA_ONEDNN | 1 (default) or 0 | Enable the oneDNN fused SDPA (flash-attention) path on supported GPUs. Set to 0 to always use the native SYCL flash-attention kernel. |
+| GGML_SYCL_FA_ONEDNN_MAX_KV | 0 (default) or positive integer | Max KV length (in tokens) the oneDNN fused SDPA handles; longer sequences fall back to the native kernel. 0 = no limit. If GPU driver watchdog resets (DEVICE_LOST) occur during long-context inference, set this near the context depth where they start, e.g. 24576. |


How to disable this feature?

0 (default) means there is no max KV length, and GGML_SYCL_FA_ONEDNN_MAX_KV is thus disabled. This should only do something when a value is set that is not zero, so disabled by default.

ive updated the wording to reflect that it is 0 (default, disabled)

How to understand: "0 = no limit"?

replaced this clunky wording as well, thank you. Now reads that default (0) means oneDNN fused SPDA handles all sequences regardless of length

Assisted-by: Claude Fable 5

arthw · 2026-07-21T05:32:55Z

+        auto sit = scale_cache.find(sckey);
+        if (sit == scale_cache.end()) {
+            scale_dev = sycl::malloc_device<sycl::half>(1, *stream);
+            stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)).wait();


Suggested change

stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)).wait();

stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half));

With Qwen3.6-27B-UD-Q4_K_XL.gguf on B60

Test fa Base t/s Primary t/s Increase Rate (Primary vs Base)

pp512 0 53.38 52.16 -2.29%

pp512 1 49.21 48.02 -2.42%

tg128 0 3.57 3.51 -1.68%

tg128 1 3.59 3.53 -1.67%

after remove wait()

Test fa Base t/s Primary t/s Increase Rate (Primary vs Base)

pp512 0 52.10 53.42 2.53%

pp512 1 48.05 49.16 2.31%

tg128 0 3.48 3.58 2.87%

tg128 1 3.51 3.62 3.13%

i intend to test this on b70 (and maybe b50), I don't think its an issue, but i would want to verify

Assisted-by: Claude Fable 5

meatposes requested review from a team and ggerganov as code owners July 19, 2026 02:36

meatposes mentioned this pull request Jul 19, 2026

sycl: single-device still requires sync #25741

Open

github-actions Bot added testing Everything test related ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels Jul 19, 2026

arthw reviewed Jul 20, 2026

View reviewed changes

github-actions Bot added the documentation Improvements or additions to documentation label Jul 20, 2026

arthw reviewed Jul 21, 2026

View reviewed changes

sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled

5158ff7

Assisted-by: Claude Fable 5

arthw reviewed Jul 21, 2026

View reviewed changes

sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly

87e1588

Assisted-by: Claude Fable 5

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path#25880

sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path#25880
meatposes wants to merge 4 commits into
ggml-org:masterfrom
meatposes:sycl/fattn-onednn-scale-uaf

meatposes commented Jul 19, 2026 •

edited

Loading

Uh oh!

ggml-gh-bot Bot commented Jul 19, 2026

Uh oh!

arthw Jul 20, 2026

Uh oh!

meatposes Jul 20, 2026

Uh oh!

arthw Jul 21, 2026

Uh oh!

meatposes Jul 21, 2026

Uh oh!

meatposes Jul 21, 2026

Uh oh!

arthw Jul 21, 2026

Uh oh!

meatposes Jul 21, 2026 •

edited

Loading

Uh oh!

arthw Jul 21, 2026

Uh oh!

meatposes Jul 21, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

	stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)).wait();
	stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half));

Test	fa	Base t/s	Primary t/s	Increase Rate (Primary vs Base)
pp512	0	53.38	52.16	-2.29%
pp512	1	49.21	48.02	-2.42%
tg128	0	3.57	3.51	-1.68%
tg128	1	3.59	3.53	-1.67%

Test	fa	Base t/s	Primary t/s	Increase Rate (Primary vs Base)
pp512	0	52.10	53.42	2.53%
pp512	1	48.05	49.16	2.31%
tg128	0	3.48	3.58	2.87%
tg128	1	3.51	3.62	3.13%

Uh oh!

Conversation

meatposes commented Jul 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Requirements

Uh oh!

ggml-gh-bot Bot commented Jul 19, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

meatposes Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

meatposes commented Jul 19, 2026 •

edited

Loading

meatposes Jul 21, 2026 •

edited

Loading