From b0b76ae577326e73108a177bb21f8efa2dd6d97e Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Sat, 18 Jul 2026 15:50:27 -0500 Subject: [PATCH 1/2] sycl: extend oneDNN SDPA to Q4_0-Q8_0 and F32 KV caches Extends the oneDNN SDPA path (PR #25222) to handle non-F16 KV caches by dequantizing or converting K/V to dense FP16 on-device before feeding them into the SDPA graph. The fused systolic kernel then runs identically to the native FP16 path. Supported KV types: - Q4_0, Q4_1, Q5_0, Q5_1, Q8_0: to_fp16_sycl / to_fp16_nc_sycl - F32: cont_to_f16_sycl - BF16 and IQ types are excluded (no conversion kernel available) Gate: non-F16 requires K >= 1024 and Q >= 32 (prefill only). F16 KV runs at any length (existing behavior). Also includes the stream sync fix (stream->wait_and_throw() unconditional, PR #25741 by @malsbat) and removal of V_is_K_view aliasing (K and V are always dequantized to separate buffers). Co-Authored-By: Claude --- docs/backend/SYCL.md | 1 + ggml/src/ggml-sycl/fattn-onednn.cpp | 149 ++++++++++++++++++++++++---- ggml/src/ggml-sycl/fattn.cpp | 30 +++++- 3 files changed, 154 insertions(+), 26 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 0814ceb60f92..157c46c011fe 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -796,6 +796,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. | | 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). | +| GGML_SYCL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for flash attention dispatch: kernel name, head_dim, n_kv, and n_kv delta. | | 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 | | UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS | 0 (default) or 1 | Allow SYCL/Unified Runtime Level Zero device allocations larger than 4 GiB. llama.cpp's direct Level Zero allocation path requests the relaxed maximum-size limit itself when GGML_SYCL_ENABLE_LEVEL_ZERO=1. | | GGML_SYCL_USM_SYSTEM | 0 (default) or 1 | Enable experimental support for [USM system allocations](https://github.khronos.org/SYCL_Reference/iface/usm_basic_concept.html#system-allocations) for large GPU buffers. This requires enough host memory for model weights and caches, an Intel Xe2+ GPU such as BMG or newer and supported on Linux only, with CONFIG_DRM_XE_GPUSVM enabled. | diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index f2e12ef1aeff..16ec05be1414 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -2,11 +2,13 @@ #include #include #include +#include #include #include #include "fattn-onednn.hpp" #include "fattn-tile.hpp" +#include "convert.hpp" // set minimum query length to treat as prefill (32) #define GGML_SYCL_FA_ONEDNN_MIN_Q 32 @@ -33,10 +35,30 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { const ggml_tensor * mask = dst->src[3]; const ggml_tensor * sinks = dst->src[4]; - // gate for f16 KV only for now - // need to implement quantized KV + // F16 KV: native SDPA at any KV length. + // Non-F16: dequant to F16 then SDPA at prefill lengths. Only the + // standard quantized KV cache types (Q4_0-Q8_0) and F32 are accepted + // because their to_fp16_sycl conversion is verified. BF16 and IQ* + // are excluded: BF16 needs a strided conversion kernel that does not + // exist yet; IQ types are model-weight-only quants with no dequant + // registration and are never used as KV caches. if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) { - return false; + auto kt = K->type, vt = V->type; + bool k_ok = kt == GGML_TYPE_F32 || kt == GGML_TYPE_Q4_0 || kt == GGML_TYPE_Q4_1 || + kt == GGML_TYPE_Q5_0 || kt == GGML_TYPE_Q5_1 || kt == GGML_TYPE_Q8_0; + bool v_ok = vt == GGML_TYPE_F32 || vt == GGML_TYPE_Q4_0 || vt == GGML_TYPE_Q4_1 || + vt == GGML_TYPE_Q5_0 || vt == GGML_TYPE_Q5_1 || vt == GGML_TYPE_Q8_0; + if (!k_ok || !v_ok) { + return false; + } + if (Q->ne[1] < 32 || K->ne[1] < 1024) { + return false; + } + for (const ggml_tensor * t : {K, V}) { + if (t->type == GGML_TYPE_F16 && t->nb[1] % (t->ne[0] * 2) != 0) { + return false; + } + } } // gate for the following cases // 1. if the oneDNN graph Add node has no input --> skip @@ -199,13 +221,101 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso dnnl::engine eng = ctx.engine_dnnl(stream); dnnl::stream strm = ctx.stream_dnnl(stream); - // cont/cast inputs to contiguous f16 (head-major) -- the layout the fast systolic path wants. - ggml_sycl_pool_alloc Qf(ctx.pool(), (size_t) H * q * d); - ggml_sycl_pool_alloc Kf(ctx.pool(), (size_t) Hkv * seq * d); - ggml_sycl_pool_alloc Vf(ctx.pool(), (size_t) Hkv * seq * d); - cont_to_f16_sycl ((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream); - cont_to_f16_sycl((const char *) K->data, Kf.get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream); - cont_to_f16_sycl((const char *) V->data, Vf.get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream); + // Q: always f32 -- copy to dense f16. + ggml_sycl_pool_alloc Qf(ctx.pool(), (size_t) H * q * d); + cont_to_f16_sycl((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream); + + // K/V: use pool-alloc for both F16 and dequant paths. + sycl::half * K_ptr = nullptr; + sycl::half * V_ptr = nullptr; + std::optional> Kf_pool; + std::optional> Vf_pool; + + if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) { + Kf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d); + Vf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d); + cont_to_f16_sycl((const char *) K->data, Kf_pool->get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream); + cont_to_f16_sycl((const char *) V->data, Vf_pool->get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream); + K_ptr = Kf_pool->get(); + V_ptr = Vf_pool->get(); + } else if (ggml_is_quantized(K->type)) { + // Quantized K/V: dequant to dense F16 using pool, same lifetime as F16 path. + Kf_pool.emplace(ctx.pool(), ggml_nelements(K)); + K_ptr = Kf_pool->get(); + { + const char * K_data = (const char *)K->data; + const bool k_non_dense = ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]) && K->ne[2] > 1; + const bool k_gemma = k_non_dense && + ((int64_t)K->nb[2] < (int64_t)K->ne[1] * (int64_t)K->nb[1]); + if (ggml_is_contiguously_allocated(K) && !k_non_dense) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst); + to_fp16(K_data, K_ptr, ggml_nelements(K), stream); + } else { + const size_t bs = ggml_blck_size(K->type); + const size_t ts = ggml_type_size(K->type); + to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(K->type); + int64_t s01, s02, s03; + if (k_gemma) { + const int64_t blk_per_row = (int64_t)K->ne[0] / bs; + s01 = (int64_t)Hkv * blk_per_row; + s02 = blk_per_row; + s03 = (int64_t)K->ne[1] * s01; + } else { + s01 = (int64_t)K->nb[1] / ts; + s02 = (int64_t)K->nb[2] / ts; + s03 = (int64_t)K->nb[3] / ts; + } + to_fp16(K_data, K_ptr, + K->ne[0], K->ne[1], K->ne[2], K->ne[3], + s01, s02, s03, stream); + } + } + // Quantized V: always dequant separately. Even when K and V share + // the same underlying allocation (V is a view of K with the same + // data pointer), their logical values differ because the quantized + // elements at different positions/offsets represent different K/V + // data. Master's F16 path also never aliases K and V. + Vf_pool.emplace(ctx.pool(), ggml_nelements(V)); + V_ptr = Vf_pool->get(); + { + const char * V_data = (const char *)V->data; + const bool v_non_dense = ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1; + const bool v_gemma = v_non_dense && + ((int64_t)V->nb[2] < (int64_t)V->ne[1] * (int64_t)V->nb[1]); + if (ggml_is_contiguously_allocated(V) && !v_non_dense) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst); + to_fp16(V_data, V_ptr, ggml_nelements(V), stream); + } else { + const size_t bs = ggml_blck_size(V->type); + const size_t ts = ggml_type_size(V->type); + to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(V->type); + int64_t s01, s02, s03; + if (v_gemma) { + const int64_t blk_per_row = (int64_t)V->ne[0] / bs; + s01 = (int64_t)V->ne[2] * blk_per_row; + s02 = blk_per_row; + s03 = (int64_t)V->ne[1] * s01; + } else { + s01 = (int64_t)V->nb[1] / ts; + s02 = (int64_t)V->nb[2] / ts; + s03 = (int64_t)V->nb[3] / ts; + } + to_fp16(V_data, V_ptr, + V->ne[0], V->ne[1], V->ne[2], V->ne[3], + s01, s02, s03, stream); + } + } + } else { + // F32: strided copy to dense F16 via cont_to_f16_sycl. + Kf_pool.emplace(ctx.pool(), ggml_nelements(K)); + K_ptr = Kf_pool->get(); + cont_to_f16_sycl((const char *) K->data, K_ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], + K->nb[1], K->nb[2], K->nb[3], stream); + Vf_pool.emplace(ctx.pool(), ggml_nelements(V)); + V_ptr = Vf_pool->get(); + cont_to_f16_sycl((const char *) V->data, V_ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], + 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. const sycl::half scale_h = (sycl::half) (1.0f / kq_scale); @@ -230,8 +340,8 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso auto id2ptr = [&](size_t r) -> void * { 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_k) return K_ptr; + if (r == E.id_v) return V_ptr; if (r == E.id_scale) return scbuf.get(); if (r == E.id_mask) return (void *) mask->data; return nullptr; @@ -245,16 +355,11 @@ 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): - 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(); - } + // 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(); } catch (const std::exception & e) { // any oneDNN/SYCL failure is non-fatal: fall back to the existing kernel (strictly additive). diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 1772b9c8584d..5d773aeccd5b 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -97,10 +97,11 @@ static void ggml_sycl_flash_attn_ext_vec(ggml_backend_sycl_context & ctx, ggml_t enum best_fattn_kernel { BEST_FATTN_KERNEL_NONE = 0, BEST_FATTN_KERNEL_VEC = 100, - BEST_FATTN_KERNEL_ONEDNN = 150, // added enum for onednn==150 + BEST_FATTN_KERNEL_ONEDNN = 150, // oneDNN SDPA: native F16 (PR #25222) BEST_FATTN_KERNEL_TILE = 200, }; + static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const ggml_tensor * dst) { GGML_UNUSED(device); #ifndef SYCL_FLASH_ATTN @@ -115,6 +116,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; const ggml_tensor * mask = dst->src[3]; + const ggml_tensor * sinks = dst->src[4]; const int gqa_ratio = Q->ne[2] / K->ne[2]; GGML_ASSERT(Q->ne[2] % K->ne[2] == 0); @@ -122,7 +124,20 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const float max_bias = 0.0f; memcpy(&max_bias, (const float *) KQV->op_params + 1, sizeof(float)); + float logit_softcap = 0.0f; + memcpy(&logit_softcap, (const float *) KQV->op_params + 2, sizeof(float)); + bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; + + // XMX-accelerated path: oneDNN SDPA (native F16 and dequant+non-F16). + // ONEDNN requires min 32 query tokens — short-circuit decode to avoid + // calling _supported() on every decode FA call. + if (Q->ne[1] >= 32 + && ggml_sycl_flash_attn_ext_onednn_supported(dst)) { + return BEST_FATTN_KERNEL_ONEDNN; + } + + for (const ggml_tensor * t : {Q, K, V, mask}) { if (t == nullptr || ggml_is_quantized(t->type)) { continue; @@ -170,6 +185,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const switch (K->type) { case GGML_TYPE_F32: case GGML_TYPE_F16: + case GGML_TYPE_BF16: break; case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -188,8 +204,11 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const return BEST_FATTN_KERNEL_NONE; } - // For small batch sizes the vector kernel may be preferable over the kernels optimized for large batch sizes: - const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0; + // For small batch sizes the vector kernel may be preferable over the kernels optimized for large batch sizes. + // BF16 is excluded: the VEC kernel has no BF16 template (it needs GGML_SYCL_FA_ALL_QUANTS for non-F16/Q4_0/Q8_0). + const bool has_bf16 = (K->type == GGML_TYPE_BF16 || V->type == GGML_TYPE_BF16); + const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0 + && !has_bf16; // Fused-XMX path: oneDNN Graph SDPA (flash attention). Strictly // additive -- taken only when statically supported, otherwise falls through to VEC/TILE below. @@ -216,7 +235,9 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { ggml_sycl_set_device(ctx.device); - switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) { + + const best_fattn_kernel fk = ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst); + switch (fk) { case BEST_FATTN_KERNEL_NONE: GGML_ABORT("Not support Flash-Attention"); case BEST_FATTN_KERNEL_ONEDNN: @@ -233,6 +254,7 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst ggml_sycl_flash_attn_ext_vec(ctx, dst); break; } + } bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst) { From 6e9257c1064f2212b7c6e3a74a6a8052740a36a4 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Thu, 23 Jul 2026 00:15:29 -0500 Subject: [PATCH 2/2] docs: drop GGML_SYCL_FA_DEBUG from SYCL.md (not shipped in this PR) Co-Authored-By: Claude --- docs/backend/SYCL.md | 1 - 1 file changed, 1 deletion(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 157c46c011fe..0814ceb60f92 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -796,7 +796,6 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. | | 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). | -| GGML_SYCL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for flash attention dispatch: kernel name, head_dim, n_kv, and n_kv delta. | | 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 | | UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS | 0 (default) or 1 | Allow SYCL/Unified Runtime Level Zero device allocations larger than 4 GiB. llama.cpp's direct Level Zero allocation path requests the relaxed maximum-size limit itself when GGML_SYCL_ENABLE_LEVEL_ZERO=1. | | GGML_SYCL_USM_SYSTEM | 0 (default) or 1 | Enable experimental support for [USM system allocations](https://github.khronos.org/SYCL_Reference/iface/usm_basic_concept.html#system-allocations) for large GPU buffers. This requires enough host memory for model weights and caches, an Intel Xe2+ GPU such as BMG or newer and supported on Linux only, with CONFIG_DRM_XE_GPUSVM enabled. |