From 763c40f0942ddeb9a6308235115e50baa1566a49 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Thu, 25 Jun 2026 20:12:26 -0500 Subject: [PATCH 01/11] SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing --- ggml/src/ggml-sycl/fattn-mkl.cpp | 567 +++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/fattn.cpp | 14 + ggml/src/ggml-sycl/fattn.hpp | 2 + 3 files changed, 583 insertions(+) create mode 100644 ggml/src/ggml-sycl/fattn-mkl.cpp diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp new file mode 100644 index 000000000000..a145082286fc --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -0,0 +1,567 @@ +// +// MIT license +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: MIT +// + +// +// Flash attention via oneMKL GEMM (XMX-accelerated). +// Uses column_major::gemm for Q*K^T and S*V matmuls +// with an online softmax SYCL kernel. +// +// All GQA query heads sharing a KV head are batched into single +// GEMM calls, amortizing MKL launch overhead across K and V reuse. +// + +#include "common.hpp" +#include "fattn-common.hpp" +#include "fattn-buffers.hpp" +#include "convert.hpp" +#include "fattn.hpp" + +#include +#include +#include +#include + +#define MKL_FA_CHUNK_SIZE_KV 8192 + +using oneapi::mkl::transpose; +using oneapi::mkl::blas::column_major::gemm; + +// --------------------------------------------------------------------------- +// Helpers +// --------------------------------------------------------------------------- + +// Pack all GQA Q heads for one KV head into fp16, applying q_scale. +// One ND-range over gqa * n_queries * DKQ elements — each work-item +// computes its GQA group from the global index. +static void mkl_fa_pack_q_fp16( + dpct::queue_ptr stream, + sycl::half * __restrict dst, + const float * __restrict q_src, + int n_queries, int n_query_rows, int DKQ, + int gqa_ratio, int kvh_base_head, + float q_scale, int64_t wg_size) { + + const int64_t q_per_head = (int64_t)n_queries * DKQ; + const int64_t total = (int64_t)n_query_rows * DKQ; + const int64_t wg = ((total + wg_size - 1) / wg_size) * wg_size; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int64_t e = item.get_global_id(0); + if (e >= total) return; + + int iqg = (int)(e / q_per_head); + int64_t off = e - (int64_t)iqg * q_per_head; + int iqh = kvh_base_head + iqg; + + dst[e] = sycl::half( + q_src[(int64_t)iqh * q_per_head + off] * q_scale); + }); + }); +} + +// Zero-initialize the online softmax state arrays. +// KQ_max → -inf, KQ_sum → 0, VKQ_accum → 0. +// Merged into one kernel to avoid per-array launch overhead. +static void mkl_fa_init_softmax_state( + dpct::queue_ptr stream, + float * kmax, float * ksum, float * vacc, + int n_query_rows, int DV, int64_t wg_size) { + + const float neg_inf = -1e30f; + const int64_t n_maxsum = n_query_rows; + const int64_t n_vacc = (int64_t)n_query_rows * DV; + const int64_t total = (n_vacc > n_maxsum) ? n_vacc : n_maxsum; + const int64_t wg = ((total + wg_size - 1) / wg_size) * wg_size; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int64_t i = item.get_global_id(0); + if (i < n_maxsum) { + kmax[i] = neg_inf; + ksum[i] = 0.0f; + } + if (i < n_vacc) { + vacc[i] = 0.0f; + } + }); + }); +} + +// Online softmax over one KV chunk for all GQA query rows. +// For each row: find local max → rescale previous VKQ_accum → +// compute exp(s - max) → write S_f16 → update running max/sum. +// Runs per-row independently (rows don't interact). +static void mkl_fa_online_softmax_chunk( + dpct::queue_ptr stream, + float * __restrict KQ_f32, + sycl::half * __restrict S_f16, + float * __restrict KQ_max, + float * __restrict KQ_sum, + float * __restrict VKQ_accum, + int n_query_rows, int n_queries, int DV, + int chunk_size, int chunk_start, + int kvh_head, int gqa_ratio, + const sycl::half * mask_data, int64_t mask_head_stride, + int64_t mask_row_stride, int mask_n_heads, + float logit_softcap, int64_t wg_size) { + + const int64_t wg = ((n_query_rows + wg_size - 1) / wg_size) * wg_size; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int jc = item.get_global_id(0); + if (jc >= n_query_rows) return; + + const int gqa_group = jc / n_queries; + const int q_row = jc % n_queries; + + const float * __restrict KQ_row = KQ_f32 + + jc * (int64_t)chunk_size; + float * __restrict vkq = VKQ_accum + + jc * (int64_t)DV; + + // Per-GQA-group mask resolution + const sycl::half * mask_h = nullptr; + int64_t m_stride = 0; + if (mask_data) { + int m_head = (mask_n_heads > 1) + ? (kvh_head + gqa_group) : 0; + mask_h = mask_data + (int64_t)m_head * mask_head_stride; + m_stride = mask_row_stride; + } + + // Row-wise local maximum + float local_max = -1e30f; + for (int i = 0; i < chunk_size; i++) { + float s = KQ_row[i]; + if (mask_h) { + s += (float)mask_h[q_row * m_stride + + (chunk_start + i)]; + } + if (s > local_max) local_max = s; + } + + // Rescale previous accumulator by exp(old_max - new_max) + float old_max = KQ_max[jc]; + float new_max = (old_max > local_max) ? old_max : local_max; + float rescale = (old_max < -1e29f) ? 1.0f + : sycl::native::exp(old_max - new_max); + + for (int v = 0; v < DV; v++) { + vkq[v] *= rescale; + } + + // Softmax and write S_f16 + float local_sum = 0.0f; + sycl::half * __restrict S_row = S_f16 + + jc * (int64_t)chunk_size; + + for (int i = 0; i < chunk_size; i++) { + float s = KQ_row[i]; + if (mask_h) { + s += (float)mask_h[q_row * m_stride + + (chunk_start + i)]; + } + if (logit_softcap != 0.0f) { + s = logit_softcap * sycl::tanh(s); + } + float val = sycl::native::exp(s - new_max); + S_row[i] = sycl::half(val); + local_sum += val; + } + + KQ_sum[jc] = KQ_sum[jc] * rescale + local_sum; + KQ_max[jc] = new_max; + }); + }); +} + +// Write one GQA group's normalized output to its destination head. +static void mkl_fa_normalize_head( + dpct::queue_ptr stream, + float * __restrict dst_batch, + const float * __restrict VKQ_accum, + const float * __restrict KQ_sum, + int iqh, int n_queries, int DV, + int64_t src_offset, int64_t dst_row_stride, + int64_t wg_size) { + + const int64_t wg = ((n_queries + wg_size - 1) / wg_size) * wg_size; + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int jc = item.get_global_id(0); + if (jc >= n_queries) return; + + int ksum_idx = (int)(src_offset / DV) + jc; + float inv_sum = 1.0f / KQ_sum[ksum_idx]; + const float * __restrict src = VKQ_accum + + src_offset + jc * (int64_t)DV; + float * __restrict dst_row = dst_batch + + iqh * n_queries * (int64_t)dst_row_stride + + jc * dst_row_stride; + + for (int v = 0; v < DV; v++) { + dst_row[v] = src[v] * inv_sum; + } + }); + }); +} + +// --------------------------------------------------------------------------- +// MKL Flash Attention orchestrator +// +// Pipeline: dequantize K/V → for each KV head: +// pack GQA Q heads → MKL GEMM KQ → online softmax → +// MKL GEMM VKQ → accumulate → normalize → scatter to dst +// --------------------------------------------------------------------------- +void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + + 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]; + // dst->src[4] = sinks (not yet supported) + ggml_tensor * KQV = dst; + + GGML_ASSERT(Q->type == GGML_TYPE_F32); + GGML_ASSERT(KQV->type == GGML_TYPE_F32); + + // --- Op params --- + float scale = 1.0f, max_bias = 0.0f, logit_softcap = 0.0f; + memcpy(&scale, (const float *)KQV->op_params + 0, sizeof(float)); + memcpy(&max_bias, (const float *)KQV->op_params + 1, sizeof(float)); + memcpy(&logit_softcap, (const float *)KQV->op_params + 2, sizeof(float)); + + const bool use_logit_softcap = (logit_softcap != 0.0f); + const float q_scale = use_logit_softcap ? (scale / logit_softcap) : scale; + + // --- Dimensions --- + const int DKQ = (int)K->ne[0]; + const int DV = (int)V->ne[0]; + const int n_queries = (int)Q->ne[1]; + const int n_q_heads = (int)Q->ne[2]; + const int n_kv_heads = (int)K->ne[2]; + const int n_batch = (int)Q->ne[3]; + const int n_kv = (int)K->ne[1]; + const int gqa_ratio = n_q_heads / n_kv_heads; + const int n_query_rows = n_queries * gqa_ratio; + + GGML_ASSERT(n_q_heads % n_kv_heads == 0); + GGML_ASSERT(max_bias == 0.0f); // ALiBi not supported + GGML_ASSERT(Q->ne[3] == K->ne[3] || K->ne[3] == 1); + + const int chunk_size = std::min(MKL_FA_CHUNK_SIZE_KV, n_kv); + const int64_t wg_size = 256; + + // --- Debug output (gated by MKL_FA_DEBUG=1) --- + static int mkl_call_count = 0; + mkl_call_count++; + static int mkl_debug = -1; + if (mkl_debug < 0) { + const char * e = getenv("MKL_FA_DEBUG"); + mkl_debug = (e && e[0] == '1') ? 1 : 0; + } + const bool do_print = (mkl_debug == 1); + + if (do_print) { + fprintf(stderr, "[MKL-FA] #%d D=%d DV=%d n_q=%d n_kv=%d " + "n_qh=%d n_kvh=%d gqa=%d batch=%d K=%s V=%s " + "chunk=%d KQ_buf=%.1fMB\n", + mkl_call_count, DKQ, DV, n_queries, n_kv, + n_q_heads, n_kv_heads, gqa_ratio, n_batch, + ggml_type_name(K->type), ggml_type_name(V->type), + chunk_size, + (double)((int64_t)n_query_rows * chunk_size * sizeof(float)) + / (1024.0 * 1024.0)); + fflush(stderr); + } + + // --- Stream and allocators --- + dpct::queue_ptr stream = ctx.stream(); + stream->wait(); // drain pending work before MKL operations + + ggml_sycl_pool & pool = ctx.pool(); + ggml_sycl_fattn_kv_buffers & fbuf = ctx.fattn_buffers(); + + // Timing helpers (local to this function) +#define MKL_TAKE_TIME(t0) auto t0 = std::chrono::steady_clock::now() +#define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \ + (std::chrono::steady_clock::now() - (t0)).count() + + int64_t gemm_kq_time_us = 0; + int64_t gemm_vkq_time_us = 0; + int64_t softmax_time_us = 0; + int64_t dequant_time_us = 0; + + MKL_TAKE_TIME(t_deq); + + // --- Dequantize K, V to fp16 --- + ggml_sycl_fattn_alloc K_f16(fbuf.K); + ggml_sycl_fattn_alloc V_f16(fbuf.V); + + const char * K_data = (const char *)K->data; + const char * V_data = (const char *)V->data; + + // Strides in the fp16 dequant buffer (row-major within each head) + size_t k_row_stride = DKQ * sizeof(sycl::half); + size_t k_head_stride = (size_t)n_kv * k_row_stride; + size_t v_row_stride = DV * sizeof(sycl::half); + size_t v_head_stride = (size_t)n_kv * v_row_stride; + + const bool V_is_K_view = V->view_src + && (V->view_src == K || (V->view_src == K->view_src + && V->view_offs == K->view_offs)); + + if (K->type != GGML_TYPE_F16) { + K_f16.alloc(ggml_nelements(K)); + if (ggml_is_contiguously_allocated(K)) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst); + to_fp16(K->data, K_f16.ptr, ggml_nelements(K), stream); + } else { + to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(K->type); + const int64_t s01 = K->nb[1] / ggml_type_size(K->type); + const int64_t s02 = K->nb[2] / ggml_type_size(K->type); + const int64_t s03 = K->nb[3] / ggml_type_size(K->type); + to_fp16(K->data, K_f16.ptr, + K->ne[0], K->ne[1], K->ne[2], K->ne[3], + s01, s02, s03, stream); + } + K_data = (char *)K_f16.ptr; + } else { + k_row_stride = K->nb[1]; + k_head_stride = K->nb[2]; + } + + if (V->type != GGML_TYPE_F16) { + if (V_is_K_view) { + V_data = K_data; + v_row_stride = k_row_stride; + v_head_stride = k_head_stride; + } else { + V_f16.alloc(ggml_nelements(V)); + if (ggml_is_contiguously_allocated(V)) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst); + to_fp16(V->data, V_f16.ptr, ggml_nelements(V), stream); + } else { + to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(V->type); + const int64_t s01 = V->nb[1] / ggml_type_size(V->type); + const int64_t s02 = V->nb[2] / ggml_type_size(V->type); + const int64_t s03 = V->nb[3] / ggml_type_size(V->type); + to_fp16(V->data, V_f16.ptr, + V->ne[0], V->ne[1], V->ne[2], V->ne[3], + s01, s02, s03, stream); + } + V_data = (char *)V_f16.ptr; + } + } else { + if (V_is_K_view) { + V_data = K_data; + v_row_stride = k_row_stride; + v_head_stride = k_head_stride; + } else { + v_row_stride = V->nb[1]; + v_head_stride = V->nb[2]; + } + } + + stream->wait(); // dequant complete + MKL_ACCUM(dequant_time_us, t_deq); + + // --- Resolve mask pointers once (valid for all chunks) --- + const sycl::half * mask_data = nullptr; + int64_t mask_head_stride = 0; + int64_t mask_row_stride = 0; + int mask_n_heads = 0; + if (mask) { + // Mask is resolved per-batch below; store strides here + mask_head_stride = mask->nb[2] / sizeof(sycl::half); + mask_row_stride = mask->nb[1] / sizeof(sycl::half); + mask_n_heads = (int)mask->ne[2]; + } + + // --- Allocate intermediates (sized for all GQA query rows) --- + ggml_sycl_pool_alloc KQ_f32(pool); + ggml_sycl_pool_alloc S_f16(pool); + ggml_sycl_pool_alloc VKQ_chunk(pool); + ggml_sycl_pool_alloc VKQ_accum(pool); + ggml_sycl_pool_alloc KQ_max(pool); + ggml_sycl_pool_alloc KQ_sum(pool); + ggml_sycl_pool_alloc Q_head_f16(pool); + + KQ_f32.alloc((size_t)n_query_rows * chunk_size); + S_f16.alloc((size_t)n_query_rows * chunk_size); + VKQ_chunk.alloc((size_t)n_query_rows * DV); + VKQ_accum.alloc((size_t)n_query_rows * DV); + KQ_max.alloc(n_query_rows); + KQ_sum.alloc(n_query_rows); + Q_head_f16.alloc((size_t)n_query_rows * DKQ); + + // Extract raw pointers — ggml_sycl_pool_alloc is not device-copyable + sycl::half * Q_head_f16_ptr = Q_head_f16.ptr; + float * KQ_f32_ptr = KQ_f32.ptr; + sycl::half * S_f16_ptr = S_f16.ptr; + float * VKQ_chunk_ptr = VKQ_chunk.ptr; + float * VKQ_accum_ptr = VKQ_accum.ptr; + float * KQ_max_ptr = KQ_max.ptr; + float * KQ_sum_ptr = KQ_sum.ptr; + + const int64_t dst_row_stride = KQV->nb[1] / sizeof(float); + + const float alpha = 1.0f; + const float beta = 0.0f; + + for (int ib = 0; ib < n_batch; ib++) { + const float * Q_batch = (const float *)Q->data + + ib * (Q->nb[3] / sizeof(float)); + float * dst_batch = (float *)KQV->data + + ib * (KQV->nb[3] / sizeof(float)); + + // Per-batch mask base + const sycl::half * mask_batch = nullptr; + if (mask) { + int m_batch = (mask->ne[3] > 1) ? ib : 0; + mask_batch = (const sycl::half *)mask->data + + m_batch * (mask->nb[3] / sizeof(sycl::half)); + } + + for (int ikvh = 0; ikvh < n_kv_heads; ikvh++) { + int kvh_base_head = ikvh * gqa_ratio; + + // 1. Pack all GQA Q heads into fp16 + mkl_fa_pack_q_fp16(stream, + Q_head_f16_ptr, Q_batch, + n_queries, n_query_rows, DKQ, + gqa_ratio, kvh_base_head, + q_scale, wg_size); + + // 2. Initialize softmax state + mkl_fa_init_softmax_state(stream, + KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr, + n_query_rows, DV, wg_size); + + // Sync before MKL GEMM (MKL may use an internal queue) + stream->wait(); + + // K/V base for this KV head + const sycl::half * K_head = (const sycl::half *)K_data + + ikvh * (k_head_stride / sizeof(sycl::half)); + const sycl::half * V_head = (const sycl::half *)V_data + + ikvh * (v_head_stride / sizeof(sycl::half)); + + // 3. Chunked KV loop + for (int chunk_start = 0; chunk_start < n_kv; chunk_start += chunk_size) { + int this_chunk = std::min(chunk_size, n_kv - chunk_start); + const sycl::half * K_chunk = K_head + + (int64_t)chunk_start * DKQ; + const sycl::half * V_chunk = V_head + + (int64_t)chunk_start * DV; + + // 3a. GEMM: KQ = Q_batched × K_chunk^T + { + MKL_TAKE_TIME(t0); + sycl::event ev = gemm(*stream, + transpose::trans, transpose::nontrans, + this_chunk, n_query_rows, DKQ, + alpha, + K_chunk, DKQ, + Q_head_f16_ptr, DKQ, + beta, + KQ_f32_ptr, this_chunk); + stream->wait(); + try { ev.wait_and_throw(); } catch (sycl::exception & e) { + fprintf(stderr, "[MKL-FA] GEMM KQ: %s\n", e.what()); + GGML_ABORT("MKL GEMM KQ failed"); + } + MKL_ACCUM(gemm_kq_time_us, t0); + } + + // 3b. Online softmax over this chunk + { + MKL_TAKE_TIME(t0); + mkl_fa_online_softmax_chunk(stream, + KQ_f32_ptr, S_f16_ptr, + KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr, + n_query_rows, n_queries, DV, + this_chunk, chunk_start, + kvh_base_head, gqa_ratio, + mask_batch, mask_head_stride, + mask_row_stride, mask_n_heads, + logit_softcap, wg_size); + stream->wait(); // S_f16 must be ready for GEMM + MKL_ACCUM(softmax_time_us, t0); + } + + // 3c. GEMM: VKQ_chunk = S × V_chunk + { + MKL_TAKE_TIME(t0); + sycl::event ev = gemm(*stream, + transpose::nontrans, transpose::nontrans, + DV, n_query_rows, this_chunk, + alpha, + V_chunk, DV, + S_f16_ptr, this_chunk, + beta, + VKQ_chunk_ptr, DV); + stream->wait(); + try { ev.wait_and_throw(); } catch (sycl::exception & e) { + fprintf(stderr, "[MKL-FA] GEMM VKQ: %s\n", e.what()); + GGML_ABORT("MKL GEMM VKQ failed"); + } + MKL_ACCUM(gemm_vkq_time_us, t0); + } + + // 3d. VKQ_accum += VKQ_chunk + { + const int64_t n_total = (int64_t)n_query_rows * DV; + const int64_t wg = ((n_total + wg_size - 1) / wg_size) + * wg_size; + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int64_t i = item.get_global_id(0); + if (i < n_total) { + VKQ_accum_ptr[i] += VKQ_chunk_ptr[i]; + } + }); + }); + // In-order queue: accumulate finishes before next + // chunk's GEMM or final normalize. + } + } + + // 4. Normalize and scatter each GQA head to dst + for (int iqg = 0; iqg < gqa_ratio; iqg++) { + int iqh = kvh_base_head + iqg; + int64_t src_offset = (int64_t)iqg * n_queries * DV; + mkl_fa_normalize_head(stream, + dst_batch, VKQ_accum_ptr, KQ_sum_ptr, + iqh, n_queries, DV, + src_offset, dst_row_stride, wg_size); + } + // In-order queue: normalize writes ordered before next KV head. + } + } + +#undef MKL_TAKE_TIME +#undef MKL_ACCUM + + if (do_print) { + fprintf(stderr, "[MKL-FA] #%d n_kv=%d n_q=%d time_us: " + "dequant=%lld GEMM_KQ=%lld softmax=%lld GEMM_VKQ=%lld\n", + mkl_call_count, n_kv, n_queries, + (long long)dequant_time_us, + (long long)gemm_kq_time_us, + (long long)softmax_time_us, + (long long)gemm_vkq_time_us); + fflush(stderr); + } +} diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 7c6e6112fdcd..7869897a9936 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -97,6 +97,7 @@ enum best_fattn_kernel { BEST_FATTN_KERNEL_NONE = 0, BEST_FATTN_KERNEL_VEC = 100, BEST_FATTN_KERNEL_TILE = 200, + BEST_FATTN_KERNEL_MKL = 300, }; static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const ggml_tensor * dst) { @@ -121,6 +122,16 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const memcpy(&max_bias, (const float *) KQV->op_params + 1, sizeof(float)); bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; + + // MKL path: XMX-accelerated GEMM for prompt processing with quantized KV cache. + // Activates at n_kv >= 1024 (matching the --batch-size) to cover the full prompt. + // Each call takes ~13ms at 8K — GEMMs ~3ms, softmax kernel ~9ms. + // FIXME: MKL GEMM calls are incompatible with SYCL graph capture replay. + if (Q->ne[1] >= 128 && K->ne[1] >= 1024 + && (ggml_is_quantized(K->type) || ggml_is_quantized(V->type))) { + return BEST_FATTN_KERNEL_MKL; + } + for (const ggml_tensor * t : {Q, K, V, mask}) { if (t == nullptr || ggml_is_quantized(t->type)) { continue; @@ -219,6 +230,9 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst case BEST_FATTN_KERNEL_VEC: ggml_sycl_flash_attn_ext_vec(ctx, dst); break; + case BEST_FATTN_KERNEL_MKL: + ggml_sycl_flash_attn_ext_mkl(ctx, dst); + break; } } diff --git a/ggml/src/ggml-sycl/fattn.hpp b/ggml/src/ggml-sycl/fattn.hpp index f2a8ffc97dee..c093970a3fed 100644 --- a/ggml/src/ggml-sycl/fattn.hpp +++ b/ggml/src/ggml-sycl/fattn.hpp @@ -19,4 +19,6 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst); +void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + #endif // GGML_SYCL_FATTN_HPP From ab0048e1d1d93986f50c87168c62e7e0bf35af59 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Sun, 28 Jun 2026 02:16:36 -0500 Subject: [PATCH 02/11] fattn-mkl: fix interleaved dst layout in normalize kernel MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro --- ggml/src/ggml-sycl/fattn-mkl.cpp | 295 ++++++++++++++++++++----------- ggml/src/ggml-sycl/fattn.cpp | 93 +++++++++- 2 files changed, 287 insertions(+), 101 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index a145082286fc..25cfab8cb1c8 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -34,34 +34,45 @@ using oneapi::mkl::blas::column_major::gemm; // --------------------------------------------------------------------------- // Pack all GQA Q heads for one KV head into fp16, applying q_scale. -// One ND-range over gqa * n_queries * DKQ elements — each work-item -// computes its GQA group from the global index. +// Launches one kernel per GQA group — each kernel copies exactly +// n_queries * DKQ elements using the per-group dst offset and +// per-head source stride. static void mkl_fa_pack_q_fp16( dpct::queue_ptr stream, sycl::half * __restrict dst, const float * __restrict q_src, int n_queries, int n_query_rows, int DKQ, int gqa_ratio, int kvh_base_head, - float q_scale, int64_t wg_size) { + float q_scale, int64_t q_row_stride, int64_t q_head_stride, + int64_t wg_size) { - const int64_t q_per_head = (int64_t)n_queries * DKQ; - const int64_t total = (int64_t)n_query_rows * DKQ; - const int64_t wg = ((total + wg_size - 1) / wg_size) * wg_size; + for (int iqg = 0; iqg < gqa_ratio; iqg++) { + int iqh = kvh_base_head + iqg; + sycl::half * dst_g = dst + (int64_t)iqg * n_queries * DKQ; - stream->submit([&](sycl::handler & cgh) { - cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), - [=](sycl::nd_item<1> item) { - int64_t e = item.get_global_id(0); - if (e >= total) return; + const int64_t n_elem = (int64_t)n_queries * DKQ; + const int64_t wg = ((n_elem + wg_size - 1) / wg_size) * wg_size; - int iqg = (int)(e / q_per_head); - int64_t off = e - (int64_t)iqg * q_per_head; - int iqh = kvh_base_head + iqg; + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(wg, wg_size), + [=](sycl::nd_item<1> item) { + int64_t e = item.get_global_id(0); + if (e >= n_elem) return; - dst[e] = sycl::half( - q_src[(int64_t)iqh * q_per_head + off] * q_scale); - }); - }); + int64_t q = e / DKQ; + int64_t d = e - q * DKQ; + + // Stride-aware source offset: handles permuted, + // sliced, or contiguous Q tensor layouts. + int64_t src_off = d + + q * q_row_stride + + (int64_t)iqh * q_head_stride; + + dst_g[e] = sycl::half( + q_src[src_off] * q_scale); + }); + }); + } } // Zero-initialize the online softmax state arrays. @@ -96,7 +107,6 @@ static void mkl_fa_init_softmax_state( // Online softmax over one KV chunk for all GQA query rows. // For each row: find local max → rescale previous VKQ_accum → // compute exp(s - max) → write S_f16 → update running max/sum. -// Runs per-row independently (rows don't interact). static void mkl_fa_online_softmax_chunk( dpct::queue_ptr stream, float * __restrict KQ_f32, @@ -127,7 +137,6 @@ static void mkl_fa_online_softmax_chunk( float * __restrict vkq = VKQ_accum + jc * (int64_t)DV; - // Per-GQA-group mask resolution const sycl::half * mask_h = nullptr; int64_t m_stride = 0; if (mask_data) { @@ -137,10 +146,13 @@ static void mkl_fa_online_softmax_chunk( m_stride = mask_row_stride; } - // Row-wise local maximum + // Row-wise local maximum (softcap before mask) float local_max = -1e30f; for (int i = 0; i < chunk_size; i++) { float s = KQ_row[i]; + if (logit_softcap != 0.0f) { + s = logit_softcap * sycl::tanh(s); + } if (mask_h) { s += (float)mask_h[q_row * m_stride + (chunk_start + i)]; @@ -165,13 +177,13 @@ static void mkl_fa_online_softmax_chunk( for (int i = 0; i < chunk_size; i++) { float s = KQ_row[i]; + if (logit_softcap != 0.0f) { + s = logit_softcap * sycl::tanh(s); + } if (mask_h) { s += (float)mask_h[q_row * m_stride + (chunk_start + i)]; } - if (logit_softcap != 0.0f) { - s = logit_softcap * sycl::tanh(s); - } float val = sycl::native::exp(s - new_max); S_row[i] = sycl::half(val); local_sum += val; @@ -189,9 +201,8 @@ static void mkl_fa_normalize_head( float * __restrict dst_batch, const float * __restrict VKQ_accum, const float * __restrict KQ_sum, - int iqh, int n_queries, int DV, - int64_t src_offset, int64_t dst_row_stride, - int64_t wg_size) { + int iqh, int n_queries, int DV, int n_q_heads, + int64_t src_offset, int64_t wg_size) { const int64_t wg = ((n_queries + wg_size - 1) / wg_size) * wg_size; @@ -205,9 +216,11 @@ static void mkl_fa_normalize_head( float inv_sum = 1.0f / KQ_sum[ksum_idx]; const float * __restrict src = VKQ_accum + src_offset + jc * (int64_t)DV; + // Interleaved dst layout (matching TILE): + // rows alternate between heads, then increment query. + // offset = (query * n_q_heads + head) * DV float * __restrict dst_row = dst_batch - + iqh * n_queries * (int64_t)dst_row_stride - + jc * dst_row_stride; + + ((int64_t)jc * n_q_heads + iqh) * (int64_t)DV; for (int v = 0; v < DV; v++) { dst_row[v] = src[v] * inv_sum; @@ -229,7 +242,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; const ggml_tensor * mask = dst->src[3]; - // dst->src[4] = sinks (not yet supported) ggml_tensor * KQV = dst; GGML_ASSERT(Q->type == GGML_TYPE_F32); @@ -241,8 +253,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * memcpy(&max_bias, (const float *)KQV->op_params + 1, sizeof(float)); memcpy(&logit_softcap, (const float *)KQV->op_params + 2, sizeof(float)); - const bool use_logit_softcap = (logit_softcap != 0.0f); - const float q_scale = use_logit_softcap ? (scale / logit_softcap) : scale; + const float q_scale = scale; // --- Dimensions --- const int DKQ = (int)K->ne[0]; @@ -272,27 +283,48 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * } const bool do_print = (mkl_debug == 1); + const int64_t q_row_stride = Q->nb[1] / sizeof(float); + const int64_t q_head_stride = Q->nb[2] / sizeof(float); + + const bool V_is_K_view = V->view_src + && (V->view_src == K || (V->view_src == K->view_src + && V->view_offs == K->view_offs)); + + // Early interleaved detection for debug output. + // True interleaved detection happens after dequant (nb12_fp16 == nb11_fp16), + // but we can pre-detect on the original tensor strides. + const bool k_early_interleaved = + ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]); + const bool v_early_interleaved = + !V_is_K_view && ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]); + if (do_print) { fprintf(stderr, "[MKL-FA] #%d D=%d DV=%d n_q=%d n_kv=%d " "n_qh=%d n_kvh=%d gqa=%d batch=%d K=%s V=%s " - "chunk=%d KQ_buf=%.1fMB\n", + "chunk=%d buf=%.1fMB%s%s\n", mkl_call_count, DKQ, DV, n_queries, n_kv, n_q_heads, n_kv_heads, gqa_ratio, n_batch, ggml_type_name(K->type), ggml_type_name(V->type), chunk_size, (double)((int64_t)n_query_rows * chunk_size * sizeof(float)) - / (1024.0 * 1024.0)); + / (1024.0 * 1024.0), + k_early_interleaved ? " K_ILV" : "", + v_early_interleaved ? " V_ILV" : ""); + fprintf(stderr, "[MKL-FA] #%d Q-nb1=%lld Q-nb2=%lld " + "q_rs=%lld q_hs=%lld dst_rs=%lld dst_hs=%lld\n", + mkl_call_count, + (long long)Q->nb[1], (long long)Q->nb[2], + (long long)q_row_stride, (long long)q_head_stride, + (long long)(KQV->nb[1] / sizeof(float)), + (long long)(KQV->nb[2] / sizeof(float))); fflush(stderr); } // --- Stream and allocators --- dpct::queue_ptr stream = ctx.stream(); - stream->wait(); // drain pending work before MKL operations - ggml_sycl_pool & pool = ctx.pool(); ggml_sycl_fattn_kv_buffers & fbuf = ctx.fattn_buffers(); - // Timing helpers (local to this function) #define MKL_TAKE_TIME(t0) auto t0 = std::chrono::steady_clock::now() #define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \ (std::chrono::steady_clock::now() - (t0)).count() @@ -305,90 +337,152 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * MKL_TAKE_TIME(t_deq); // --- Dequantize K, V to fp16 --- + // + // Output is ALWAYS dense row-major (contiguous rows per head). + // For interleaved source layouts (Gemma: nb[2]==nb[1], heads + // interleave at the row level), we reconstruct the PHYSICAL + // source strides so to_fp16_nc writes the correct data into a + // dense buffer. The logical view strides (nb##/ts) are wrong + // for interleaved layouts — e.g. s01=8 aliases Row1/Head0 to + // the same blocks as Row0/Head1. + // + // Detection: ne[1]*nb[1] != nb[2] means heads are interleaved. ggml_sycl_fattn_alloc K_f16(fbuf.K); ggml_sycl_fattn_alloc V_f16(fbuf.V); const char * K_data = (const char *)K->data; + size_t nb11 = K->nb[1]; + size_t nb12 = K->nb[2]; + size_t nb13 = K->nb[3]; + const char * V_data = (const char *)V->data; + size_t nb21 = V->nb[1]; + size_t nb22 = V->nb[2]; + size_t nb23 = V->nb[3]; - // Strides in the fp16 dequant buffer (row-major within each head) - size_t k_row_stride = DKQ * sizeof(sycl::half); - size_t k_head_stride = (size_t)n_kv * k_row_stride; - size_t v_row_stride = DV * sizeof(sycl::half); - size_t v_head_stride = (size_t)n_kv * v_row_stride; - - const bool V_is_K_view = V->view_src - && (V->view_src == K || (V->view_src == K->view_src - && V->view_offs == K->view_offs)); + const bool k_interleaved = + ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]) && K->ne[2] > 1; + const bool v_interleaved = + ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1; if (K->type != GGML_TYPE_F16) { + const size_t bs = ggml_blck_size(K->type); + const size_t ts = ggml_type_size(K->type); + K_f16.alloc(ggml_nelements(K)); - if (ggml_is_contiguously_allocated(K)) { + + if (ggml_is_contiguously_allocated(K) && !k_interleaved) { + // Dense source — flat to_fp16 preserves layout to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst); - to_fp16(K->data, K_f16.ptr, ggml_nelements(K), stream); + to_fp16(K_data, K_f16.ptr, ggml_nelements(K), stream); + + nb11 = nb11 * bs * sizeof(sycl::half) / ts; + nb12 = nb12 * bs * sizeof(sycl::half) / ts; + nb13 = nb13 * bs * sizeof(sycl::half) / ts; } else { + // Interleaved or non-contiguous source — use to_fp16_nc + // with PHYSICAL source strides so the output is dense. to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(K->type); - const int64_t s01 = K->nb[1] / ggml_type_size(K->type); - const int64_t s02 = K->nb[2] / ggml_type_size(K->type); - const int64_t s03 = K->nb[3] / ggml_type_size(K->type); - to_fp16(K->data, K_f16.ptr, + + int64_t s01, s02, s03; + if (k_interleaved) { + // Interleaved: heads alternate at the row level. + // Row r of head h starts at block + // h * blk_per_row + r * n_kv_heads * blk_per_row. + const int64_t blk_per_row = (int64_t)K->ne[0] / bs; + s01 = (int64_t)n_kv_heads * blk_per_row; // row stride + s02 = blk_per_row; // head stride + s03 = (int64_t)K->ne[1] * s01; // batch stride + } else { + s01 = nb11 / ts; + s02 = nb12 / ts; + s03 = nb13 / ts; + } + + to_fp16(K_data, K_f16.ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, stream); + + // Output is dense row-major + nb11 = K->ne[0] * sizeof(sycl::half); + nb12 = K->ne[1] * nb11; + nb13 = K->ne[2] * nb12; } - K_data = (char *)K_f16.ptr; - } else { - k_row_stride = K->nb[1]; - k_head_stride = K->nb[2]; + K_data = (char *) K_f16.ptr; } if (V->type != GGML_TYPE_F16) { if (V_is_K_view) { - V_data = K_data; - v_row_stride = k_row_stride; - v_head_stride = k_head_stride; + V_data = K_data; + nb21 = nb11; + nb22 = nb12; + nb23 = nb13; } else { + const size_t bs = ggml_blck_size(V->type); + const size_t ts = ggml_type_size(V->type); + V_f16.alloc(ggml_nelements(V)); - if (ggml_is_contiguously_allocated(V)) { + + if (ggml_is_contiguously_allocated(V) && !v_interleaved) { to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst); - to_fp16(V->data, V_f16.ptr, ggml_nelements(V), stream); + to_fp16(V_data, V_f16.ptr, ggml_nelements(V), stream); + V_data = (char *) V_f16.ptr; + + nb21 = nb21 * bs * sizeof(sycl::half) / ts; + nb22 = nb22 * bs * sizeof(sycl::half) / ts; + nb23 = nb23 * bs * sizeof(sycl::half) / ts; } else { to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(V->type); - const int64_t s01 = V->nb[1] / ggml_type_size(V->type); - const int64_t s02 = V->nb[2] / ggml_type_size(V->type); - const int64_t s03 = V->nb[3] / ggml_type_size(V->type); - to_fp16(V->data, V_f16.ptr, + + int64_t s01, s02, s03; + if (v_interleaved) { + 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 = nb21 / ts; + s02 = nb22 / ts; + s03 = nb23 / ts; + } + + to_fp16(V_data, V_f16.ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, stream); + + nb21 = V->ne[0] * sizeof(sycl::half); + nb22 = V->ne[1] * nb21; + nb23 = V->ne[2] * nb22; + + V_data = (char *) V_f16.ptr; } - V_data = (char *)V_f16.ptr; - } - } else { - if (V_is_K_view) { - V_data = K_data; - v_row_stride = k_row_stride; - v_head_stride = k_head_stride; - } else { - v_row_stride = V->nb[1]; - v_head_stride = V->nb[2]; } } - stream->wait(); // dequant complete MKL_ACCUM(dequant_time_us, t_deq); - // --- Resolve mask pointers once (valid for all chunks) --- + // All paths produce dense row-major fp16 — GEMM lda = D. + // Head pointers use dense strides (nb12 = ne[1]*ne[0]*sizeof(half)). + const int64_t nb12_fp16 = nb12 / sizeof(sycl::half); + const int64_t nb22_fp16 = nb22 / sizeof(sycl::half); + + // --- Resolve mask pointers --- const sycl::half * mask_data = nullptr; int64_t mask_head_stride = 0; int64_t mask_row_stride = 0; int mask_n_heads = 0; + if (mask) { - // Mask is resolved per-batch below; store strides here - mask_head_stride = mask->nb[2] / sizeof(sycl::half); - mask_row_stride = mask->nb[1] / sizeof(sycl::half); + // Use actual fp16 device size (2 bytes), NOT sizeof(sycl::half) + // which may be 4 on the host in oneAPI. + mask_head_stride = mask->nb[2] / 2; + mask_row_stride = mask->nb[1] / 2; mask_n_heads = (int)mask->ne[2]; } - // --- Allocate intermediates (sized for all GQA query rows) --- + // --- Allocate intermediates from pool --- + ggml_sycl_pool & pool = ctx.pool(); + ggml_sycl_pool_alloc KQ_f32(pool); ggml_sycl_pool_alloc S_f16(pool); ggml_sycl_pool_alloc VKQ_chunk(pool); @@ -405,7 +499,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * KQ_sum.alloc(n_query_rows); Q_head_f16.alloc((size_t)n_query_rows * DKQ); - // Extract raw pointers — ggml_sycl_pool_alloc is not device-copyable sycl::half * Q_head_f16_ptr = Q_head_f16.ptr; float * KQ_f32_ptr = KQ_f32.ptr; sycl::half * S_f16_ptr = S_f16.ptr; @@ -414,8 +507,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * float * KQ_max_ptr = KQ_max.ptr; float * KQ_sum_ptr = KQ_sum.ptr; - const int64_t dst_row_stride = KQV->nb[1] / sizeof(float); - const float alpha = 1.0f; const float beta = 0.0f; @@ -425,12 +516,11 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * float * dst_batch = (float *)KQV->data + ib * (KQV->nb[3] / sizeof(float)); - // Per-batch mask base const sycl::half * mask_batch = nullptr; if (mask) { int m_batch = (mask->ne[3] > 1) ? ib : 0; mask_batch = (const sycl::half *)mask->data - + m_batch * (mask->nb[3] / sizeof(sycl::half)); + + m_batch * (mask->nb[3] / 2); // 2 = actual fp16 device size } for (int ikvh = 0; ikvh < n_kv_heads; ikvh++) { @@ -441,7 +531,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * Q_head_f16_ptr, Q_batch, n_queries, n_query_rows, DKQ, gqa_ratio, kvh_base_head, - q_scale, wg_size); + q_scale, q_row_stride, q_head_stride, wg_size); // 2. Initialize softmax state mkl_fa_init_softmax_state(stream, @@ -451,11 +541,12 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * // Sync before MKL GEMM (MKL may use an internal queue) stream->wait(); - // K/V base for this KV head + // K/V base for this KV head. + // Dense row-major: nb12_fp16 = n_kv * DKQ (head stride in half-elements). const sycl::half * K_head = (const sycl::half *)K_data - + ikvh * (k_head_stride / sizeof(sycl::half)); + + ikvh * nb12_fp16; const sycl::half * V_head = (const sycl::half *)V_data - + ikvh * (v_head_stride / sizeof(sycl::half)); + + ikvh * nb22_fp16; // 3. Chunked KV loop for (int chunk_start = 0; chunk_start < n_kv; chunk_start += chunk_size) { @@ -483,7 +574,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * } MKL_ACCUM(gemm_kq_time_us, t0); } - // 3b. Online softmax over this chunk { MKL_TAKE_TIME(t0); @@ -518,7 +608,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * } MKL_ACCUM(gemm_vkq_time_us, t0); } - // 3d. VKQ_accum += VKQ_chunk { const int64_t n_total = (int64_t)n_query_rows * DV; @@ -533,8 +622,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * } }); }); - // In-order queue: accumulate finishes before next - // chunk's GEMM or final normalize. } } @@ -544,10 +631,9 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * int64_t src_offset = (int64_t)iqg * n_queries * DV; mkl_fa_normalize_head(stream, dst_batch, VKQ_accum_ptr, KQ_sum_ptr, - iqh, n_queries, DV, - src_offset, dst_row_stride, wg_size); + iqh, n_queries, DV, n_q_heads, + src_offset, wg_size); } - // In-order queue: normalize writes ordered before next KV head. } } @@ -555,13 +641,24 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * #undef MKL_ACCUM if (do_print) { + double total_mb = (double)( + (int64_t)n_query_rows * chunk_size * sizeof(float) + + (int64_t)n_query_rows * chunk_size * sizeof(sycl::half) + + (int64_t)n_query_rows * DV * sizeof(float) + + (int64_t)n_query_rows * DV * sizeof(float) + + (int64_t)n_query_rows * sizeof(float) + + (int64_t)n_query_rows * sizeof(float) + + (int64_t)n_query_rows * DKQ * sizeof(sycl::half) + ) / (1024.0 * 1024.0); fprintf(stderr, "[MKL-FA] #%d n_kv=%d n_q=%d time_us: " - "dequant=%lld GEMM_KQ=%lld softmax=%lld GEMM_VKQ=%lld\n", + "dequant=%lld GEMM_KQ=%lld softmax=%lld GEMM_VKQ=%lld " + "buf_mb=%.1f\n", mkl_call_count, n_kv, n_queries, (long long)dequant_time_us, (long long)gemm_kq_time_us, (long long)softmax_time_us, - (long long)gemm_vkq_time_us); + (long long)gemm_vkq_time_us, + total_mb); fflush(stderr); } } diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 7869897a9936..164b4c2510bc 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -100,6 +100,7 @@ enum best_fattn_kernel { BEST_FATTN_KERNEL_MKL = 300, }; + static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const ggml_tensor * dst) { GGML_UNUSED(device); #ifndef SYCL_FLASH_ATTN @@ -125,9 +126,14 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // MKL path: XMX-accelerated GEMM for prompt processing with quantized KV cache. // Activates at n_kv >= 1024 (matching the --batch-size) to cover the full prompt. - // Each call takes ~13ms at 8K — GEMMs ~3ms, softmax kernel ~9ms. // FIXME: MKL GEMM calls are incompatible with SYCL graph capture replay. - if (Q->ne[1] >= 128 && K->ne[1] >= 1024 + // Set MKL_FA_DISABLE=1 to force TILE/VEC path for A/B testing. + static int mkl_disable = -1; + if (mkl_disable < 0) { + const char * e = getenv("MKL_FA_DISABLE"); + mkl_disable = (e && e[0] == '1') ? 1 : 0; + } + if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024 && (ggml_is_quantized(K->type) || ggml_is_quantized(V->type))) { return BEST_FATTN_KERNEL_MKL; } @@ -221,6 +227,42 @@ 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); + + // n_kv watchdog: log when n_kv differs from the last FA call with + // the same D — helps detect cache-truncation issues. + static int nkv_debug = -1; + if (nkv_debug < 0) { + const char * e = getenv("MKL_FA_DEBUG"); + nkv_debug = (e && e[0] == '1') ? 1 : 0; + } + if (nkv_debug == 1) { + const ggml_tensor * K_dbg = dst->src[1]; + const ggml_tensor * V_dbg = dst->src[2]; + static int64_t last_nkv_d256 = 0, last_nkv_d512 = 0; + static int fa_call_seq = 0; + fa_call_seq++; + int64_t cur_nkv = K_dbg->ne[1]; + int Dk = (int)K_dbg->ne[0]; + const char * kname = "TILE"; + best_fattn_kernel k = ggml_sycl_get_best_fattn_kernel(ctx.device, dst); + if (k == BEST_FATTN_KERNEL_MKL) kname = "MKL"; + if (k == BEST_FATTN_KERNEL_VEC) kname = "VEC"; + int64_t delta = 0; + if (Dk == 256) { + delta = cur_nkv - last_nkv_d256; + last_nkv_d256 = cur_nkv; + } else if (Dk == 512) { + delta = cur_nkv - last_nkv_d512; + last_nkv_d512 = cur_nkv; + } + fprintf(stderr, "[FA-DISP] #%d %s D=%d n_kv=%lld delta=%lld " + "V_ne1=%lld\n", + fa_call_seq, kname, Dk, + (long long)cur_nkv, (long long)delta, + (long long)V_dbg->ne[1]); + fflush(stderr); + } + switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) { case BEST_FATTN_KERNEL_NONE: GGML_ABORT("Not support Flash-Attention"); @@ -234,6 +276,53 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst ggml_sycl_flash_attn_ext_mkl(ctx, dst); break; } + + // --- Output fingerprint (MKL_FA_DIAG=1) --- + // Copy first 64 float output values to host for fingerprinting. + // Compare MKL vs TILE (MKL_FA_DISABLE=1) to detect divergence. + // Only fingerprints the first 3 FA calls with n_kv >= 1024. + static int fa_diag = -1; + static int fa_diag_count = 0; + if (fa_diag < 0) { + const char * e = getenv("MKL_FA_DIAG"); + fa_diag = (e && e[0] == '1') ? 1 : 0; + } + if (fa_diag == 1 && fa_diag_count < 6) { + const ggml_tensor * K_diag = dst->src[1]; + const ggml_tensor * V_diag = dst->src[2]; + const ggml_tensor * Q_diag = dst->src[0]; + if (K_diag->ne[1] >= 1024) { + fa_diag_count++; + float diag_buf[64]; + dpct::queue_ptr q = ctx.stream(); + q->memcpy(diag_buf, dst->data, 64 * sizeof(float)); + q->wait(); + const char * kname = "???"; + best_fattn_kernel kb = ggml_sycl_get_best_fattn_kernel(ctx.device, dst); + if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL"; + if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE"; + if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC"; + fprintf(stderr, "[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld " + "n_qh=%lld n_kvh=%lld K=%s V=%s " + "nb1=%zu nb2=%zu first 64 floats:\n", + fa_diag_count, kname, + (int)K_diag->ne[0], (long long)K_diag->ne[1], + (long long)Q_diag->ne[1], + (long long)Q_diag->ne[2], (long long)K_diag->ne[2], + ggml_type_name(K_diag->type), + ggml_type_name(V_diag->type), + K_diag->nb[1], K_diag->nb[2]); + for (int i = 0; i < 64; i += 8) { + fprintf(stderr, " [%2d] %08x %08x %08x %08x %08x %08x %08x %08x\n", + i, + *(unsigned *)&diag_buf[i+0], *(unsigned *)&diag_buf[i+1], + *(unsigned *)&diag_buf[i+2], *(unsigned *)&diag_buf[i+3], + *(unsigned *)&diag_buf[i+4], *(unsigned *)&diag_buf[i+5], + *(unsigned *)&diag_buf[i+6], *(unsigned *)&diag_buf[i+7]); + } + fflush(stderr); + } + } } bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst) { From 5a81c11140f2502400cb150d5c40a4566dc36ac1 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Sun, 28 Jun 2026 22:08:49 -0500 Subject: [PATCH 03/11] Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro --- docs/backend/SYCL.md | 3 +++ ggml/src/ggml-sycl/fattn-mkl.cpp | 16 +++++++--------- ggml/src/ggml-sycl/fattn.cpp | 30 +++++++++++++++--------------- 3 files changed, 25 insertions(+), 24 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 8b0b9a18691a..2df3b764faa1 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -795,6 +795,9 @@ 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_DISABLE_DNN | 0 (default) or 1 | Disable 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_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Set to 0 to force the TILE kernel path for A/B testing. Activated at runtime when flash-attn is enabled (`-fa` or `--flash-attn on`), KV cache is quantized (`--cache-type-k/q *_0/*_1`), and KV length ≥ 1024. | +| GGML_SYCL_MKL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for MKL flash attention: GEMM/softmax timings, interleaved-head detection, and buffer memory usage. | +| GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. | | 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-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index 25cfab8cb1c8..42437c3121da 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -273,12 +273,12 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * const int chunk_size = std::min(MKL_FA_CHUNK_SIZE_KV, n_kv); const int64_t wg_size = 256; - // --- Debug output (gated by MKL_FA_DEBUG=1) --- + // --- Debug output (gated by GGML_SYCL_MKL_FA_DEBUG=1) --- static int mkl_call_count = 0; mkl_call_count++; static int mkl_debug = -1; if (mkl_debug < 0) { - const char * e = getenv("MKL_FA_DEBUG"); + const char * e = getenv("GGML_SYCL_MKL_FA_DEBUG"); mkl_debug = (e && e[0] == '1') ? 1 : 0; } const bool do_print = (mkl_debug == 1); @@ -299,7 +299,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * !V_is_K_view && ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]); if (do_print) { - fprintf(stderr, "[MKL-FA] #%d D=%d DV=%d n_q=%d n_kv=%d " + GGML_LOG_INFO("[MKL-FA] #%d D=%d DV=%d n_q=%d n_kv=%d " "n_qh=%d n_kvh=%d gqa=%d batch=%d K=%s V=%s " "chunk=%d buf=%.1fMB%s%s\n", mkl_call_count, DKQ, DV, n_queries, n_kv, @@ -310,14 +310,13 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * / (1024.0 * 1024.0), k_early_interleaved ? " K_ILV" : "", v_early_interleaved ? " V_ILV" : ""); - fprintf(stderr, "[MKL-FA] #%d Q-nb1=%lld Q-nb2=%lld " + GGML_LOG_INFO("[MKL-FA] #%d Q-nb1=%lld Q-nb2=%lld " "q_rs=%lld q_hs=%lld dst_rs=%lld dst_hs=%lld\n", mkl_call_count, (long long)Q->nb[1], (long long)Q->nb[2], (long long)q_row_stride, (long long)q_head_stride, (long long)(KQV->nb[1] / sizeof(float)), (long long)(KQV->nb[2] / sizeof(float))); - fflush(stderr); } // --- Stream and allocators --- @@ -569,7 +568,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * KQ_f32_ptr, this_chunk); stream->wait(); try { ev.wait_and_throw(); } catch (sycl::exception & e) { - fprintf(stderr, "[MKL-FA] GEMM KQ: %s\n", e.what()); + GGML_LOG_INFO("[MKL-FA] GEMM KQ: %s\n", e.what()); GGML_ABORT("MKL GEMM KQ failed"); } MKL_ACCUM(gemm_kq_time_us, t0); @@ -603,7 +602,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * VKQ_chunk_ptr, DV); stream->wait(); try { ev.wait_and_throw(); } catch (sycl::exception & e) { - fprintf(stderr, "[MKL-FA] GEMM VKQ: %s\n", e.what()); + GGML_LOG_INFO("[MKL-FA] GEMM VKQ: %s\n", e.what()); GGML_ABORT("MKL GEMM VKQ failed"); } MKL_ACCUM(gemm_vkq_time_us, t0); @@ -650,7 +649,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * + (int64_t)n_query_rows * sizeof(float) + (int64_t)n_query_rows * DKQ * sizeof(sycl::half) ) / (1024.0 * 1024.0); - fprintf(stderr, "[MKL-FA] #%d n_kv=%d n_q=%d time_us: " + GGML_LOG_INFO("[MKL-FA] #%d n_kv=%d n_q=%d time_us: " "dequant=%lld GEMM_KQ=%lld softmax=%lld GEMM_VKQ=%lld " "buf_mb=%.1f\n", mkl_call_count, n_kv, n_queries, @@ -659,6 +658,5 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * (long long)softmax_time_us, (long long)gemm_vkq_time_us, total_mb); - fflush(stderr); } } diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 164b4c2510bc..23bb0da42513 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -125,13 +125,15 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; // MKL path: XMX-accelerated GEMM for prompt processing with quantized KV cache. - // Activates at n_kv >= 1024 (matching the --batch-size) to cover the full prompt. - // FIXME: MKL GEMM calls are incompatible with SYCL graph capture replay. - // Set MKL_FA_DISABLE=1 to force TILE/VEC path for A/B testing. + // Activates automatically when flash-attn is enabled (--flash-attn on or -fa), + // KV cache is quantized (--cache-type-k/q *_0/*_1), batch size >= 1024, and + // n_kv >= 1024. Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B + // testing. Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... + // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. static int mkl_disable = -1; if (mkl_disable < 0) { - const char * e = getenv("MKL_FA_DISABLE"); - mkl_disable = (e && e[0] == '1') ? 1 : 0; + const char * e = getenv("GGML_SYCL_ENABLE_MKL_FA"); + mkl_disable = (e && e[0] == '0') ? 1 : 0; } if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024 && (ggml_is_quantized(K->type) || ggml_is_quantized(V->type))) { @@ -232,7 +234,7 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst // the same D — helps detect cache-truncation issues. static int nkv_debug = -1; if (nkv_debug < 0) { - const char * e = getenv("MKL_FA_DEBUG"); + const char * e = getenv("GGML_SYCL_MKL_FA_DEBUG"); nkv_debug = (e && e[0] == '1') ? 1 : 0; } if (nkv_debug == 1) { @@ -255,12 +257,11 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst delta = cur_nkv - last_nkv_d512; last_nkv_d512 = cur_nkv; } - fprintf(stderr, "[FA-DISP] #%d %s D=%d n_kv=%lld delta=%lld " + GGML_LOG_INFO("[FA-DISP] #%d %s D=%d n_kv=%lld delta=%lld " "V_ne1=%lld\n", fa_call_seq, kname, Dk, (long long)cur_nkv, (long long)delta, (long long)V_dbg->ne[1]); - fflush(stderr); } switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) { @@ -277,14 +278,14 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst break; } - // --- Output fingerprint (MKL_FA_DIAG=1) --- + // --- Output fingerprint (GGML_SYCL_MKL_FA_DIAG=1) --- // Copy first 64 float output values to host for fingerprinting. - // Compare MKL vs TILE (MKL_FA_DISABLE=1) to detect divergence. - // Only fingerprints the first 3 FA calls with n_kv >= 1024. + // Compare MKL vs TILE (GGML_SYCL_ENABLE_MKL_FA=0) to detect divergence. + // Only fingerprints the first 6 FA calls with n_kv >= 1024. static int fa_diag = -1; static int fa_diag_count = 0; if (fa_diag < 0) { - const char * e = getenv("MKL_FA_DIAG"); + const char * e = getenv("GGML_SYCL_MKL_FA_DIAG"); fa_diag = (e && e[0] == '1') ? 1 : 0; } if (fa_diag == 1 && fa_diag_count < 6) { @@ -302,7 +303,7 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL"; if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE"; if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC"; - fprintf(stderr, "[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld " + GGML_LOG_INFO("[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld " "n_qh=%lld n_kvh=%lld K=%s V=%s " "nb1=%zu nb2=%zu first 64 floats:\n", fa_diag_count, kname, @@ -313,14 +314,13 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst ggml_type_name(V_diag->type), K_diag->nb[1], K_diag->nb[2]); for (int i = 0; i < 64; i += 8) { - fprintf(stderr, " [%2d] %08x %08x %08x %08x %08x %08x %08x %08x\n", + GGML_LOG_INFO(" [%2d] %08x %08x %08x %08x %08x %08x %08x %08x\n", i, *(unsigned *)&diag_buf[i+0], *(unsigned *)&diag_buf[i+1], *(unsigned *)&diag_buf[i+2], *(unsigned *)&diag_buf[i+3], *(unsigned *)&diag_buf[i+4], *(unsigned *)&diag_buf[i+5], *(unsigned *)&diag_buf[i+6], *(unsigned *)&diag_buf[i+7]); } - fflush(stderr); } } } From bc757b125250ddb7bcd62123c98d20f6296dbef3 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Mon, 29 Jun 2026 10:12:45 -0500 Subject: [PATCH 04/11] Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro --- docs/backend/SYCL.md | 2 +- ggml/src/ggml-sycl/fattn-mkl.cpp | 19 +++++-------------- ggml/src/ggml-sycl/fattn.cpp | 9 +++------ 3 files changed, 9 insertions(+), 21 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 2df3b764faa1..baf4dafe9847 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_DISABLE_DNN | 0 (default) or 1 | Disable 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_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Set to 0 to force the TILE kernel path for A/B testing. Activated at runtime when flash-attn is enabled (`-fa` or `--flash-attn on`), KV cache is quantized (`--cache-type-k/q *_0/*_1`), and KV length ≥ 1024. | +| GGML_SYCL_ENABLE_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Automatically activates during prefill (prompt processing) when all conditions are met: (1) flash-attn enabled (`-fa` or `--flash-attn on`), (2) KV cache quantized (`--cache-type-k q8_0 --cache-type-v q8_0` or other `*_0/*_1` types), (3) batch size ≥ 1024 (`--batch-size 1024`), (4) prompt length ≥ 1024 tokens. Set to 0 to force the TILE kernel for A/B testing. Example minimum command: `llama-cli -m model.gguf -fa -ngl 99 --cache-type-k q8_0 --cache-type-v q8_0 --batch-size 1024 -p "your prompt"` | | GGML_SYCL_MKL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for MKL flash attention: GEMM/softmax timings, interleaved-head detection, and buffer memory usage. | | GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. | | 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/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index 42437c3121da..757f190900ba 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -1,10 +1,3 @@ -// -// MIT license -// Copyright (C) 2025 Intel Corporation -// SPDX-License-Identifier: MIT -// - -// // Flash attention via oneMKL GEMM (XMX-accelerated). // Uses column_major::gemm for Q*K^T and S*V matmuls // with an online softmax SYCL kernel. @@ -21,7 +14,6 @@ #include #include -#include #include #define MKL_FA_CHUNK_SIZE_KV 8192 @@ -278,8 +270,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * mkl_call_count++; static int mkl_debug = -1; if (mkl_debug < 0) { - const char * e = getenv("GGML_SYCL_MKL_FA_DEBUG"); - mkl_debug = (e && e[0] == '1') ? 1 : 0; + mkl_debug = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DEBUG", 0); } const bool do_print = (mkl_debug == 1); @@ -325,8 +316,10 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * ggml_sycl_fattn_kv_buffers & fbuf = ctx.fattn_buffers(); #define MKL_TAKE_TIME(t0) auto t0 = std::chrono::steady_clock::now() -#define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \ - (std::chrono::steady_clock::now() - (t0)).count() +#define MKL_ACCUM(acc, t0) do { if (do_print) { \ + acc += (int64_t)std::chrono::duration_cast \ + (std::chrono::steady_clock::now() - (t0)).count(); \ +} } while(0) int64_t gemm_kq_time_us = 0; int64_t gemm_vkq_time_us = 0; @@ -566,7 +559,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * Q_head_f16_ptr, DKQ, beta, KQ_f32_ptr, this_chunk); - stream->wait(); try { ev.wait_and_throw(); } catch (sycl::exception & e) { GGML_LOG_INFO("[MKL-FA] GEMM KQ: %s\n", e.what()); GGML_ABORT("MKL GEMM KQ failed"); @@ -600,7 +592,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * S_f16_ptr, this_chunk, beta, VKQ_chunk_ptr, DV); - stream->wait(); try { ev.wait_and_throw(); } catch (sycl::exception & e) { GGML_LOG_INFO("[MKL-FA] GEMM VKQ: %s\n", e.what()); GGML_ABORT("MKL GEMM VKQ failed"); diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 23bb0da42513..eb0138ae1b71 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -132,8 +132,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. static int mkl_disable = -1; if (mkl_disable < 0) { - const char * e = getenv("GGML_SYCL_ENABLE_MKL_FA"); - mkl_disable = (e && e[0] == '0') ? 1 : 0; + mkl_disable = !ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024 && (ggml_is_quantized(K->type) || ggml_is_quantized(V->type))) { @@ -234,8 +233,7 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst // the same D — helps detect cache-truncation issues. static int nkv_debug = -1; if (nkv_debug < 0) { - const char * e = getenv("GGML_SYCL_MKL_FA_DEBUG"); - nkv_debug = (e && e[0] == '1') ? 1 : 0; + nkv_debug = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DEBUG", 0); } if (nkv_debug == 1) { const ggml_tensor * K_dbg = dst->src[1]; @@ -285,8 +283,7 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst static int fa_diag = -1; static int fa_diag_count = 0; if (fa_diag < 0) { - const char * e = getenv("GGML_SYCL_MKL_FA_DIAG"); - fa_diag = (e && e[0] == '1') ? 1 : 0; + fa_diag = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DIAG", 0); } if (fa_diag == 1 && fa_diag_count < 6) { const ggml_tensor * K_diag = dst->src[1]; From b5b58ae5923a92ad3789d7d36c3804389bb18a10 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Mon, 29 Jun 2026 16:04:31 -0500 Subject: [PATCH 05/11] fattn-mkl: enable MKL FA for all KV cache types MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types. --- ggml/src/ggml-sycl/fattn.cpp | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index eb0138ae1b71..d1ca1bdc05ba 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -124,18 +124,19 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0; - // MKL path: XMX-accelerated GEMM for prompt processing with quantized KV cache. + // MKL path: XMX-accelerated GEMM for prompt processing (all KV cache types). + // The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before GEMM, + // so quantized, F16, BF16, and F32 caches all benefit from XMX acceleration. // Activates automatically when flash-attn is enabled (--flash-attn on or -fa), - // KV cache is quantized (--cache-type-k/q *_0/*_1), batch size >= 1024, and - // n_kv >= 1024. Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B - // testing. Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... + // batch size >= 1024, and n_kv >= 1024. + // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. + // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. static int mkl_disable = -1; if (mkl_disable < 0) { mkl_disable = !ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } - if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024 - && (ggml_is_quantized(K->type) || ggml_is_quantized(V->type))) { + if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024) { return BEST_FATTN_KERNEL_MKL; } @@ -186,6 +187,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: From e8b207e2fffde3b6a3c9244866ef7b7c9ce38faf Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Tue, 30 Jun 2026 08:14:42 -0500 Subject: [PATCH 06/11] fattn-mkl: rename mkl_disable -> mkl_enable for clarity --- ggml/src/ggml-sycl/fattn.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index d1ca1bdc05ba..72ae0b412ced 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -132,11 +132,11 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. - static int mkl_disable = -1; - if (mkl_disable < 0) { - mkl_disable = !ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); + static int mkl_enable = -1; + if (mkl_enable < 0) { + mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } - if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024) { + if (mkl_enable == 1 && Q->ne[1] >= 128 && K->ne[1] >= 1024) { return BEST_FATTN_KERNEL_MKL; } From 7083be7f9fb1d70f7cb655867ac6af7e80091a43 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Tue, 30 Jun 2026 14:34:45 -0500 Subject: [PATCH 07/11] fattn-mkl: refine MKL FA dispatch gates Three changes: 1. Remove quantized-only restriction - MKL FA activates for all KV cache types (F16 default, BF16, F32, quantized). The MKL kernel converts non-F16 K/V via to_fp16_sycl before GEMM. 2. Rename mkl_disable -> mkl_enable to match env var (GGML_SYCL_ENABLE_MKL_FA). 3. Replace batch-size threshold with Q->ne[1] >= 32 gate. Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where fused kernel beats MKL launch overhead. Routes all multi-token prefill through XMX-accelerated GEMM. Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse, 128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s. --- ggml/src/ggml-sycl/fattn.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 72ae0b412ced..f45d9edca6e0 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -127,8 +127,8 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // MKL path: XMX-accelerated GEMM for prompt processing (all KV cache types). // The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before GEMM, // so quantized, F16, BF16, and F32 caches all benefit from XMX acceleration. - // Activates automatically when flash-attn is enabled (--flash-attn on or -fa), - // batch size >= 1024, and n_kv >= 1024. + // Activates automatically when flash-attn is enabled (--flash-attn on or -fa) + // and n_kv >= 1024. // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. @@ -136,7 +136,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const if (mkl_enable < 0) { mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } - if (mkl_enable == 1 && Q->ne[1] >= 128 && K->ne[1] >= 1024) { + if (mkl_enable == 1 && Q->ne[1] >= 32 && K->ne[1] >= 1024) { return BEST_FATTN_KERNEL_MKL; } From 13f503255767e62ef39f6b8a0037fb28ddfa0337 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Thu, 2 Jul 2026 14:31:24 -0500 Subject: [PATCH 08/11] ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (#25222) which surfaced the same insight: always normalize inputs to contiguous F16 before GEMM. Co-Authored-By: Claude Code using DeepSeek-V4-Pro --- ggml/src/ggml-sycl/fattn-mkl.cpp | 102 ++++++++++++++++++++++++------- ggml/src/ggml-sycl/fattn.cpp | 25 +++++++- 2 files changed, 102 insertions(+), 25 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index 757f190900ba..6fc719fcddd3 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -357,14 +357,45 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * const bool v_interleaved = ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1; - if (K->type != GGML_TYPE_F16) { - const size_t bs = ggml_blck_size(K->type); - const size_t ts = ggml_type_size(K->type); - + { + // Normalize K to dense row-major F16 for GEMM. + // F16 is copied via memcpy (strides guaranteed standard by the gate). K_f16.alloc(ggml_nelements(K)); - if (ggml_is_contiguously_allocated(K) && !k_interleaved) { - // Dense source — flat to_fp16 preserves layout + if (K->type == GGML_TYPE_F16) { + // F16: always copy to dense for GEMM. + // Standard strides guaranteed by gate (nb1 == ne0 * sizeof(half)). + if (!k_interleaved) { + // Dense row-major F16 — flat memcpy + stream->memcpy(K_f16.ptr, K_data, ggml_nelements(K) * sizeof(sycl::half)); + } else { + // Interleaved F16 — strided copy to dense + const int64_t ne0 = K->ne[0]; + const int64_t ne1 = K->ne[1]; + const int64_t ne23 = (int64_t)K->ne[2] * K->ne[3]; + const int64_t src_nb1 = (int64_t)nb11; + const int64_t src_nb2 = (int64_t)nb12; + const sycl::half * src = (const sycl::half *)K_data; + sycl::half * dst = K_f16.ptr; + + stream->parallel_for( + sycl::range<3>((size_t)ne23, (size_t)ne1, (size_t)ne0), + [=](sycl::item<3> it) { + int64_t hb = it.get_id(0); + int64_t r = it.get_id(1); + int64_t c = it.get_id(2); + const sycl::half * src_row = (const sycl::half *)( + (const char *)src + hb * src_nb2 + r * src_nb1); + dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; + }); + + nb11 = K->ne[0] * sizeof(sycl::half); + nb12 = K->ne[1] * nb11; + nb13 = K->ne[2] * nb12; + } + } else if (ggml_is_contiguously_allocated(K) && !k_interleaved) { + const size_t bs = ggml_blck_size(K->type); + const size_t ts = ggml_type_size(K->type); to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst); to_fp16(K_data, K_f16.ptr, ggml_nelements(K), stream); @@ -372,30 +403,25 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * nb12 = nb12 * bs * sizeof(sycl::half) / ts; nb13 = nb13 * bs * sizeof(sycl::half) / ts; } else { - // Interleaved or non-contiguous source — use to_fp16_nc - // with PHYSICAL source strides so the output is dense. + 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_interleaved) { - // Interleaved: heads alternate at the row level. - // Row r of head h starts at block - // h * blk_per_row + r * n_kv_heads * blk_per_row. const int64_t blk_per_row = (int64_t)K->ne[0] / bs; - s01 = (int64_t)n_kv_heads * blk_per_row; // row stride - s02 = blk_per_row; // head stride - s03 = (int64_t)K->ne[1] * s01; // batch stride + s01 = (int64_t)n_kv_heads * blk_per_row; + s02 = blk_per_row; + s03 = (int64_t)K->ne[1] * s01; } else { s01 = nb11 / ts; s02 = nb12 / ts; s03 = nb13 / ts; } - to_fp16(K_data, K_f16.ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, stream); - // Output is dense row-major nb11 = K->ne[0] * sizeof(sycl::half); nb12 = K->ne[1] * nb11; nb13 = K->ne[2] * nb12; @@ -403,19 +429,50 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * K_data = (char *) K_f16.ptr; } - if (V->type != GGML_TYPE_F16) { + { + // Normalize V to dense row-major F16 for GEMM. if (V_is_K_view) { + // V shares K's buffer — reuse the already-normalized dense K data V_data = K_data; nb21 = nb11; nb22 = nb12; nb23 = nb13; } else { - const size_t bs = ggml_blck_size(V->type); - const size_t ts = ggml_type_size(V->type); - V_f16.alloc(ggml_nelements(V)); - if (ggml_is_contiguously_allocated(V) && !v_interleaved) { + if (V->type == GGML_TYPE_F16) { + if (!v_interleaved) { + // Dense row-major F16 — flat memcpy + stream->memcpy(V_f16.ptr, V_data, ggml_nelements(V) * sizeof(sycl::half)); + } else { + // Interleaved F16 — strided copy to dense + const int64_t ne0 = V->ne[0]; + const int64_t ne1 = V->ne[1]; + const int64_t ne23 = (int64_t)V->ne[2] * V->ne[3]; + const int64_t src_nb1 = (int64_t)nb21; + const int64_t src_nb2 = (int64_t)nb22; + const sycl::half * src = (const sycl::half *)V_data; + sycl::half * dst = V_f16.ptr; + + stream->parallel_for( + sycl::range<3>((size_t)ne23, (size_t)ne1, (size_t)ne0), + [=](sycl::item<3> it) { + int64_t hb = it.get_id(0); + int64_t r = it.get_id(1); + int64_t c = it.get_id(2); + const sycl::half * src_row = (const sycl::half *)( + (const char *)src + hb * src_nb2 + r * src_nb1); + dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; + }); + + nb21 = V->ne[0] * sizeof(sycl::half); + nb22 = V->ne[1] * nb21; + nb23 = V->ne[2] * nb22; + } + V_data = (char *) V_f16.ptr; + } else if (ggml_is_contiguously_allocated(V) && !v_interleaved) { + const size_t bs = ggml_blck_size(V->type); + const size_t ts = ggml_type_size(V->type); to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst); to_fp16(V_data, V_f16.ptr, ggml_nelements(V), stream); V_data = (char *) V_f16.ptr; @@ -424,6 +481,8 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * nb22 = nb22 * bs * sizeof(sycl::half) / ts; nb23 = nb23 * bs * sizeof(sycl::half) / ts; } 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; @@ -437,7 +496,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * s02 = nb22 / ts; s03 = nb23 / ts; } - to_fp16(V_data, V_f16.ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, stream); diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index f45d9edca6e0..33b583c4610a 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -122,13 +122,17 @@ 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; // MKL path: XMX-accelerated GEMM for prompt processing (all KV cache types). // The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before GEMM, // so quantized, F16, BF16, and F32 caches all benefit from XMX acceleration. // Activates automatically when flash-attn is enabled (--flash-attn on or -fa) - // and n_kv >= 1024. + // and n_kv >= 1024. Falls through to TILE/VEC for ALiBi, logit softcap, + // and mismatched batch dimensions (unsupported by the MKL kernel). // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. @@ -136,8 +140,23 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const if (mkl_enable < 0) { mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } - if (mkl_enable == 1 && Q->ne[1] >= 32 && K->ne[1] >= 1024) { - return BEST_FATTN_KERNEL_MKL; + if (mkl_enable == 1 && Q->ne[1] >= 32 && K->ne[1] >= 1024 && + max_bias == 0.0f && logit_softcap == 0.0f && + (Q->ne[3] == K->ne[3] || K->ne[3] == 1)) { + // F16 K/V strides must be a multiple of ne[0]*2 (the natural row size + // in bytes). This passes both dense (nb1 == ne0*2) and interleaved + // (nb1 == H * ne0*2). Only pathological test strides like nb1=32 or + // nb1=75 for ne0=40 fall through to TILE. + bool kv_strides_ok = true; + for (const ggml_tensor * t : {K, V}) { + if (t->type == GGML_TYPE_F16 && t->nb[1] % (t->ne[0] * 2) != 0) { + kv_strides_ok = false; + break; + } + } + if (kv_strides_ok) { + return BEST_FATTN_KERNEL_MKL; + } } for (const ggml_tensor * t : {Q, K, V, mask}) { From 6516c33e1586a2249dc6254e8d095f521e67668b Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Sun, 5 Jul 2026 11:38:46 -0500 Subject: [PATCH 09/11] SYCL: oneDNN SDPA Flash Attention for Quantized KV Caches Add four-tier XMX-accelerated flash attention dispatch: 1. ONEDNN: native F16 oneDNN SDPA (from PR #25222 by @hmscider) 2. HYBRID: dequant K/V to F16, then oneDNN SDPA (new) 3. MKL: oneMKL GEMM fallback for SDPA-incompatible shapes 4. TILE/VEC: existing scalar fallback for decode Fixes: permute formula (h/t swap), pool-alloc stream sync, MKL gate guards (mask, sinks, head_dim >= 64). Co-Authored-By: Claude Code on DeepSeek-v4-Pro --- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/fattn-hybrid.cpp | 339 ++++++++++++++++++++++++++++ ggml/src/ggml-sycl/fattn-hybrid.hpp | 26 +++ ggml/src/ggml-sycl/fattn-onednn.cpp | 233 +++++++++++++++++++ ggml/src/ggml-sycl/fattn-onednn.hpp | 50 ++++ ggml/src/ggml-sycl/fattn.cpp | 80 +++---- ggml/src/ggml-sycl/ggml-sycl.cpp | 3 + 7 files changed, 684 insertions(+), 48 deletions(-) create mode 100644 ggml/src/ggml-sycl/fattn-hybrid.cpp create mode 100644 ggml/src/ggml-sycl/fattn-hybrid.hpp 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 8534bd3581e5..57d1ec974e21 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -63,6 +63,7 @@ extern int g_ggml_sycl_disable_optimize; 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-hybrid.cpp b/ggml/src/ggml-sycl/fattn-hybrid.cpp new file mode 100644 index 000000000000..c670949d3a0e --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-hybrid.cpp @@ -0,0 +1,339 @@ +// +// MIT license +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: MIT +// + +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// + +// Hybrid flash attention: dequantize K/V to F16 → oneDNN SDPA (XMX fused kernel). +// Combines the dequant pipeline from PR #25025 with the hardware SDPA from PR #25222. +// Covers quantized, BF16, and F32 KV caches that oneDNN's native F16-only gate rejects. + +#include +#include +#include +#include + +#include "fattn-hybrid.hpp" +#include "fattn-onednn.hpp" +#include "fattn.hpp" // for fallback: ggml_sycl_flash_attn_ext_mkl +#include "fattn-tile.hpp" +#include "convert.hpp" +#include "fattn-buffers.hpp" + +#ifdef GGML_SYCL_DNNL +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" +#endif + +// --------------------------------------------------------------------------- +// Gate: eligible shapes for dequant + oneDNN SDPA +// --------------------------------------------------------------------------- +bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst) { +#if !GGML_SYCL_DNNL + GGML_UNUSED(dst); + return false; +#else + // Respect the global kill switch + if (!g_ggml_sycl_fa_onednn) { + 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]; + + // HYBRID handles non-F16 types (native F16 goes to ONEDNN). + if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) { + return false; + } + + // MKL prefill gate conditions + if (Q->ne[1] < 32 || K->ne[1] < 1024) { + 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; + } + if (!(Q->ne[3] == K->ne[3] || K->ne[3] == 1)) { + return false; + } + + // F16 stride alignment (from MKL gate) + for (const ggml_tensor * t : {K, V}) { + if (t->type == GGML_TYPE_F16 && t->nb[1] % (t->ne[0] * 2) != 0) { + return false; + } + } + + // SDPA conditions (same as fattn-onednn.cpp except F16-only KV) + if (!mask || mask->type != GGML_TYPE_F16 || mask->ne[2] != 1 || mask->ne[3] != 1) { + return false; + } + if (sinks) { + return false; + } + const int64_t d = K->ne[0]; + if (V->ne[0] != d || Q->ne[3] != 1) { + return false; + } + if (K->ne[2] == 0 || Q->ne[2] % K->ne[2] != 0) { + return false; + } + + return true; +#endif +} + +// --------------------------------------------------------------------------- +// Implementation: dequant K/V → oneDNN SDPA +// --------------------------------------------------------------------------- +#if GGML_SYCL_DNNL + +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" + +using namespace dnnl; +using namespace dnnl::graph; + +void ggml_sycl_flash_attn_ext_hybrid(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]; + ggml_tensor * KQV = dst; + + const int64_t d = K->ne[0]; + const int64_t seq = K->ne[1]; + const int64_t Hkv = K->ne[2]; + const int64_t H = Q->ne[2]; + const int64_t q = Q->ne[1]; + const int64_t mb = Q->ne[3]; + + float kq_scale = 1.0f; + memcpy(&kq_scale, (const float *) KQV->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); + + // --- Step 1: Dequant K to dense F16 --- + ggml_sycl_fattn_alloc K_f16(ctx.fattn_buffers().K); + { + const char * K_data = (const char *)K->data; + const bool k_interleaved = + ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]) && K->ne[2] > 1; + + K_f16.alloc(ggml_nelements(K)); + + if (K->type == GGML_TYPE_F16) { + if (!k_interleaved) { + stream->memcpy(K_f16.ptr, K_data, ggml_nelements(K) * sizeof(sycl::half)); + } else { + const int64_t ne0 = K->ne[0], ne1 = K->ne[1]; + const int64_t ne23 = (int64_t)K->ne[2] * K->ne[3]; + const int64_t src_nb1 = (int64_t)K->nb[1]; + const int64_t src_nb2 = (int64_t)K->nb[2]; + const sycl::half * src = (const sycl::half *)K_data; + sycl::half * dst = K_f16.ptr; + stream->parallel_for( + sycl::range<3>((size_t)ne23, (size_t)ne1, (size_t)ne0), + [=](sycl::item<3> it) { + int64_t hb = it.get_id(0), r = it.get_id(1), c = it.get_id(2); + const sycl::half * src_row = (const sycl::half *)( + (const char *)src + hb * src_nb2 + r * src_nb1); + dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; + }); + } + } else if (ggml_is_contiguously_allocated(K) && !k_interleaved) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst); + to_fp16(K_data, K_f16.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_interleaved) { + 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_f16.ptr, + K->ne[0], K->ne[1], K->ne[2], K->ne[3], + s01, s02, s03, stream); + } + } + + // --- Step 2: Dequant V to dense F16 --- + // K and V share buffers in quantized KV cache (V is a view of K). + // Detect it: same data pointer means reuse K_f16. + ggml_sycl_fattn_alloc V_f16(ctx.fattn_buffers().V); + bool V_is_K_view = (K->type != GGML_TYPE_F32 && V->type != GGML_TYPE_F32 && + K->data == V->data); + if (V_is_K_view) { + // V shares K's buffer — reuse the already-normalized dense K data + V_f16.ptr = K_f16.ptr; // don't allocate, just alias + } else { + const char * V_data = (const char *)V->data; + const bool v_interleaved = + ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1; + + V_f16.alloc(ggml_nelements(V)); + + if (V->type == GGML_TYPE_F16) { + if (!v_interleaved) { + stream->memcpy(V_f16.ptr, V_data, ggml_nelements(V) * sizeof(sycl::half)); + } else { + const int64_t ne0 = V->ne[0], ne1 = V->ne[1]; + const int64_t ne23 = (int64_t)V->ne[2] * V->ne[3]; + const int64_t src_nb1 = (int64_t)V->nb[1]; + const int64_t src_nb2 = (int64_t)V->nb[2]; + const sycl::half * src = (const sycl::half *)V_data; + sycl::half * dst = V_f16.ptr; + stream->parallel_for( + sycl::range<3>((size_t)ne23, (size_t)ne1, (size_t)ne0), + [=](sycl::item<3> it) { + int64_t hb = it.get_id(0), r = it.get_id(1), c = it.get_id(2); + const sycl::half * src_row = (const sycl::half *)( + (const char *)src + hb * src_nb2 + r * src_nb1); + dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; + }); + } + } else if (ggml_is_contiguously_allocated(V) && !v_interleaved) { + to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst); + to_fp16(V_data, V_f16.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_interleaved) { + 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_f16.ptr, + V->ne[0], V->ne[1], V->ne[2], V->ne[3], + s01, s02, s03, stream); + } + } + + // --- Step 3: Copy Q to dense F16 --- + ggml_sycl_pool_alloc Qf(ctx.pool(), (size_t) H * q * d); + { + const char * Q_data = (const char *)Q->data; + const int64_t n = H * q * d; + sycl::half * Qf_ptr = Qf.get(); + size_t Q_nb1 = Q->nb[1]; + size_t Q_nb2 = Q->nb[2]; + size_t Q_nb3 = Q->nb[3]; + 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 % d; i /= d; + const int64_t i1 = i % q; i /= q; + const int64_t i2 = i % H; + const int64_t i3 = i / H; + const float * p = (const float *) (Q_data + i1 * Q_nb1 + i2 * Q_nb2 + i3 * Q_nb3) + i0; + Qf_ptr[gid] = (sycl::half) (*p); + }); + } + + // --- Step 4: Build/run oneDNN SDPA --- + 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); + + // Compile once per (device, shape), reuse across layers. + 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; + if (!E.ok) { + // Partition failed — fall back to MKL + ggml_sycl_flash_attn_ext_mkl(ctx, dst); + return; + } + + auto id2ptr = [&](size_t r) -> void * { + if (r == E.id_q) return Qf.get(); + if (r == E.id_k) return K_f16.ptr; + if (r == E.id_v) return V_f16.ptr; + 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}); + + // --- Step 5: Permute SDPA output → ggml dst --- + const int64_t n = mb * H * q * d; + float * dst_data = (float *) dst->data; + const sycl::half * out_data = outf.get(); + const size_t dst_nb0 = sizeof(float); + const size_t dst_nb1 = dst->nb[1]; + const size_t dst_nb2 = dst->nb[2]; + const size_t dst_nb3 = dst->nb[3]; + 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; + int64_t off = (e * dst_nb0 + h * dst_nb1 + t * dst_nb2 + b * dst_nb3) / sizeof(float); + dst_data[off] = (float) out_data[gid]; + }); + + // Wait for all GPU work to finish before pool allocs go out of scope. + // Without this, the pool frees and reuses memory while the GPU is still + // reading it, corrupting the output. + stream->wait(); +} +catch (const std::exception & e) { + GGML_LOG_WARN("%s: hybrid SDPA failed (%s); falling back to MKL\n", __func__, e.what()); + ggml_sycl_flash_attn_ext_mkl(ctx, dst); +} + +#else // !GGML_SYCL_DNNL + +void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + // oneDNN not compiled — fall back to MKL + ggml_sycl_flash_attn_ext_mkl(ctx, dst); +} + +#endif // GGML_SYCL_DNNL diff --git a/ggml/src/ggml-sycl/fattn-hybrid.hpp b/ggml/src/ggml-sycl/fattn-hybrid.hpp new file mode 100644 index 000000000000..ebc437937c5b --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-hybrid.hpp @@ -0,0 +1,26 @@ +// +// MIT license +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: MIT +// + +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// + +#ifndef GGML_SYCL_FATTN_HYBRID_HPP +#define GGML_SYCL_FATTN_HYBRID_HPP + +#include "common.hpp" + +// Check whether this FA call qualifies for the hybrid path: +// dequantize K/V to f16, then route through oneDNN's fused SDPA kernel. +bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst); + +// Run FA: dequant K/V → oneDNN SDPA → permute to ggml dst. +// Falls back to MKL GEMM on failure. +void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + +#endif // GGML_SYCL_FATTN_HYBRID_HPP diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp new file mode 100644 index 000000000000..c7055e5e149a --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -0,0 +1,233 @@ +// +// MIT license +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: MIT +// + +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// + +#include +#include +#include +#include +#include +#include + +#include "fattn-onednn.hpp" +#include "fattn-tile.hpp" + +// Minimum query length to treat as prefill. +#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; } + 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]; + + if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) { return false; } + 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; } + const int64_t d = K->ne[0]; + if (V->ne[0] != d || Q->ne[3] != 1) { return false; } + if (K->ne[2] == 0 || Q->ne[2] % K->ne[2] != 0) { return false; } + if (Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) { return false; } + if (Q->ne[0] < 64) { return false; } // precision loss on small head dims + return true; +#endif +} + +#if GGML_SYCL_DNNL + +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" + +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_fp16_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]). +// Uses actual dst strides (nb1, nb2, nb3) to handle both dense and interleaved layouts. +static void permute_sdpa_out_sycl(const sycl::half * out, float * dst, + int64_t mb, int64_t H, int64_t q, int64_t d, + size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream) { + const int64_t n = mb * H * q * d; + const size_t nb0 = sizeof(float); + 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; + int64_t off = (e * nb0 + h * nb1 + t * nb2 + b * nb3) / sizeof(float); + dst[off] = (float) out[gid]; + }); +} + +// Build + compile the contiguous-input GQA SDPA graph. +// MatMul → Divide → Add(mask) → SoftMax → MatMul, f16 out. +// Partitions=1, sdp_primitive_kernel_t (the systolic XMX kernel). +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); + 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"); + 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 required to hit sdp_primitive_kernel_t. + auto output = logical_tensor(id++, t, o_sz, ltype::strided); + 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; + } + 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]; + const int64_t seq = K->ne[1]; + const int64_t Hkv = K->ne[2]; + const int64_t H = Q->ne[2]; + const int64_t q = Q->ne[1]; + const int64_t mb = Q->ne[3]; + + 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); + + // Copy inputs to contiguous f16 (head-major layout for the SDPA graph). + 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_fp16_sycl ((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream); + cont_to_fp16_sycl((const char *) K->data, Kf.get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream); + cont_to_fp16_sycl((const char *) V->data, Vf.get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream); + + // Scale as reciprocal: score / (1/kq_scale) == score * kq_scale. + 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); + + // 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; + 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, + dst->nb[1], dst->nb[2], dst->nb[3], stream); + + // Wait for all GPU work to finish before pool allocs go out of scope. + stream->wait(); +} +catch (const std::exception & e) { + 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 000000000000..4b0e4eebf968 --- /dev/null +++ b/ggml/src/ggml-sycl/fattn-onednn.hpp @@ -0,0 +1,50 @@ +// +// MIT license +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: MIT +// + +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// + +#ifndef GGML_SYCL_FATTN_ONEDNN_HPP +#define GGML_SYCL_FATTN_ONEDNN_HPP + +#include "common.hpp" + +// Static-only check: oneDNN Graph SDPA path for native f16 KV FA. +// (f16 KV, mask required, no softcap/ALiBi, single stream, prefill-sized q.) +bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst); + +// Run flash attention through oneDNN's fused XMX SDPA kernel. +// Falls back to the TILE kernel on any failure. +void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + +#if GGML_SYCL_DNNL +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" + +// Shared SDPA partition type and builder — used by both the native-F16 +// onednn path and the hybrid dequant-then-SDPA path. +struct sdpa_partition { + dnnl::graph::compiled_partition cp; + std::vector ins; + dnnl::graph::logical_tensor out; + size_t id_q = 0, id_k = 0, id_v = 0, id_scale = 0, id_mask = 0; + bool ok = false; +}; + +sdpa_partition build_sdpa(const dnnl::engine & eng, int H, int Hkv, int q, int seq, int d); + +// Permute SDPA output (f16 [mb,H,q,d]) → ggml dst (f32 [d,H,q,mb]). +// Uses actual dst strides for interleaved head layout support. +void permute_sdpa_out_sycl(const sycl::half * out, float * dst, + int64_t mb, int64_t H, int64_t q, int64_t d, + size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream); +#endif + +#endif // GGML_SYCL_FATTN_ONEDNN_HPP diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 33b583c4610a..900fee9e5f68 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -18,6 +18,8 @@ #include "fattn-tile.hpp" #include "fattn-vec.hpp" #include "fattn.hpp" +#include "fattn-onednn.hpp" +#include "fattn-hybrid.hpp" #define FATTN_VEC_CASE(D, type_K, type_V) \ @@ -96,7 +98,9 @@ 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, // oneDNN SDPA: native F16 (PR #25222) BEST_FATTN_KERNEL_TILE = 200, + BEST_FATTN_KERNEL_HYBRID = 250, // dequant + oneDNN SDPA (frankenmerge) BEST_FATTN_KERNEL_MKL = 300, }; @@ -115,6 +119,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); @@ -136,11 +141,27 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. + + // XMX-accelerated paths (prefill only, Q >= 32). + // Dispatch order: native-F16 oneDNN SDPA → hybrid dequant+SDPA → MKL GEMM. + + // Path 1: oneDNN SDPA for native F16 KV (from PR #25222). + if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) { + return BEST_FATTN_KERNEL_ONEDNN; + } + + // Path 2: Hybrid — dequantize K/V to F16, then oneDNN SDPA. + if (ggml_sycl_flash_attn_ext_hybrid_supported(dst)) { + return BEST_FATTN_KERNEL_HYBRID; + } + + // Path 3: MKL GEMM — handles SDPA-incompatible shapes (no mask, ALiBi, etc.). static int mkl_enable = -1; if (mkl_enable < 0) { mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); } - if (mkl_enable == 1 && Q->ne[1] >= 32 && K->ne[1] >= 1024 && + if (mkl_enable == 1 && Q->ne[0] >= 64 && mask && !sinks && + Q->ne[1] >= 32 && K->ne[1] >= 1024 && max_bias == 0.0f && logit_softcap == 0.0f && (Q->ne[3] == K->ne[3] || K->ne[3] == 1)) { // F16 K/V strides must be a multiple of ne[0]*2 (the natural row size @@ -228,8 +249,6 @@ 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: - // If there are no tensor cores available, use the generic tile kernel: if (can_use_vector_kernel) { if (!ggml_is_quantized(K->type) && !ggml_is_quantized(V->type)) { @@ -283,7 +302,8 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst (long long)V_dbg->ne[1]); } - 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_TILE: @@ -292,55 +312,19 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst case BEST_FATTN_KERNEL_VEC: ggml_sycl_flash_attn_ext_vec(ctx, dst); break; +#if GGML_SYCL_DNNL + case BEST_FATTN_KERNEL_ONEDNN: + ggml_sycl_flash_attn_ext_onednn(ctx, dst); + break; + case BEST_FATTN_KERNEL_HYBRID: + ggml_sycl_flash_attn_ext_hybrid(ctx, dst); + break; +#endif case BEST_FATTN_KERNEL_MKL: ggml_sycl_flash_attn_ext_mkl(ctx, dst); break; } - // --- Output fingerprint (GGML_SYCL_MKL_FA_DIAG=1) --- - // Copy first 64 float output values to host for fingerprinting. - // Compare MKL vs TILE (GGML_SYCL_ENABLE_MKL_FA=0) to detect divergence. - // Only fingerprints the first 6 FA calls with n_kv >= 1024. - static int fa_diag = -1; - static int fa_diag_count = 0; - if (fa_diag < 0) { - fa_diag = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DIAG", 0); - } - if (fa_diag == 1 && fa_diag_count < 6) { - const ggml_tensor * K_diag = dst->src[1]; - const ggml_tensor * V_diag = dst->src[2]; - const ggml_tensor * Q_diag = dst->src[0]; - if (K_diag->ne[1] >= 1024) { - fa_diag_count++; - float diag_buf[64]; - dpct::queue_ptr q = ctx.stream(); - q->memcpy(diag_buf, dst->data, 64 * sizeof(float)); - q->wait(); - const char * kname = "???"; - best_fattn_kernel kb = ggml_sycl_get_best_fattn_kernel(ctx.device, dst); - if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL"; - if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE"; - if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC"; - GGML_LOG_INFO("[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld " - "n_qh=%lld n_kvh=%lld K=%s V=%s " - "nb1=%zu nb2=%zu first 64 floats:\n", - fa_diag_count, kname, - (int)K_diag->ne[0], (long long)K_diag->ne[1], - (long long)Q_diag->ne[1], - (long long)Q_diag->ne[2], (long long)K_diag->ne[2], - ggml_type_name(K_diag->type), - ggml_type_name(V_diag->type), - K_diag->nb[1], K_diag->nb[2]); - for (int i = 0; i < 64; i += 8) { - GGML_LOG_INFO(" [%2d] %08x %08x %08x %08x %08x %08x %08x %08x\n", - i, - *(unsigned *)&diag_buf[i+0], *(unsigned *)&diag_buf[i+1], - *(unsigned *)&diag_buf[i+2], *(unsigned *)&diag_buf[i+3], - *(unsigned *)&diag_buf[i+4], *(unsigned *)&diag_buf[i+5], - *(unsigned *)&diag_buf[i+6], *(unsigned *)&diag_buf[i+7]); - } - } - } } bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst) { diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 41449db665ec..e4070149a381 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -90,6 +90,7 @@ int g_ggml_sycl_use_async_mem_op_requested = 1; int g_ggml_sycl_use_level_zero_api = 0; int g_ggml_sycl_enable_flash_attention = 1; int g_ggml_sycl_dev2dev_memcpy = DEV2DEV_MEMCPY_SYCL; +int g_ggml_sycl_fa_onednn = 1; int g_ggml_sycl_usm_system = 0; static ggml_sycl_device_info ggml_sycl_init() { @@ -273,6 +274,7 @@ static void ggml_check_sycl() try { g_ggml_sycl_disable_optimize = ggml_sycl_get_env("GGML_SYCL_DISABLE_OPT", 0); g_ggml_sycl_disable_graph = ggml_sycl_get_env("GGML_SYCL_DISABLE_GRAPH", 1); g_ggml_sycl_disable_dnn = ggml_sycl_get_env("GGML_SYCL_DISABLE_DNN", 0); + 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_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0); @@ -341,6 +343,7 @@ static void ggml_check_sycl() try { #endif #if GGML_SYCL_DNNL GGML_LOG_INFO(" GGML_SYCL_DISABLE_DNN: %d\n", g_ggml_sycl_disable_dnn); + GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN: %d\n", g_ggml_sycl_fa_onednn); #else GGML_LOG_INFO(" GGML_SYCL_DISABLE_DNN: DNN disabled by compile flag\n"); #endif From 9e9524e04970824b3567fb6cb50a90af30be5794 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Sun, 5 Jul 2026 20:18:31 -0500 Subject: [PATCH 10/11] Refactor to HYBRID-only: remove ONEDNN files, move build_sdpa() into hybrid MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Remove fattn-onednn.cpp/hpp (reserved for PR #25222) - Move sdpa_partition + build_sdpa() into fattn-hybrid.cpp/hpp - Strip FA-DISP diagnostics from fattn.cpp - Comment out ONEDNN enum/dispatch (uncomment when #25222 merges) - Update gate comments for HYBRID-only dispatch order - HYBRID now ships build_sdpa() independently — no dependency on ONEDNN files Build: cmake --preset x64-windows-sycl-release -DGGML_SYCL_DNN=ON -DGGML_SYCL_F16=ON Co-Authored-By: Claude Opus 4.8 --- ggml/src/ggml-sycl/fattn-hybrid.cpp | 69 +++++++- ggml/src/ggml-sycl/fattn-hybrid.hpp | 17 ++ ggml/src/ggml-sycl/fattn-onednn.cpp | 233 ---------------------------- ggml/src/ggml-sycl/fattn-onednn.hpp | 50 ------ ggml/src/ggml-sycl/fattn.cpp | 59 ++----- 5 files changed, 94 insertions(+), 334 deletions(-) delete mode 100644 ggml/src/ggml-sycl/fattn-onednn.cpp delete mode 100644 ggml/src/ggml-sycl/fattn-onednn.hpp diff --git a/ggml/src/ggml-sycl/fattn-hybrid.cpp b/ggml/src/ggml-sycl/fattn-hybrid.cpp index c670949d3a0e..9c32e1d28bf9 100644 --- a/ggml/src/ggml-sycl/fattn-hybrid.cpp +++ b/ggml/src/ggml-sycl/fattn-hybrid.cpp @@ -20,7 +20,6 @@ #include #include "fattn-hybrid.hpp" -#include "fattn-onednn.hpp" #include "fattn.hpp" // for fallback: ggml_sycl_flash_attn_ext_mkl #include "fattn-tile.hpp" #include "convert.hpp" @@ -51,7 +50,7 @@ bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst) { const ggml_tensor * mask = dst->src[3]; const ggml_tensor * sinks = dst->src[4]; - // HYBRID handles non-F16 types (native F16 goes to ONEDNN). + // HYBRID handles non-F16 types (native F16 goes to ONEDNN when merged). if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) { return false; } @@ -77,7 +76,7 @@ bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst) { } } - // SDPA conditions (same as fattn-onednn.cpp except F16-only KV) + // SDPA conditions if (!mask || mask->type != GGML_TYPE_F16 || mask->ne[2] != 1 || mask->ne[3] != 1) { return false; } @@ -97,7 +96,7 @@ bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst) { } // --------------------------------------------------------------------------- -// Implementation: dequant K/V → oneDNN SDPA +// Implementation // --------------------------------------------------------------------------- #if GGML_SYCL_DNNL @@ -108,6 +107,68 @@ bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst) { using namespace dnnl; using namespace dnnl::graph; +// Build + compile the contiguous-input GQA SDPA graph. +// MatMul → Divide → Add(mask) → SoftMax → MatMul, f16 out. +// Partitions=1, sdp_primitive_kernel_t (the systolic XMX kernel). +// Originally from fattn-onednn.cpp (PR #25222 by @hmscider). +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); + 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"); + 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 required to hit sdp_primitive_kernel_t. + auto output = logical_tensor(id++, t, o_sz, ltype::strided); + 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; + } + 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_hybrid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) try { const ggml_tensor * Q = dst->src[0]; diff --git a/ggml/src/ggml-sycl/fattn-hybrid.hpp b/ggml/src/ggml-sycl/fattn-hybrid.hpp index ebc437937c5b..990cd6cc51b4 100644 --- a/ggml/src/ggml-sycl/fattn-hybrid.hpp +++ b/ggml/src/ggml-sycl/fattn-hybrid.hpp @@ -23,4 +23,21 @@ bool ggml_sycl_flash_attn_ext_hybrid_supported(const ggml_tensor * dst); // Falls back to MKL GEMM on failure. void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tensor * dst); +#if GGML_SYCL_DNNL +#include "dnnl.hpp" +#include "dnnl_sycl.hpp" +#include "oneapi/dnnl/dnnl_graph.hpp" + +// Compiled SDPA partition — caches the fused MatMul+SoftMax kernel. +struct sdpa_partition { + dnnl::graph::compiled_partition cp; + std::vector ins; + dnnl::graph::logical_tensor out; + size_t id_q = 0, id_k = 0, id_v = 0, id_scale = 0, id_mask = 0; + bool ok = false; +}; + +sdpa_partition build_sdpa(const dnnl::engine & eng, int H, int Hkv, int q, int seq, int d); +#endif + #endif // GGML_SYCL_FATTN_HYBRID_HPP diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp deleted file mode 100644 index c7055e5e149a..000000000000 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ /dev/null @@ -1,233 +0,0 @@ -// -// MIT license -// Copyright (C) 2025 Intel Corporation -// SPDX-License-Identifier: MIT -// - -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// - -#include -#include -#include -#include -#include -#include - -#include "fattn-onednn.hpp" -#include "fattn-tile.hpp" - -// Minimum query length to treat as prefill. -#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; } - 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]; - - if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) { return false; } - 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; } - const int64_t d = K->ne[0]; - if (V->ne[0] != d || Q->ne[3] != 1) { return false; } - if (K->ne[2] == 0 || Q->ne[2] % K->ne[2] != 0) { return false; } - if (Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) { return false; } - if (Q->ne[0] < 64) { return false; } // precision loss on small head dims - return true; -#endif -} - -#if GGML_SYCL_DNNL - -#include "dnnl.hpp" -#include "dnnl_sycl.hpp" -#include "oneapi/dnnl/dnnl_graph.hpp" - -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_fp16_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]). -// Uses actual dst strides (nb1, nb2, nb3) to handle both dense and interleaved layouts. -static void permute_sdpa_out_sycl(const sycl::half * out, float * dst, - int64_t mb, int64_t H, int64_t q, int64_t d, - size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream) { - const int64_t n = mb * H * q * d; - const size_t nb0 = sizeof(float); - 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; - int64_t off = (e * nb0 + h * nb1 + t * nb2 + b * nb3) / sizeof(float); - dst[off] = (float) out[gid]; - }); -} - -// Build + compile the contiguous-input GQA SDPA graph. -// MatMul → Divide → Add(mask) → SoftMax → MatMul, f16 out. -// Partitions=1, sdp_primitive_kernel_t (the systolic XMX kernel). -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); - 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"); - 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 required to hit sdp_primitive_kernel_t. - auto output = logical_tensor(id++, t, o_sz, ltype::strided); - 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; - } - 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]; - const int64_t seq = K->ne[1]; - const int64_t Hkv = K->ne[2]; - const int64_t H = Q->ne[2]; - const int64_t q = Q->ne[1]; - const int64_t mb = Q->ne[3]; - - 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); - - // Copy inputs to contiguous f16 (head-major layout for the SDPA graph). - 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_fp16_sycl ((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream); - cont_to_fp16_sycl((const char *) K->data, Kf.get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream); - cont_to_fp16_sycl((const char *) V->data, Vf.get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream); - - // Scale as reciprocal: score / (1/kq_scale) == score * kq_scale. - 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); - - // 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; - 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, - dst->nb[1], dst->nb[2], dst->nb[3], stream); - - // Wait for all GPU work to finish before pool allocs go out of scope. - stream->wait(); -} -catch (const std::exception & e) { - 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 deleted file mode 100644 index 4b0e4eebf968..000000000000 --- a/ggml/src/ggml-sycl/fattn-onednn.hpp +++ /dev/null @@ -1,50 +0,0 @@ -// -// MIT license -// Copyright (C) 2025 Intel Corporation -// SPDX-License-Identifier: MIT -// - -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// - -#ifndef GGML_SYCL_FATTN_ONEDNN_HPP -#define GGML_SYCL_FATTN_ONEDNN_HPP - -#include "common.hpp" - -// Static-only check: oneDNN Graph SDPA path for native f16 KV FA. -// (f16 KV, mask required, no softcap/ALiBi, single stream, prefill-sized q.) -bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst); - -// Run flash attention through oneDNN's fused XMX SDPA kernel. -// Falls back to the TILE kernel on any failure. -void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst); - -#if GGML_SYCL_DNNL -#include "dnnl.hpp" -#include "dnnl_sycl.hpp" -#include "oneapi/dnnl/dnnl_graph.hpp" - -// Shared SDPA partition type and builder — used by both the native-F16 -// onednn path and the hybrid dequant-then-SDPA path. -struct sdpa_partition { - dnnl::graph::compiled_partition cp; - std::vector ins; - dnnl::graph::logical_tensor out; - size_t id_q = 0, id_k = 0, id_v = 0, id_scale = 0, id_mask = 0; - bool ok = false; -}; - -sdpa_partition build_sdpa(const dnnl::engine & eng, int H, int Hkv, int q, int seq, int d); - -// Permute SDPA output (f16 [mb,H,q,d]) → ggml dst (f32 [d,H,q,mb]). -// Uses actual dst strides for interleaved head layout support. -void permute_sdpa_out_sycl(const sycl::half * out, float * dst, - int64_t mb, int64_t H, int64_t q, int64_t d, - size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream); -#endif - -#endif // GGML_SYCL_FATTN_ONEDNN_HPP diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 900fee9e5f68..d04e3578fbe8 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -18,7 +18,7 @@ #include "fattn-tile.hpp" #include "fattn-vec.hpp" #include "fattn.hpp" -#include "fattn-onednn.hpp" +// #include "fattn-onednn.hpp" // from PR #25222 — add when merged #include "fattn-hybrid.hpp" @@ -98,9 +98,9 @@ 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, // oneDNN SDPA: native F16 (PR #25222) + // BEST_FATTN_KERNEL_ONEDNN = 150, // native F16 SDPA — reserved for PR #25222 BEST_FATTN_KERNEL_TILE = 200, - BEST_FATTN_KERNEL_HYBRID = 250, // dequant + oneDNN SDPA (frankenmerge) + BEST_FATTN_KERNEL_HYBRID = 250, // dequant K/V → oneDNN SDPA BEST_FATTN_KERNEL_MKL = 300, }; @@ -143,19 +143,16 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. // XMX-accelerated paths (prefill only, Q >= 32). - // Dispatch order: native-F16 oneDNN SDPA → hybrid dequant+SDPA → MKL GEMM. + // Dispatch order: HYBRID (dequant→SDPA) → MKL GEMM. + // When PR #25222 merges, add ONEDNN path above HYBRID: + // if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) return BEST_FATTN_KERNEL_ONEDNN; - // Path 1: oneDNN SDPA for native F16 KV (from PR #25222). - if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) { - return BEST_FATTN_KERNEL_ONEDNN; - } - - // Path 2: Hybrid — dequantize K/V to F16, then oneDNN SDPA. + // Path 1: Hybrid — dequantize K/V to F16, then oneDNN SDPA. if (ggml_sycl_flash_attn_ext_hybrid_supported(dst)) { return BEST_FATTN_KERNEL_HYBRID; } - // Path 3: MKL GEMM — handles SDPA-incompatible shapes (no mask, ALiBi, etc.). + // Path 2: MKL GEMM — handles SDPA-incompatible shapes (no mask, ALiBi, etc.). static int mkl_enable = -1; if (mkl_enable < 0) { mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); @@ -269,39 +266,6 @@ 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); - // n_kv watchdog: log when n_kv differs from the last FA call with - // the same D — helps detect cache-truncation issues. - static int nkv_debug = -1; - if (nkv_debug < 0) { - nkv_debug = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DEBUG", 0); - } - if (nkv_debug == 1) { - const ggml_tensor * K_dbg = dst->src[1]; - const ggml_tensor * V_dbg = dst->src[2]; - static int64_t last_nkv_d256 = 0, last_nkv_d512 = 0; - static int fa_call_seq = 0; - fa_call_seq++; - int64_t cur_nkv = K_dbg->ne[1]; - int Dk = (int)K_dbg->ne[0]; - const char * kname = "TILE"; - best_fattn_kernel k = ggml_sycl_get_best_fattn_kernel(ctx.device, dst); - if (k == BEST_FATTN_KERNEL_MKL) kname = "MKL"; - if (k == BEST_FATTN_KERNEL_VEC) kname = "VEC"; - int64_t delta = 0; - if (Dk == 256) { - delta = cur_nkv - last_nkv_d256; - last_nkv_d256 = cur_nkv; - } else if (Dk == 512) { - delta = cur_nkv - last_nkv_d512; - last_nkv_d512 = cur_nkv; - } - GGML_LOG_INFO("[FA-DISP] #%d %s D=%d n_kv=%lld delta=%lld " - "V_ne1=%lld\n", - fa_call_seq, kname, Dk, - (long long)cur_nkv, (long long)delta, - (long long)V_dbg->ne[1]); - } - const best_fattn_kernel fk = ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst); switch (fk) { case BEST_FATTN_KERNEL_NONE: @@ -313,9 +277,10 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst ggml_sycl_flash_attn_ext_vec(ctx, dst); break; #if GGML_SYCL_DNNL - case BEST_FATTN_KERNEL_ONEDNN: - ggml_sycl_flash_attn_ext_onednn(ctx, dst); - break; + // ONEDNN dispatch reserved for PR #25222: + // case BEST_FATTN_KERNEL_ONEDNN: + // ggml_sycl_flash_attn_ext_onednn(ctx, dst); + // break; case BEST_FATTN_KERNEL_HYBRID: ggml_sycl_flash_attn_ext_hybrid(ctx, dst); break; From 66993fde3a8890c60098b570eaf5801644ac4528 Mon Sep 17 00:00:00 2001 From: Ozymandias_EBON Date: Fri, 10 Jul 2026 07:14:25 -0500 Subject: [PATCH 11/11] hybrid: port seq-view stride fix from MKL, env-var one-liner --- ggml/src/ggml-sycl/fattn-hybrid.cpp | 30 +++++++++++++++++++---------- ggml/src/ggml-sycl/fattn.cpp | 5 +---- 2 files changed, 21 insertions(+), 14 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-hybrid.cpp b/ggml/src/ggml-sycl/fattn-hybrid.cpp index 9c32e1d28bf9..d78159cf6221 100644 --- a/ggml/src/ggml-sycl/fattn-hybrid.cpp +++ b/ggml/src/ggml-sycl/fattn-hybrid.cpp @@ -195,13 +195,18 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso ggml_sycl_fattn_alloc K_f16(ctx.fattn_buffers().K); { const char * K_data = (const char *)K->data; - const bool k_interleaved = - ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]) && K->ne[2] > 1; + + // True Gemma interleave packs heads within a row (nb[2] < ne[1]*nb[1]). + // Padded seq-views have nb[2] > ne[1]*nb[1] — use physical strides. + // Both satisfy ne[1]*nb[1] != nb[2] but need different dequant paths. + 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]); K_f16.alloc(ggml_nelements(K)); if (K->type == GGML_TYPE_F16) { - if (!k_interleaved) { + if (!k_non_dense) { stream->memcpy(K_f16.ptr, K_data, ggml_nelements(K) * sizeof(sycl::half)); } else { const int64_t ne0 = K->ne[0], ne1 = K->ne[1]; @@ -219,7 +224,7 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; }); } - } else if (ggml_is_contiguously_allocated(K) && !k_interleaved) { + } else 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_f16.ptr, ggml_nelements(K), stream); } else { @@ -227,7 +232,7 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso 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_interleaved) { + 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; @@ -254,13 +259,18 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso V_f16.ptr = K_f16.ptr; // don't allocate, just alias } else { const char * V_data = (const char *)V->data; - const bool v_interleaved = - ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1; + + // True Gemma interleave packs heads within a row (nb[2] < ne[1]*nb[1]). + // Padded seq-views have nb[2] > ne[1]*nb[1] — use physical strides. + // Both satisfy ne[1]*nb[1] != nb[2] but need different dequant paths. + 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]); V_f16.alloc(ggml_nelements(V)); if (V->type == GGML_TYPE_F16) { - if (!v_interleaved) { + if (!v_non_dense) { stream->memcpy(V_f16.ptr, V_data, ggml_nelements(V) * sizeof(sycl::half)); } else { const int64_t ne0 = V->ne[0], ne1 = V->ne[1]; @@ -278,7 +288,7 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso dst[(hb * ne1 + r) * ne0 + c] = src_row[c]; }); } - } else if (ggml_is_contiguously_allocated(V) && !v_interleaved) { + } else 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_f16.ptr, ggml_nelements(V), stream); } else { @@ -286,7 +296,7 @@ void ggml_sycl_flash_attn_ext_hybrid(ggml_backend_sycl_context & ctx, ggml_tenso 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_interleaved) { + 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; diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index d04e3578fbe8..988149c1225f 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -153,10 +153,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const } // Path 2: MKL GEMM — handles SDPA-incompatible shapes (no mask, ALiBi, etc.). - static int mkl_enable = -1; - if (mkl_enable < 0) { - mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); - } + static int mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1); if (mkl_enable == 1 && Q->ne[0] >= 64 && mask && !sinks && Q->ne[1] >= 32 && K->ne[1] >= 1024 && max_bias == 0.0f && logit_softcap == 0.0f &&