From ef912575d5d15b328711d6764fdd96636ae207da Mon Sep 17 00:00:00 2001 From: nullraptor Date: Sun, 19 Jul 2026 02:22:53 +0000 Subject: [PATCH 1/6] sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention 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 #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/src/ggml-sycl/fattn-onednn.cpp | 47 +++++++++++++++++++++++------ tests/test-backend-ops.cpp | 9 ++++++ 2 files changed, 46 insertions(+), 10 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index f2e12ef1aeff..cf2627e1be0c 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -1,5 +1,6 @@ #include #include +#include #include #include #include @@ -38,6 +39,16 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) { return false; } + // Optional KV-length ceiling (GGML_SYCL_FA_ONEDNN_MAX_KV, 0 = unlimited). Escape hatch: + // very long sequences make the fused SDPA slow enough to risk the xe driver watchdog on + // some stacks; past the cap we fall back to the native FA kernel instead. + static const int64_t max_kv = [] { + const char * v = getenv("GGML_SYCL_FA_ONEDNN_MAX_KV"); + return v ? atoll(v) : (long long) 0; + }(); + if (max_kv > 0 && K->ne[1] > max_kv) { + return false; + } // gate for the following cases // 1. if the oneDNN graph Add node has no input --> skip // 2. types other than f16 need different logical_tensor declaration @@ -208,9 +219,27 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso cont_to_f16_sycl((const char *) V->data, Vf.get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream); // divide-by-(1/scale) reproduces ggml's score *= kq_scale on the proven probe graph. + // + // The scale upload must not be an async memcpy from a stack local: on the in-order queue + // the copy waits behind the K/V staging kernels, and once those take long enough + // (n_kv >= ~26k on B70) the frame is recycled before the copy runs, feeding the SDPA a + // garbage scale (output collapses to a repeated token). Cache one device scalar per + // (device, value) -- the scale is constant per model -- and upload it synchronously once. const sycl::half scale_h = (sycl::half) (1.0f / kq_scale); - ggml_sycl_pool_alloc scbuf(ctx.pool(), 1); - stream->memcpy(scbuf.get(), &scale_h, sizeof(sycl::half)); + static std::unordered_map scale_cache; + char sckey[64]; + snprintf(sckey, sizeof(sckey), "%d:%a", ggml_sycl_get_device(), (double) (float) scale_h); + sycl::half * scale_dev; + { + auto sit = scale_cache.find(sckey); + if (sit == scale_cache.end()) { + scale_dev = sycl::malloc_device(1, *stream); + stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)).wait(); + scale_cache.emplace(sckey, scale_dev); + } else { + scale_dev = sit->second; + } + } ggml_sycl_pool_alloc outf(ctx.pool(), (size_t) H * q * d); // f16 contiguous SDPA out [mb,H,q,d] @@ -232,7 +261,7 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso if (r == E.id_q) return Qf.get(); if (r == E.id_k) return Kf.get(); if (r == E.id_v) return Vf.get(); - if (r == E.id_scale) return scbuf.get(); + if (r == E.id_scale) return scale_dev; if (r == E.id_mask) return (void *) mask->data; return nullptr; }; @@ -245,14 +274,12 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso E.cp.execute(strm, ti, {to}); permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream); - // Single device: no sync is required, and actually PP perf is ~6% > wait_and_throw() (tested on llama-3.1-8b & qwen3.6-27b, both Q8_0, with Arc B70). - // Any future multi-GPU refactor MUST re-measure this single-device path and keep the best - // single-device PP speed. Otherwise (multiple devices/streams can race the reuse): + // Single device needs no sync: the dnnl stream wraps this same in-order queue, so the SDPA + // serializes with the staging kernels before it and the permute/pool reuse after it. The + // garbage output formerly blamed on the missing sync here was the scale use-after-return + // fixed above. Keep the conservative wait for multi-GPU, where other devices' streams can + // race the pool: if (ggml_sycl_info().device_count > 1) { - // cont_to_f16 -> oneDNN execute -> permute is async on this stream, but the - // pool_alloc*s above free their device buffers at host return. Without this wait the next - // scheduler op re-acquires those bytes while the GPU is still computing the SDPA, turning - // it into garbage and collapsing multi-turn output to a single repeated token ("GGGGG..."). stream->wait_and_throw(); } } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index e3553c5e9ca1..a600f149976f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9542,6 +9542,15 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_flash_attn_ext(64, 128, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q1_0)); test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 64, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_F16)); + // large-KV F16 cases (Qwen3.6-27B geometry and a llama-class control): the upstream matrix + // stops at kv=1024, blind to long-context FA bugs (e.g. the oneDNN SDPA ordering race on BMG). + for (int64_t kv : { 4096, 16384, 32768, 65536 }) { + test_cases.emplace_back(new test_flash_attn_ext(256, 256, 4, {6, 1}, kv, 512, true, false, 0, 0, + GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + test_cases.emplace_back(new test_flash_attn_ext(128, 128, 8, {4, 1}, kv, 512, true, false, 0, 0, + GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + } + test_cases.emplace_back(new test_cross_entropy_loss (GGML_TYPE_F32, { 10, 5, 4, 3})); test_cases.emplace_back(new test_cross_entropy_loss (GGML_TYPE_F32, {30000, 1, 1, 1})); test_cases.emplace_back(new test_cross_entropy_loss_back(GGML_TYPE_F32, { 10, 5, 4, 3})); From 79bfc18aff658317d54a807e53919a5cf500151a Mon Sep 17 00:00:00 2001 From: nullraptor Date: Mon, 20 Jul 2026 04:46:46 +0000 Subject: [PATCH 2/6] sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it Review feedback on #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 --- docs/backend/SYCL.md | 2 ++ ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/fattn-onednn.cpp | 6 +----- ggml/src/ggml-sycl/ggml-sycl.cpp | 3 +++ tests/test-backend-ops.cpp | 2 +- 5 files changed, 8 insertions(+), 6 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 0814ceb60f92..4292f8e6e0d9 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -794,6 +794,8 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. | | 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. | | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. | | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). | | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.
Recommended to use when --split-mode = layer | diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index e5d9ee89dd86..f27ec5dd6260 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -65,6 +65,7 @@ extern int g_ggml_sycl_prioritize_dmmv; extern int g_ggml_sycl_enable_flash_attention; extern int g_ggml_sycl_dev2dev_memcpy; extern int g_ggml_sycl_fa_onednn; +extern int g_ggml_sycl_fa_onednn_max_kv; #if defined(__clang__) && __has_builtin(__builtin_expect) diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index cf2627e1be0c..597d28b678b0 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -42,11 +42,7 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { // Optional KV-length ceiling (GGML_SYCL_FA_ONEDNN_MAX_KV, 0 = unlimited). Escape hatch: // very long sequences make the fused SDPA slow enough to risk the xe driver watchdog on // some stacks; past the cap we fall back to the native FA kernel instead. - static const int64_t max_kv = [] { - const char * v = getenv("GGML_SYCL_FA_ONEDNN_MAX_KV"); - return v ? atoll(v) : (long long) 0; - }(); - if (max_kv > 0 && K->ne[1] > max_kv) { + if (g_ggml_sycl_fa_onednn_max_kv > 0 && K->ne[1] > g_ggml_sycl_fa_onednn_max_kv) { return false; } // gate for the following cases diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index cb8974eedb75..3b807c7cbd54 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -85,6 +85,7 @@ int g_ggml_sycl_enable_optimize = 1; int g_ggml_sycl_enable_graph = 0; int g_ggml_sycl_enable_dnn = 1; int g_ggml_sycl_fa_onednn = 1; +int g_ggml_sycl_fa_onednn_max_kv = 0; int g_ggml_sycl_enable_vmm = 1; int g_ggml_sycl_enable_fusion = 1; int g_ggml_sycl_prioritize_dmmv = 0; @@ -287,6 +288,7 @@ static void ggml_check_sycl() try { g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0); g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1); g_ggml_sycl_fa_onednn = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN", 1); + g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0); g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1); g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1); g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0); @@ -359,6 +361,7 @@ static void ggml_check_sycl() try { GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: DNN disabled by compile flag\n"); GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN: %d\n", g_ggml_sycl_fa_onednn); #endif + GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN_MAX_KV: %d\n", g_ggml_sycl_fa_onednn_max_kv); #ifdef SYCL_FLASH_ATTN GGML_LOG_INFO(" GGML_SYCL_ENABLE_FLASH_ATTN: %d\n", g_ggml_sycl_enable_flash_attention); #else diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index a600f149976f..63b7ab6215d6 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9544,7 +9544,7 @@ static std::vector> make_test_cases_eval() { // large-KV F16 cases (Qwen3.6-27B geometry and a llama-class control): the upstream matrix // stops at kv=1024, blind to long-context FA bugs (e.g. the oneDNN SDPA ordering race on BMG). - for (int64_t kv : { 4096, 16384, 32768, 65536 }) { + for (int64_t kv : { 4096, 16384 }) { test_cases.emplace_back(new test_flash_attn_ext(256, 256, 4, {6, 1}, kv, 512, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); test_cases.emplace_back(new test_flash_attn_ext(128, 128, 8, {4, 1}, kv, 512, true, false, 0, 0, From 5158ff7d3e787611b41b03a65525bf11438e7516 Mon Sep 17 00:00:00 2001 From: nullraptor Date: Tue, 21 Jul 2026 04:51:21 +0000 Subject: [PATCH 3/6] sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled Assisted-by: Claude Fable 5 --- docs/backend/SYCL.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 4292f8e6e0d9..4526d498fdcc 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -795,7 +795,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | 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. | +| GGML_SYCL_FA_ONEDNN_MAX_KV | 0 (default, disabled) 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. | | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. | | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). | | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.
Recommended to use when --split-mode = layer | From 87e1588c6fd94534ef727c393eb3e5eddd51a34c Mon Sep 17 00:00:00 2001 From: nullraptor Date: Tue, 21 Jul 2026 05:38:35 +0000 Subject: [PATCH 4/6] sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly Assisted-by: Claude Fable 5 --- docs/backend/SYCL.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 4526d498fdcc..c72dc3b3ee02 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -795,7 +795,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | 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, disabled) 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. | +| GGML_SYCL_FA_ONEDNN_MAX_KV | 0 (default, disabled) or positive integer | By default (0), all sequences are handled by the oneDNN fused SDPA path, regardless of KV length; a positive value caps that length, past which sequences fall back to the native kernel. If GPU driver watchdog resets (DEVICE_LOST) occur during long-context inference, set this near the context depth where they start, e.g. 24576. | | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. | | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). | | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.
Recommended to use when --split-mode = layer | From 00b19ec2e1f3172f7de345247e400f1777849ebf Mon Sep 17 00:00:00 2001 From: meatposes Date: Tue, 21 Jul 2026 09:04:57 -0500 Subject: [PATCH 5/6] Update ggml/src/ggml-sycl/fattn-onednn.cpp Co-authored-by: Neo Zhang --- ggml/src/ggml-sycl/fattn-onednn.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index 597d28b678b0..18f4ab75b565 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -230,7 +230,7 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso auto sit = scale_cache.find(sckey); if (sit == scale_cache.end()) { scale_dev = sycl::malloc_device(1, *stream); - stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)).wait(); + stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)); scale_cache.emplace(sckey, scale_dev); } else { scale_dev = sit->second; From 6ecb74788d9a82a772f1c943cf00f18450effd4f Mon Sep 17 00:00:00 2001 From: nullraptor Date: Mon, 27 Jul 2026 18:40:18 +0000 Subject: [PATCH 6/6] sycl: write the SDPA scale from a kernel instead of caching it The per-(device, value) scale cache was a function-local static unordered_map with no synchronization, so concurrent backend instances could access and rehash it at the same time. Write the scalar with a single_task instead. The value is captured into the command, so no host memory has to outlive the call -- which is what the use-after-return fix needed in the first place. That removes the shared container, the leaked device allocation and the string key, and it also closes the remaining async-memcpy-from-a-stack-local on the first flash-attention call. Ordering does not rely on timing: the queue is created with sycl::property::queue::in_order and the dnnl stream wraps that same queue, so the write completes before the SDPA reads the scalar. The multi-GPU wait_and_throw() branch is unchanged. Also drop the include, which is unused. Assisted-by: Claude Opus 5 --- ggml/src/ggml-sycl/fattn-onednn.cpp | 29 +++++++++-------------------- 1 file changed, 9 insertions(+), 20 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index 18f4ab75b565..8465e12248f4 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -1,6 +1,5 @@ #include #include -#include #include #include #include @@ -216,26 +215,16 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso // divide-by-(1/scale) reproduces ggml's score *= kq_scale on the proven probe graph. // - // The scale upload must not be an async memcpy from a stack local: on the in-order queue - // the copy waits behind the K/V staging kernels, and once those take long enough - // (n_kv >= ~26k on B70) the frame is recycled before the copy runs, feeding the SDPA a - // garbage scale (output collapses to a repeated token). Cache one device scalar per - // (device, value) -- the scale is constant per model -- and upload it synchronously once. + // The scale must not be uploaded with an async memcpy from a stack local: on the in-order + // queue that copy waits behind the K/V staging kernels, and once those take long enough + // (n_kv >= ~26k on B70) the host frame is recycled before the copy runs, feeding the SDPA a + // garbage scale (output collapses to a repeated token). Write the scalar from a kernel + // instead -- the value is captured into the command, so no host memory has to outlive the + // call, and the enqueue stays async. const sycl::half scale_h = (sycl::half) (1.0f / kq_scale); - static std::unordered_map scale_cache; - char sckey[64]; - snprintf(sckey, sizeof(sckey), "%d:%a", ggml_sycl_get_device(), (double) (float) scale_h); - sycl::half * scale_dev; - { - auto sit = scale_cache.find(sckey); - if (sit == scale_cache.end()) { - scale_dev = sycl::malloc_device(1, *stream); - stream->memcpy(scale_dev, &scale_h, sizeof(sycl::half)); - scale_cache.emplace(sckey, scale_dev); - } else { - scale_dev = sit->second; - } - } + ggml_sycl_pool_alloc scbuf(ctx.pool(), 1); + sycl::half * const scale_dev = scbuf.get(); + stream->single_task([=]() { *scale_dev = scale_h; }); ggml_sycl_pool_alloc outf(ctx.pool(), (size_t) H * q * d); // f16 contiguous SDPA out [mb,H,q,d]