From 32b741c336decea914e4c1c24a9c9815485901b2 Mon Sep 17 00:00:00 2001 From: hmscider <201289679+hmscider@users.noreply.github.com> Date: Wed, 15 Jul 2026 03:26:53 -0400 Subject: [PATCH 1/6] [SYCL] Flash Attention with XMX engine via oneDNN (#25222) * [SYCL] F16 (default) Flash Attention with XMX engine via oneDNN graph API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k * [SYCL] Address review on FA oneDNN path. Result: llama-bench---pp512; 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf * PR-25222 revision v2: addressed audits * [SYCL] flash-attn oneDNN SDPA KV F16 rev 3.0: add BMG gate + multi-device sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious. Co-authored-by: maxious <81432+maxious@users.noreply.github.com> * updated comment on bmg gate, noted the issue --------- Co-authored-by: scientist3 Co-authored-by: hmscider Co-authored-by: maxious <81432+maxious@users.noreply.github.com> --- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/fattn-onednn.cpp | 265 ++++++++++++++++++++++++++++ ggml/src/ggml-sycl/fattn-onednn.hpp | 14 ++ ggml/src/ggml-sycl/fattn.cpp | 15 +- ggml/src/ggml-sycl/ggml-sycl.cpp | 4 + 5 files changed, 298 insertions(+), 1 deletion(-) create mode 100644 ggml/src/ggml-sycl/fattn-onednn.cpp create mode 100644 ggml/src/ggml-sycl/fattn-onednn.hpp diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index be0db6b2f..e5d9ee89d 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -64,6 +64,7 @@ extern int g_ggml_sycl_enable_fusion; 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; #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 new file mode 100644 index 000000000..f2e12ef1a --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -0,0 +1,265 @@ +#include +#include +#include +#include +#include +#include + +#include "fattn-onednn.hpp" +#include "fattn-tile.hpp" + +// set minimum query length to treat as prefill (32) +#define GGML_SYCL_FA_ONEDNN_MIN_Q 32 + +bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { +#if !GGML_SYCL_DNNL + GGML_UNUSED(dst); + return false; +#else + if (!g_ggml_sycl_fa_onednn) { + return false; + } + // Battlemage (Xe2) only, for now. On other Intel archs oneDNN's fused SDPA returns wrong results + // for some shapes (e.g. head_dim=64 on Arc / xe_hpg) -- an oneDNN bug tracked upstream at + // https://github.com/uxlfoundation/oneDNN/issues/5510. Remove this hardware limitation once that + // is fixed; until then non-BMG archs fall back to the existing FA kernel. + const gpu_arch arch = ggml_sycl_info().devices[ggml_sycl_get_device()].hw_info.arch; + if (arch != gpu_arch::intel_gpu_bmg_g21 && arch != gpu_arch::intel_gpu_bmg_g31) { + return false; + } + const ggml_tensor * Q = dst->src[0]; + 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]; + + // gate for f16 KV only for now + // need to implement quantized KV + if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) { + 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 + // 3. the mask must be shape [1, 1, q, seq] + // 4. sinks: excludes attention sink (Xiao et al., 2024) that can't be modeled by oneDNN graph + if (!mask || mask->type != GGML_TYPE_F16 || mask->ne[2] != 1 || mask->ne[3] != 1 || sinks) { + return false; + } + float max_bias = 0.0f, logit_softcap = 0.0f; + memcpy(&max_bias, (const float *) dst->op_params + 1, sizeof(float)); + memcpy(&logit_softcap, (const float *) dst->op_params + 2, sizeof(float)); + if (max_bias != 0.0f || logit_softcap != 0.0f) { + return false; + } + // K and V must share head_dim: the SDPA graph uses a single `d` for both. + const int64_t d = K->ne[0]; + if (V->ne[0] != d || Q->ne[3] != 1) { + return false; + } + // GQA must divide evenly. + if (K->ne[2] == 0 || Q->ne[2] % K->ne[2] != 0) { + return false; + } + // Prefill only. + if (Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) { + return false; + } + return true; +#endif +} + +#if GGML_SYCL_DNNL + +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" // graph API lives only under oneapi/dnnl/, not at the include root + +using namespace dnnl; +using namespace dnnl::graph; + +// strided src (f16 or f32) -> contiguous f16 [ne0,ne1,ne2,ne3] (ne0 innermost). nb* are BYTE strides. +template +static void cont_to_f16_sycl(const char * src, sycl::half * dst, + int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, + size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream) { + const int64_t n = ne0 * ne1 * ne2 * ne3; + stream->parallel_for(sycl::range<1>(n), [=](sycl::id<1> ix) { + const int64_t gid = ix[0]; + int64_t i = gid; + const int64_t i0 = i % ne0; i /= ne0; + const int64_t i1 = i % ne1; i /= ne1; + const int64_t i2 = i % ne2; const int64_t i3 = i / ne2; + const src_t * p = (const src_t *) (src + i1 * nb1 + i2 * nb2 + i3 * nb3) + i0; + dst[gid] = (sycl::half) (*p); + }); +} + +// oneDNN SDPA out (f16 contiguous [mb,H,q,d]) -> ggml dst (f32 [head_dim,H,n_tok,mb], contiguous). +static void permute_sdpa_out_sycl(const sycl::half * out, float * dst, + int64_t mb, int64_t H, int64_t q, int64_t d, dpct::queue_ptr stream) { + const int64_t n = mb * H * q * d; + stream->parallel_for(sycl::range<1>(n), [=](sycl::id<1> ix) { + const int64_t gid = ix[0]; + int64_t i = gid; + const int64_t e = i % d; i /= d; + const int64_t t = i % q; i /= q; + const int64_t h = i % H; const int64_t b = i / H; + dst[e + h * d + t * d * H + b * d * H * q] = (float) out[gid]; + }); +} + +struct sdpa_partition { + compiled_partition cp; + std::vector ins; + logical_tensor out; + size_t id_q = 0, id_k = 0, id_v = 0, id_scale = 0, id_mask = 0; + bool ok = false; +}; + +// Build + compile the contiguous-input GQA SDPA graph (MatMul->Divide->Add->SoftMax->MatMul), f32 out. +// Mirrors the hardware-verified scratch/onednn_sdpa_probe.cpp build_gqa (partitions=1, sdp_primitive_kernel_t). +static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int seq, int d) { + using ltype = logical_tensor::layout_type; + using dt = logical_tensor::data_type; + using ldims = logical_tensor::dims; + const dt fi = dt::f32, t = dt::f16; + const int rep = H / Hkv; + const ldims q_sz = {1, Hkv, rep, q, d}, kv_sz = {1, Hkv, 1, seq, d}, s_sz = {1, Hkv, rep, q, seq}, + sc = {1, 1, 1, 1, 1}, msk = {1, 1, 1, q, seq}, o_sz = {1, Hkv, rep, q, d}; + int64_t id = 0; + sdpa_partition E; + + auto query = logical_tensor(id++, t, q_sz, ltype::strided); + auto key = logical_tensor(id++, t, kv_sz, ltype::strided); + auto score = logical_tensor(id++, fi, s_sz, ltype::strided); + auto bmm1 = op(id++, op::kind::MatMul, "bmm1"); + bmm1.set_attr(op::attr::transpose_b, true); // key is [.., seq, d] + bmm1.add_inputs({query, key}); bmm1.add_outputs({score}); + + auto scale = logical_tensor(id++, t, sc, ltype::strided); + auto scaled = logical_tensor(id++, fi, s_sz, ltype::strided); + auto sdiv = op(id++, op::kind::Divide, "scale_div"); // score / (1/kq_scale) == score * kq_scale + sdiv.add_inputs({score, scale}); sdiv.add_outputs({scaled}); + + auto mask = logical_tensor(id++, t, msk, ltype::strided); + auto masked = logical_tensor(id++, fi, s_sz, ltype::strided); + auto madd = op(id++, op::kind::Add, "mask_add"); + madd.add_inputs({scaled, mask}); madd.add_outputs({masked}); + + auto probs = logical_tensor(id++, t, s_sz, ltype::strided); + auto smax = op(id++, op::kind::SoftMax, "softmax"); + smax.set_attr(op::attr::axis, -1); + smax.set_attr(op::attr::mode, "inf_as_zero"); + smax.add_inputs({masked}); smax.add_outputs({probs}); + + auto value = logical_tensor(id++, t, kv_sz, ltype::strided); + // f16 output is REQUIRED to hit sdp_primitive_kernel_t (the systolic micro-kernel); an f32 output + // falls to larger_partition_kernel_t which materializes N^2 (confirmed: scratch/onednn_sdpa_kernel_probe.cpp). + // converted to the f32 ggml dst in the permute below. + auto output = logical_tensor(id++, t, o_sz, ltype::strided); // f16 contiguous [mb,Hkv,rep,q,d] + auto bmm2 = op(id++, op::kind::MatMul, "bmm2"); + bmm2.add_inputs({probs, value}); bmm2.add_outputs({output}); + + dnnl::graph::graph g(eng.get_kind()); + g.add_op(bmm1); g.add_op(sdiv); g.add_op(madd); g.add_op(smax); g.add_op(bmm2); + g.finalize(); + + auto parts = g.get_partitions(); + if (parts.size() != 1 || !parts[0].is_supported()) { + return E; // ok stays false -> caller falls back to TILE + } + E.ins = parts[0].get_input_ports(); + E.out = parts[0].get_output_ports()[0]; + E.cp = parts[0].compile(E.ins, {E.out}, eng); + E.out = E.cp.query_logical_tensor(E.out.get_id()); + E.id_q = query.get_id(); E.id_k = key.get_id(); E.id_v = value.get_id(); + E.id_scale = scale.get_id(); E.id_mask = mask.get_id(); + E.ok = true; + return E; +} + +void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) try { + const ggml_tensor * Q = dst->src[0]; + const ggml_tensor * K = dst->src[1]; + const ggml_tensor * V = dst->src[2]; + const ggml_tensor * mask = dst->src[3]; + + const int64_t d = K->ne[0]; // head_dim + const int64_t seq = K->ne[1]; // n_kv + const int64_t Hkv = K->ne[2]; // n_head_kv + const int64_t H = Q->ne[2]; // n_head + const int64_t q = Q->ne[1]; // n_tok + const int64_t mb = Q->ne[3]; // batch (== 1, gated) + + float kq_scale = 1.0f; + memcpy(&kq_scale, (const float *) dst->op_params + 0, sizeof(float)); + + dpct::queue_ptr stream = ctx.stream(); + 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); + + // 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); + ggml_sycl_pool_alloc scbuf(ctx.pool(), 1); + stream->memcpy(scbuf.get(), &scale_h, sizeof(sycl::half)); + + ggml_sycl_pool_alloc outf(ctx.pool(), (size_t) H * q * d); // f16 contiguous SDPA out [mb,H,q,d] + + // compile once per (device, shape), reuse across layers/calls. + static std::unordered_map cache; + char keyb[96]; + snprintf(keyb, sizeof(keyb), "%d:%lld:%lld:%lld:%lld:%lld", ggml_sycl_get_device(), + (long long) H, (long long) Hkv, (long long) q, (long long) seq, (long long) d); + auto it = cache.find(keyb); + if (it == cache.end()) { + it = cache.emplace(keyb, build_sdpa(eng, (int) H, (int) Hkv, (int) q, (int) seq, (int) d)).first; + } + sdpa_partition & E = it->second; + // _supported() is authoritative: if it accepted this op the partition must build. + // A failure here is a gap in _supported() -- surface it, don't mask it with a fallback. + GGML_ASSERT(E.ok && "oneDNN SDPA partition failed to build for a _supported() shape"); + + 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_scale) return scbuf.get(); + if (r == E.id_mask) return (void *) mask->data; + return nullptr; + }; + std::vector ti; + ti.reserve(E.ins.size()); + for (auto & lt : E.ins) { + ti.emplace_back(lt, eng, id2ptr(lt.get_id())); + } + tensor to(E.out, eng, outf.get()); + 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(); + } +} +catch (const std::exception & e) { + // any oneDNN/SYCL failure is non-fatal: fall back to the existing kernel (strictly additive). + GGML_LOG_WARN("%s: oneDNN SDPA failed (%s); falling back to TILE kernel\n", __func__, e.what()); + ggml_sycl_flash_attn_ext_tile(ctx, dst); +} + +#endif // GGML_SYCL_DNNL diff --git a/ggml/src/ggml-sycl/fattn-onednn.hpp b/ggml/src/ggml-sycl/fattn-onednn.hpp new file mode 100644 index 000000000..d3019e876 --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-onednn.hpp @@ -0,0 +1,14 @@ +#ifndef GGML_SYCL_FATTN_ONEDNN_HPP +#define GGML_SYCL_FATTN_ONEDNN_HPP + +#include "common.hpp" + +// Static-only check: fused-XMX oneDNN Graph SDPA path==flash-attn op +// (f16 KV, no softcap/ALiBi, single stream, tuned head_dim, prefill-sized q.) +bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst); + +// Run flash attention through oneDNN's fused xmx SDPA +// execute the cached SDPA partition, write the f32 dst. Falls back to the TILE kernel on any failure. +void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + +#endif // GGML_SYCL_FATTN_ONEDNN_HPP diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 7c6e6112f..1772b9c85 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -18,6 +18,7 @@ #include "fattn-tile.hpp" #include "fattn-vec.hpp" #include "fattn.hpp" +#include "fattn-onednn.hpp" #define FATTN_VEC_CASE(D, type_K, type_V) \ @@ -96,6 +97,7 @@ 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_TILE = 200, }; @@ -189,7 +191,11 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // 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; - // Todo: Use the XMX kernel if possible: + // Fused-XMX path: oneDNN Graph SDPA (flash attention). Strictly + // additive -- taken only when statically supported, otherwise falls through to VEC/TILE below. + if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) { + return BEST_FATTN_KERNEL_ONEDNN; + } // If there are no tensor cores available, use the generic tile kernel: if (can_use_vector_kernel) { @@ -213,6 +219,13 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) { case BEST_FATTN_KERNEL_NONE: GGML_ABORT("Not support Flash-Attention"); + case BEST_FATTN_KERNEL_ONEDNN: + // guarded: ggml_sycl_flash_attn_ext_onednn() is only defined under GGML_SYCL_DNNL; + // the reference must be compiled out here or the GGML_SYCL_DNNL=0 build fails to link. +#if GGML_SYCL_DNNL + ggml_sycl_flash_attn_ext_onednn(ctx, dst); +#endif + break; case BEST_FATTN_KERNEL_TILE: ggml_sycl_flash_attn_ext_tile(ctx, dst); break; diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 444255cdf..df7159db5 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -84,6 +84,7 @@ int g_ggml_sycl_debug = 0; 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_enable_vmm = 1; int g_ggml_sycl_enable_fusion = 1; int g_ggml_sycl_prioritize_dmmv = 0; @@ -285,6 +286,7 @@ static void ggml_check_sycl() try { g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1); 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_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); @@ -352,8 +354,10 @@ static void ggml_check_sycl() try { #if defined(GGML_SYCL_DNNL) GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: %d\n", g_ggml_sycl_enable_dnn); + GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN: %d\n", g_ggml_sycl_fa_onednn); #else 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 #ifdef SYCL_FLASH_ATTN GGML_LOG_INFO(" GGML_SYCL_ENABLE_FLASH_ATTN: %d\n", g_ggml_sycl_enable_flash_attention); From 0e148a573f0ccbb0e9000312e2757b8998c1e070 Mon Sep 17 00:00:00 2001 From: Francois Dugast Date: Wed, 15 Jul 2026 09:28:24 +0200 Subject: [PATCH 2/6] sycl: Increase minimum buffer size for USM system allocations (#25525) Raise the threshold for minimum buffer size from 1 GiB to 4 GiB, based on real-world experiments of overcommitting device memory with model weights larger than available VRAM, for example Qwen3.5-35B-A3B-Q8 running on a B70. Also add a debug message to better track USM system allocations. Signed-off-by: Francois Dugast --- ggml/src/ggml-sycl/ggml-sycl.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index df7159db5..f534d237d 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -843,7 +843,7 @@ static const char * ggml_backend_sycl_buffer_type_get_name(ggml_backend_buffer_t } static bool check_usm_system(int device, size_t size) { - bool use_usm_system = g_ggml_sycl_usm_system && size >= MEM_SIZE_1G; + bool use_usm_system = g_ggml_sycl_usm_system && size >= ((size_t)4 * MEM_SIZE_1G); if (use_usm_system && !ggml_sycl_info().devices[device].usm_system_support) { GGML_LOG_INFO("Device does not support USM system allocations\n"); @@ -882,6 +882,7 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, void * dev_ptr; if (use_usm_system) { + GGML_SYCL_DEBUG("[SYCL] allocating %lu Bytes with USM system\n", size); dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size); if (!dev_ptr) { GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size); From 22b208b1cacb67bae191b00d795dae7cc819edb8 Mon Sep 17 00:00:00 2001 From: Andrew Smith Date: Wed, 15 Jul 2026 00:29:12 -0700 Subject: [PATCH 3/6] sycl : implement xielu op (#25550) --- docs/ops.md | 2 +- docs/ops/SYCL.csv | 8 +++--- ggml/src/ggml-sycl/element_wise.cpp | 40 +++++++++++++++++++++++++++++ ggml/src/ggml-sycl/element_wise.hpp | 2 ++ ggml/src/ggml-sycl/ggml-sycl.cpp | 4 +++ 5 files changed, 51 insertions(+), 5 deletions(-) diff --git a/docs/ops.md b/docs/ops.md index f13875385..c56015236 100644 --- a/docs/ops.md +++ b/docs/ops.md @@ -120,4 +120,4 @@ Legend: | TRI | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | | TRUNC | ❌ | ❌ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | | UPSCALE | ❌ | 🟡 | ✅ | ✅ | ❌ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ | -| XIELU | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | +| XIELU | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | diff --git a/docs/ops/SYCL.csv b/docs/ops/SYCL.csv index 8c94d14b5..b563e76a8 100644 --- a/docs/ops/SYCL.csv +++ b/docs/ops/SYCL.csv @@ -11600,10 +11600,10 @@ zjy 2 "SYCL0","CUMSUM","type=f32,ne=[242004,1,1,1]","support","1","yes","SYCL" "SYCL0","CUMSUM","type=f32,ne=[375960,1,1,1]","support","1","yes","SYCL" "SYCL0","CUMSUM","type=f32,ne=[20481,4,1,1]","support","1","yes","SYCL" -"SYCL0","XIELU","type=f32,ne=[10,5,4,3]","support","0","no","SYCL" -"SYCL0","XIELU","type=f16,ne=[10,5,4,3]","support","0","no","SYCL" -"SYCL0","XIELU","type=f32,ne=[512,16,1,1]","support","0","no","SYCL" -"SYCL0","XIELU","type=f16,ne=[512,16,1,1]","support","0","no","SYCL" +"SYCL0","XIELU","type=f32,ne=[10,5,4,3]","support","1","yes","SYCL" +"SYCL0","XIELU","type=f16,ne=[10,5,4,3]","support","1","yes","SYCL" +"SYCL0","XIELU","type=f32,ne=[512,16,1,1]","support","1","yes","SYCL" +"SYCL0","XIELU","type=f16,ne=[512,16,1,1]","support","1","yes","SYCL" "SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=3","support","1","yes","SYCL" "SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=2","support","1","yes","SYCL" "SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=1","support","1","yes","SYCL" diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp index bae157a48..b2406e11b 100644 --- a/ggml/src/ggml-sycl/element_wise.cpp +++ b/ggml/src/ggml-sycl/element_wise.cpp @@ -247,6 +247,17 @@ static __dpct_inline__ T op_leaky_relu(T x, float negative_slope) { } } +template +static __dpct_inline__ T op_xielu(T x, float alpha_n, float alpha_p, float beta, float eps) { + const float xi = static_cast(x); + const float gate_pos = (xi > 0.0f); + const float y_pos = alpha_p * xi * xi + beta * xi; + const float min_v_eps = sycl::fmin(xi, eps); + const float y_neg = (sycl::expm1(min_v_eps) - xi) * alpha_n + beta * xi; + const float out = gate_pos * y_pos + (1.0f - gate_pos) * y_neg; + return static_cast(out); +} + template static __dpct_inline__ T op_sqr(T x) { return x * x; @@ -359,6 +370,13 @@ static void unary_op_leaky_relu_kernel(const T * x, T * dst, const int k, float } } +template +static void unary_op_xielu_kernel(const T * x, T * dst, const int k, float alpha_n, float alpha_p, float beta, float eps, const sycl::nd_item<1> &item_ct1) { + SYCL_GLOBAL_ID_LOOP(k, item_ct1) { + dst[i] = op_xielu(x[i], alpha_n, alpha_p, beta, eps); + } +} + template static void unary_op_sqr_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) { SYCL_GLOBAL_ID_LOOP(k, item_ct1) { @@ -836,6 +854,23 @@ static inline void ggml_sycl_op_clamp(ggml_backend_sycl_context & ctx, ggml_tens }, min_val, max_val); } +static inline void ggml_sycl_op_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + const float alpha_n = ggml_get_op_params_f32(dst, 1); + const float alpha_p = ggml_get_op_params_f32(dst, 2); + const float beta = ggml_get_op_params_f32(dst, 3); + const float eps = ggml_get_op_params_f32(dst, 4); + ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst, + [](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream, float alpha_n_arg, float alpha_p_arg, float beta_arg, float eps_arg) { + const int num_blocks = ceil_div(k_elements, SYCL_RELU_BLOCK_SIZE); + stream->parallel_for( + sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(SYCL_RELU_BLOCK_SIZE), + sycl::range<1>(SYCL_RELU_BLOCK_SIZE)), + [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + unary_op_xielu_kernel(src, dst_ptr, k_elements, alpha_n_arg, alpha_p_arg, beta_arg, eps_arg, item_ct1); + }); + }, alpha_n, alpha_p, beta, eps); +} + static inline void ggml_sycl_op_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { ggml_sycl_detail::ggml_sycl_op_unary(ctx, dst, [](auto x) { return op_floor(x); @@ -1153,6 +1188,11 @@ void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { ggml_sycl_op_clamp(ctx, dst); } +void ggml_sycl_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1); + ggml_sycl_op_xielu(ctx, dst); +} + void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1); ggml_sycl_op_sgn(ctx, dst); diff --git a/ggml/src/ggml-sycl/element_wise.hpp b/ggml/src/ggml-sycl/element_wise.hpp index 3bdc38596..beea052cf 100644 --- a/ggml/src/ggml-sycl/element_wise.hpp +++ b/ggml/src/ggml-sycl/element_wise.hpp @@ -75,6 +75,8 @@ void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst); void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst); +void ggml_sycl_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst); void ggml_sycl_abs(ggml_backend_sycl_context & ctx, ggml_tensor * dst); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index f534d237d..cb8974eed 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5011,6 +5011,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg case GGML_UNARY_OP_ELU: ggml_sycl_elu(ctx, dst); break; + case GGML_UNARY_OP_XIELU: + ggml_sycl_xielu(ctx, dst); + break; case GGML_UNARY_OP_FLOOR: ggml_sycl_floor(ctx, dst); break; @@ -5673,6 +5676,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons case GGML_UNARY_OP_EXPM1: case GGML_UNARY_OP_SOFTPLUS: case GGML_UNARY_OP_ELU: + case GGML_UNARY_OP_XIELU: case GGML_UNARY_OP_CEIL: return true; case GGML_UNARY_OP_FLOOR: From ae9291e16b976514bf1f3d7f1616da7db3459496 Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Wed, 15 Jul 2026 15:31:10 +0800 Subject: [PATCH 4/6] sycl : support kernel type fp16 for conv2d_dw (#25653) --- ggml/src/ggml-sycl/conv2d-dw.cpp | 38 +++++++++++++++++++++----------- 1 file changed, 25 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/conv2d-dw.cpp b/ggml/src/ggml-sycl/conv2d-dw.cpp index 0a52b7917..8755a4c95 100644 --- a/ggml/src/ggml-sycl/conv2d-dw.cpp +++ b/ggml/src/ggml-sycl/conv2d-dw.cpp @@ -71,8 +71,8 @@ struct dw_cwhn_layout { } }; -template -static void conv2d_dw_kernel(const float * input, const float * kernel, float * output, +template +static void conv2d_dw_kernel(const float * input, const KernelT * kernel, float * output, const conv2d_dw_params p, const sycl::nd_item<3> & item_ct1) { const int global_idx = item_ct1.get_local_id(2) + item_ct1.get_group(2) * item_ct1.get_local_range(2); @@ -93,15 +93,15 @@ static void conv2d_dw_kernel(const float * input, const float * kernel, float * for (int kx = bounds.x_min; kx < bounds.x_max; ++kx) { const int in_x = dw_calculate_input_coord(out_x, kx, p.stride_x, p.dilation_x, p.padding_x); acc += input[Layout::input_index(n, c, in_y, in_x, p)] * - kernel[Layout::kernel_index(c, ky, kx, p)]; + static_cast(kernel[Layout::kernel_index(c, ky, kx, p)]); } } output[Layout::output_index(n, c, out_y, out_x, p)] = acc; } -template -static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d, +template +static void conv2d_dw_sycl(const float * x_d, const KernelT * w_d, float * y_d, const conv2d_dw_params p, const queue_ptr & stream) { const int total = p.batches * p.channels * p.out_h * p.out_w; const int num_blocks = (total + SYCL_CONV2D_DW_BLOCK_SIZE - 1) / SYCL_CONV2D_DW_BLOCK_SIZE; @@ -109,7 +109,7 @@ static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d, const sycl::range<3> block_nums(1, 1, num_blocks); stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) { - conv2d_dw_kernel(x_d, w_d, y_d, p, item_ct1); + conv2d_dw_kernel(x_d, w_d, y_d, p, item_ct1); }); } @@ -119,9 +119,9 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) const ggml_tensor * kernel = dst->src[0]; const ggml_tensor * input = dst->src[1]; - GGML_ASSERT(kernel->type == GGML_TYPE_F32 && input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); + GGML_ASSERT((kernel->type == GGML_TYPE_F32 || kernel->type == GGML_TYPE_F16) && + input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); - const float * w_d = (const float *) kernel->data; const float * x_d = (const float *) input->data; float * y_d = (float *) dst->data; @@ -148,11 +148,23 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) const queue_ptr stream = ctx.stream(); - if (ggml_is_contiguous(input)) { - conv2d_dw_sycl(x_d, w_d, y_d, params, stream); - } else if (ggml_is_contiguous_channels(input)) { - conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + if (kernel->type == GGML_TYPE_F16) { + const sycl::half * w_d = (const sycl::half *) kernel->data; + if (ggml_is_contiguous(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else if (ggml_is_contiguous_channels(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else { + GGML_ABORT("Unsupported memory layout for conv2d_dw"); + } } else { - GGML_ABORT("Unsupported memory layout for conv2d_dw"); + const float * w_d = (const float *) kernel->data; + if (ggml_is_contiguous(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else if (ggml_is_contiguous_channels(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else { + GGML_ABORT("Unsupported memory layout for conv2d_dw"); + } } } From d3fba0c79db8d25a76031783bf071d02234812e6 Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Wed, 15 Jul 2026 15:32:28 +0800 Subject: [PATCH 5/6] sycl : fix get_rows Q2_K, Q4_K, Q5_K (#25656) --- ggml/src/ggml-sycl/dequantize.hpp | 131 ++++++++++++++++++++++++++---- ggml/src/ggml-sycl/getrows.cpp | 83 ++++++++++++++++++- 2 files changed, 195 insertions(+), 19 deletions(-) diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 7b66c73b0..3db55319f 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -19,6 +19,7 @@ typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, dfloat2 & v); typedef void (*dequantize_kernel_t_reorder)(const void *d, const int64_t ib, const void *qs, const int iqs, dfloat2 &v); +typedef void (*dequantize_kernel_f32_t)(const void * vx, const int64_t ib, const int iqs, float & v0, float & v1); #if QK_K == 256 static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m); @@ -85,6 +86,21 @@ static __dpct_inline__ void dequantize_q1_0_reorder(const void *d_ptr, const int v.y() = (2 * bit_1 - 1) * d; } +static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib, + const int iqs, dfloat2 &v) { + const block_q1_0 * x = (const block_q1_0 *) vx; + const dfloat d = x[ib].d; + + const int bit_index_0 = iqs + 0; + const int bit_index_1 = iqs + 1; + + const int bit_0 = (x[ib].qs[bit_index_0 / 8] >> (bit_index_0 % 8)) & 1; + const int bit_1 = (x[ib].qs[bit_index_1 / 8] >> (bit_index_1 % 8)) & 1; + + v.x() = (2 * bit_0 - 1) * d; + v.y() = (2 * bit_1 - 1) * d; +} + static __dpct_inline__ void dequantize_q4_1(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { const block_q4_1 * x = (const block_q4_1 *) vx; @@ -140,6 +156,39 @@ static __dpct_inline__ void dequantize_q4_K(const void *vx, const int64_t ib, #endif } +static __dpct_inline__ void dequantize_q4_K_f32(const void *vx, const int64_t ib, + const int iqs, float &v0, float &v1) { +#if QK_K == 256 + const block_q4_K * x = (const block_q4_K *) vx; + const sycl::half2 dm = x[ib].dm; + const float dall = dm[0]; + const float dmin = dm[1]; + + auto dequantize_one = [&](const int idx) -> float { + const int il = idx / 64; + const int in = idx % 64; + const int is = 2 * il + (in >= 32 ? 1 : 0); + const int qsi = 32 * il + (in & 31); + + uint8_t sc; + uint8_t m; + get_scale_min_k4(is, x[ib].scales, sc, m); + + const float d = dall * sc; + const float mn = dmin * m; + const uint8_t q = x[ib].qs[qsi]; + const uint8_t qv = (in >= 32) ? (q >> 4) : (q & 0xF); + + return d * qv - mn; + }; + + v0 = dequantize_one(iqs + 0); + v1 = dequantize_one(iqs + 1); +#else + GGML_ABORT("Q4_K dequantize not supported for QK_K != 256"); +#endif +} + static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { #if QK_K == 256 @@ -159,7 +208,7 @@ static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib, const float d = dall * (sc & 0xF); const float m = dmin * (sc >> 4); - return sycl::fma((dfloat) ((q >> (2 * g)) & 3), (dfloat) d, (dfloat) (-m)); + return (dfloat) d * (dfloat) ((q >> (2 * g)) & 3) - (dfloat) m; }; v.x() = dequantize_one(iqs + 0); @@ -169,6 +218,35 @@ static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib, #endif } +static __dpct_inline__ void dequantize_q2_K_f32(const void *vx, const int64_t ib, + const int iqs, float &v0, float &v1) { +#if QK_K == 256 + const block_q2_K * x = (const block_q2_K *) vx; + const float dall = x[ib].dm[0]; + const float dmin = x[ib].dm[1]; + + auto dequantize_one = [&](const int idx) -> float { + const int n = idx / 128; + const int r = idx % 128; + const int g = r / 32; + const int l = r % 32; + const int is = 8 * n + l / 16; + + const uint8_t q = x[ib].qs[32 * n + l]; + const uint8_t sc = x[ib].scales[is + 2 * g]; + const float d = dall * (sc & 0xF); + const float m = dmin * (sc >> 4); + + return d * ((q >> (2 * g)) & 3) - m; + }; + + v0 = dequantize_one(iqs + 0); + v1 = dequantize_one(iqs + 1); +#else + GGML_ABORT("Q2_K dequantize not supported for QK_K != 256"); +#endif +} + static __dpct_inline__ void dequantize_q3_K(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { #if QK_K == 256 @@ -242,6 +320,42 @@ static __dpct_inline__ void dequantize_q5_K(const void *vx, const int64_t ib, #endif } +static __dpct_inline__ void dequantize_q5_K_f32(const void *vx, const int64_t ib, + const int iqs, float &v0, float &v1) { +#if QK_K == 256 + const block_q5_K * x = (const block_q5_K *) vx; + const float dall = x[ib].dm[0]; + const float dmin = x[ib].dm[1]; + + auto dequantize_one = [&](const int idx) -> float { + const int il = idx / 64; + const int in = idx % 64; + const int is = 2 * il + (in >= 32 ? 1 : 0); + const int ir = (in & 31) / 2; + const int iq = in & 1; + + const uint8_t q = x[ib].qs[32 * il + 2 * ir + iq]; + const uint8_t h = x[ib].qh[2 * ir + iq]; + const uint8_t qv = (in >= 32) ? (q >> 4) : (q & 0xF); + + uint8_t sc; + uint8_t m; + get_scale_min_k4(is, x[ib].scales, sc, m); + + const float d = dall * sc; + const float mn = dmin * m; + const uint8_t hm = 1 << (2 * il + (in >= 32 ? 1 : 0)); + + return (qv + ((h & hm) ? 16 : 0)) * d - mn; + }; + + v0 = dequantize_one(iqs + 0); + v1 = dequantize_one(iqs + 1); +#else + GGML_ABORT("Q5_K dequantize not supported for QK_K != 256"); +#endif +} + static __dpct_inline__ void dequantize_q6_K(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { #if QK_K == 256 @@ -296,21 +410,6 @@ static __dpct_inline__ void dequantize_mxfp4(const void *vx, const int64_t ib, v.y() = d * kvalues_mxfp4[q >> 4] * 0.5f; } -static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib, - const int iqs, dfloat2 &v) { - const block_q1_0 * x = (const block_q1_0 *) vx; - const dfloat d = x[ib].d; - - const int bit_index_0 = iqs + 0; - const int bit_index_1 = iqs + 1; - - const int bit_0 = (x[ib].qs[bit_index_0 / 8] >> (bit_index_0 % 8)) & 1; - const int bit_1 = (x[ib].qs[bit_index_1 / 8] >> (bit_index_1 % 8)) & 1; - - v.x() = (2 * bit_0 - 1) * d; - v.y() = (2 * bit_1 - 1) * d; -} - static __dpct_inline__ void dequantize_nvfp4(const void *vx, const int64_t ib, const int iqs, dfloat2 &v) { const block_nvfp4 & xb = ((const block_nvfp4 *) vx)[ib]; diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 298f247f8..2113f3563 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -60,6 +60,50 @@ static void k_get_rows( dst_row[iybs + iqs + y_offset] = v.y(); } +template +static void k_get_rows_f32( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, + int64_t ne12, + size_t s1, size_t s2, size_t s3, + size_t nb01, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + const sycl::nd_item<3> &item_ct1) { + + const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) + + item_ct1.get_local_id(2)) * + 2; + const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) + + item_ct1.get_local_id(1); + const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + + item_ct1.get_local_id(0)) / + ne12; + const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + + item_ct1.get_local_id(0)) % + ne12; + + if (i00 >= ne00) { + return; + } + + const int i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + const void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03; + + const int ib = i00/qk; + const int iqs = (i00%qk)/qr; + const int iybs = i00 - i00%qk; + const int y_offset = qr == 1 ? 1 : qk/2; + + float v0; + float v1; + dequantize_kernel(src0_row, ib, iqs, v0, v1); + + dst_row[iybs + iqs + 0] = (dst_t) v0; + dst_row[iybs + iqs + y_offset] = (dst_t) v1; +} + template static void k_get_rows_float( const src0_t * src0, const int32_t * src1, dst_t * dst, @@ -129,6 +173,39 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr GGML_UNUSED(ctx); } +template +static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, + ggml_tensor *dst, const void *src0_dd, + const int32_t *src1_dd, float *dst_dd, + queue_ptr stream) { + + GGML_TENSOR_BINARY_OP_LOCALS + + const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE); + const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE); + const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x); + + const size_t s1 = nb1 / ggml_element_size(dst); + const size_t s2 = nb2 / ggml_element_size(dst); + const size_t s3 = nb3 / ggml_element_size(dst); + + const size_t s10 = nb10 / ggml_element_size(src1); + const size_t s11 = nb11 / ggml_element_size(src1); + const size_t s12 = nb12 / ggml_element_size(src1); + + GGML_ASSERT(ne00 % 2 == 0); + + stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_f32( + src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2, + s3, nb01, nb02, nb03, s10, s11, s12, item_ct1); + }); + + GGML_UNUSED(dst); + GGML_UNUSED(ctx); +} + template static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, @@ -244,7 +321,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q2_K: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q3_K: @@ -260,7 +337,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q4_K: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q5_0: @@ -272,7 +349,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q5_K: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q6_K: From 33a75f41c30052fd3d1c38e8ed2f86ee3c3f8fba Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Wed, 15 Jul 2026 15:47:18 +0800 Subject: [PATCH 6/6] DeepseekV4: reduce graph splits (#25702) --- src/models/deepseek4.cpp | 26 ++++++++++++++------------ 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/src/models/deepseek4.cpp b/src/models/deepseek4.cpp index e64e43451..579608b3e 100644 --- a/src/models/deepseek4.cpp +++ b/src/models/deepseek4.cpp @@ -435,27 +435,29 @@ ggml_tensor * llama_model_deepseek4::graph::build_overlap_compressed_kv_from_sta kv_state = dsv4_append_zero_row(ctx0, kv_state, false); score_state = dsv4_append_zero_row(ctx0, score_state, true); - ggml_tensor * prev_idxs = dsv4_view_1d(ctx0, state_read_idxs, ratio*n_blocks, 0); - ggml_tensor * cur_idxs = dsv4_view_1d(ctx0, state_read_idxs, ratio*n_blocks, ratio*n_blocks); + const int64_t n_read = ratio*n_blocks; - ggml_tensor * kv_prev = ggml_get_rows(ctx0, kv_state, prev_idxs); - kv_prev = ggml_cont(ctx0, ggml_view_2d(ctx0, kv_prev, n_embd_head, ratio*n_blocks, kv_prev->nb[1], 0)); + ggml_tensor * kv_rows = ggml_get_rows(ctx0, kv_state, state_read_idxs); + ggml_tensor * score_rows = ggml_get_rows(ctx0, score_state, state_read_idxs); + + ggml_tensor * kv_prev = ggml_cont(ctx0, + ggml_view_2d(ctx0, kv_rows, n_embd_head, n_read, kv_rows->nb[1], 0)); kv_prev = ggml_reshape_3d(ctx0, kv_prev, n_embd_head, ratio, n_blocks); cb(kv_prev, name, il); - ggml_tensor * score_prev = ggml_get_rows(ctx0, score_state, prev_idxs); - score_prev = ggml_cont(ctx0, ggml_view_2d(ctx0, score_prev, n_embd_head, ratio*n_blocks, score_prev->nb[1], 0)); + ggml_tensor * score_prev = ggml_cont(ctx0, + ggml_view_2d(ctx0, score_rows, n_embd_head, n_read, score_rows->nb[1], 0)); score_prev = ggml_reshape_3d(ctx0, score_prev, n_embd_head, ratio, n_blocks); cb(score_prev, name, il); - ggml_tensor * kv_cur = ggml_get_rows(ctx0, kv_state, cur_idxs); - kv_cur = ggml_cont(ctx0, ggml_view_2d(ctx0, kv_cur, n_embd_head, ratio*n_blocks, kv_cur->nb[1], - ggml_row_size(kv_cur->type, n_embd_head))); + ggml_tensor * kv_cur = ggml_cont(ctx0, + ggml_view_2d(ctx0, kv_rows, n_embd_head, n_read, kv_rows->nb[1], + n_read*kv_rows->nb[1] + ggml_row_size(kv_rows->type, n_embd_head))); kv_cur = ggml_reshape_3d(ctx0, kv_cur, n_embd_head, ratio, n_blocks); - ggml_tensor * score_cur = ggml_get_rows(ctx0, score_state, cur_idxs); - score_cur = ggml_cont(ctx0, ggml_view_2d(ctx0, score_cur, n_embd_head, ratio*n_blocks, score_cur->nb[1], - ggml_row_size(score_cur->type, n_embd_head))); + ggml_tensor * score_cur = ggml_cont(ctx0, + ggml_view_2d(ctx0, score_rows, n_embd_head, n_read, score_rows->nb[1], + n_read*score_rows->nb[1] + ggml_row_size(score_rows->type, n_embd_head))); score_cur = ggml_reshape_3d(ctx0, score_cur, n_embd_head, ratio, n_blocks); ggml_tensor * values = ggml_concat(ctx0, kv_prev, kv_cur, 1);