vulkan: add POOL_1D op (#25431 )

* vulkan : add pool1d push constants and pipeline field Declared data structures needed for POOL1D OP, which are the vk_op_pool1d_push_constants struct and pipeline_pool1d_f32 field. * vulkan : add pool1d compute shader Added pool1d.comp for Vulkan backend mirroring the existing pool2d shader. * vulkan : add full GGML_OP_POOL_1D support Added pipeline creation and op dispatch for 1D pooling in the Vulkan backend. * vulkan : fix pool1d shader logic Registered pool1d_f32 in vulkan-shaders-gen.cpp and fixed tensor dimension indices and avg pool scale. * vulkan : fix pool1d end boundary crash and expand test coverage Fixed an issue where the shader crashed when the end boundary was negative when k0 < p0. Also, added more test cases related to this fix.
vulkan: Introduce driver version check for Windows Intel GPU to mitigate crashing (#25192 )
2026-07-31 18:18:18 +02:00 · 2026-07-31 16:48:58 +02:00 · 2026-07-31 16:26:37 +02:00 · 2026-07-31 15:30:19 +02:00 · 2026-07-31 21:06:40 +08:00 · 2026-07-31 14:57:02 +02:00
32 changed files with 1115 additions and 126 deletions
@@ -121,7 +121,7 @@ jobs:
    env:
      OPENBLAS_VERSION: 0.3.23
      SDE_VERSION: 9.33.0-2024-01-07
-      VULKAN_VERSION: 1.4.313.2
+      VULKAN_VERSION: 1.4.357.0

    strategy:
      matrix:
@@ -759,7 +759,7 @@ jobs:

    env:
      OPENBLAS_VERSION: 0.3.23
-      VULKAN_VERSION: 1.4.313.2
+      VULKAN_VERSION: 1.4.357.0

    strategy:
      matrix:
@@ -1620,6 +1620,7 @@ struct llama_model_params common_model_params_to_llama(common_params & params) {
    mparams.progress_callback           = params.load_progress_callback;
    mparams.progress_callback_user_data = params.load_progress_callback_user_data;
    mparams.no_alloc                    = params.no_alloc;
+    mparams.load_mtp                    = std::find(params.speculative.types.begin(), params.speculative.types.end(), COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end();

    return mparams;
 }
@@ -797,6 +797,9 @@ use 1 SYCL GPUs: [0] with Max compute units:512
 | GGML_SYCL_FA_ONEDNN | 1 (default) or 0 | Enable the oneDNN fused SDPA (flash-attention) path on supported GPUs. Set to 0 to always use the native SYCL flash-attention kernel. |
 | GGML_SYCL_FA_ONEDNN_MAX_KV | 0 (default, disabled) or positive integer | By default (0), all sequences are handled by the oneDNN fused SDPA path, regardless of KV length; a positive value caps that length, past which sequences fall back to the native kernel. If GPU driver watchdog resets (DEVICE_LOST) occur during long-context inference, set this near the context depth where they start, e.g. 24576. |
 | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. |
+| GGML_SYCL_ENABLE_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. |
 | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). |
 | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.<br>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. |
@@ -130,43 +130,27 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
        }

        const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + kbx;
-        // Each 32-element chunk occupies 8 bytes of qs (32 elements * 2 bits = 64 bits)
-        const int qs_offset = 8*kqsx;
-        const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) |
-                        (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24);
-        const int qs1 = bxi->qs[qs_offset + 4] | (bxi->qs[qs_offset + 5] << 8) |
-                        (bxi->qs[qs_offset + 6] << 16) | (bxi->qs[qs_offset + 7] << 24);
-
-        // Unpack 32 2-bit codes into 8 int32s, each holding 4 signed int8s in {-1,0,1,2}.
-        int unpacked_bytes[8];
-#pragma unroll
-        for (int j = 0; j < 4; ++j) {
-            const int shift = j * 8;
-            const int codes = (qs0 >> shift) & 0xFF;
-            const int c0 = ((codes >> 0) & 0x3) - 1;
-            const int c1 = ((codes >> 2) & 0x3) - 1;
-            const int c2 = ((codes >> 4) & 0x3) - 1;
-            const int c3 = ((codes >> 6) & 0x3) - 1;
-            unpacked_bytes[j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
-        }
-#pragma unroll
-        for (int j = 0; j < 4; ++j) {
-            const int shift = j * 8;
-            const int codes = (qs1 >> shift) & 0xFF;
-            const int c0 = ((codes >> 0) & 0x3) - 1;
-            const int c1 = ((codes >> 2) & 0x3) - 1;
-            const int c2 = ((codes >> 4) & 0x3) - 1;
-            const int c3 = ((codes >> 6) & 0x3) - 1;
-            unpacked_bytes[4 + j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
-        }
+        const int16_t    * qxi = (const int16_t *) bxi->qs + kqsx * 4;

        const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0;
+
 #pragma unroll
-        for (int j = 0; j < 8; ++j) {
+        for (int j = 0; j < 4; ++j) {
+            const int q  = qxi[j];
+
+            // unpack even and odd crumbs into byte values
+            const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
+            const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2);
+            // unshuffle values
+            const int qx = __byte_perm(qe, qo, 0x5140);
+            const int qy = __byte_perm(qe, qo, 0x7362);
+
 #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
-            x_qs[i*sram_stride           + dst_offset + j] = unpacked_bytes[j];
+            x_qs[i*sram_stride           + dst_offset + j*2+0] = qx;
+            x_qs[i*sram_stride           + dst_offset + j*2+1] = qy;
 #else
-            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j] = unpacked_bytes[j];
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+0] = qx;
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+1] = qy;
 #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
        }
    }
@@ -734,48 +734,28 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
    // Q8_1: 32 elements per block with individual scales
    // iqs selects which of the 2 chunks of 32 elements to process (0-1)

-    const float d2 = bq2_0->d;
+    const float     d2 = bq2_0->d;
+    const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4;

    // Process only the chunk specified by iqs
    const block_q8_1 * bq8_1_chunk = bq8_1 + iqs;

-    // Load 64 bits (8 bytes) for this chunk from Q2_0: bytes [8*iqs, 8*iqs+8)
-    const int offset = iqs * 8;
-    const int v0 = bq2_0->qs[offset + 0] | (bq2_0->qs[offset + 1] << 8) |
-                   (bq2_0->qs[offset + 2] << 16) | (bq2_0->qs[offset + 3] << 24);
-    const int v1 = bq2_0->qs[offset + 4] | (bq2_0->qs[offset + 5] << 8) |
-                   (bq2_0->qs[offset + 6] << 16) | (bq2_0->qs[offset + 7] << 24);
-
-    // Unpack 32 2-bit codes into 8 int32s, each holding 4 signed int8 symbols in {-1,0,1,2}.
-    // Stored code c in {0,1,2,3} -> symbol s = c - 1.
-    int vi_bytes[8];
-#pragma unroll
-    for (int j = 0; j < 4; ++j) {
-        const int shift = j * 8;
-        const int codes = (v0 >> shift) & 0xFF;
-        const int c0 = ((codes >> 0) & 0x3) - 1;
-        const int c1 = ((codes >> 2) & 0x3) - 1;
-        const int c2 = ((codes >> 4) & 0x3) - 1;
-        const int c3 = ((codes >> 6) & 0x3) - 1;
-        vi_bytes[j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
-    }
-#pragma unroll
-    for (int j = 0; j < 4; ++j) {
-        const int shift = j * 8;
-        const int codes = (v1 >> shift) & 0xFF;
-        const int c0 = ((codes >> 0) & 0x3) - 1;
-        const int c1 = ((codes >> 2) & 0x3) - 1;
-        const int c2 = ((codes >> 4) & 0x3) - 1;
-        const int c3 = ((codes >> 6) & 0x3) - 1;
-        vi_bytes[4 + j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
-    }
-
-    // Compute dot product for this 32-element chunk
    int sumi = 0;
 #pragma unroll
-    for (int j = 0; j < 8; ++j) {
-        const int u = get_int_b4(bq8_1_chunk->qs, j);
-        sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi);
+    for (int j = 0; j < 4; ++j) {
+        const int q  = qs[j];
+        const int u  = get_int_b4(bq8_1_chunk->qs, j*2+0);
+        const int v  = get_int_b4(bq8_1_chunk->qs, j*2+1);
+
+        // unpack even and odd crumbs into byte values
+        const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
+        const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2);
+        // unshuffle values
+        const int qx = __byte_perm(qe, qo, 0x5140);
+        const int qy = __byte_perm(qe, qo, 0x7362);
+
+        sumi = ggml_cuda_dp4a(u, qx, sumi);
+        sumi = ggml_cuda_dp4a(v, qy, sumi);
    }

    // Apply Q2_0's single scale and this chunk's Q8_1 scale
@@ -0,0 +1,690 @@
+// 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 <oneapi/mkl.hpp>
+#include <cstdio>
+#include <chrono>
+
+#define MKL_FA_CHUNK_SIZE_KV 8192
+
+// Number of query rows processed per tile. The score buffers (KQ_f32, S_f16)
+// are sized q_tile_rows * chunk_size, so this bounds their footprint
+// regardless of batch size (n_query_rows = n_queries * gqa_ratio). A typical
+// single-ubatch prefill (e.g. ubatch 1024 * gqa 8 = 8192 rows) is exactly one
+// tile, so it runs with no extra iterations. Larger batches tile and stay
+// bounded. Override with GGML_SYCL_MKL_FA_Q_TILE.
+#define MKL_FA_Q_TILE 8192
+
+#define MKL_FA_WG_SIZE 256
+
+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.
+// 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 q_row_stride, int64_t q_head_stride,
+    int64_t 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;
+
+        const int64_t n_elem = (int64_t)n_queries * DKQ;
+        const int64_t wg = ((n_elem + 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 >= n_elem) return;
+
+                    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.
+// 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 a tile of GQA query rows.
+// The tile spans absolute rows [q0, q0 + q_rows). Score buffers
+// (KQ_f32/S_f16) are indexed RELATIVE to the tile; the persistent state
+// (VKQ_accum/KQ_max/KQ_sum) and mask are indexed by ABSOLUTE row.
+// For each row: find local max → rescale previous VKQ_accum →
+// compute exp(s - max) → write S_f16 → update running max/sum.
+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 q0, int q_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 = ((q_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_rel = item.get_global_id(0);
+                if (jc_rel >= q_rows) return;
+                int jc_abs = q0 + jc_rel;
+
+                const int gqa_group = jc_abs / n_queries;
+                const int q_row     = jc_abs % n_queries;
+
+                // Score buffers are tile-local (relative index).
+                const float * __restrict KQ_row = KQ_f32
+                    + jc_rel * (int64_t)chunk_size;
+                // Persistent accumulator is full-sized (absolute index).
+                float * __restrict vkq = VKQ_accum
+                    + jc_abs * (int64_t)DV;
+
+                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 (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)];
+                    }
+                    if (s > local_max) local_max = s;
+                }
+
+                // Rescale previous accumulator by exp(old_max - new_max)
+                float old_max = KQ_max[jc_abs];
+                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 (tile-local index)
+                float local_sum = 0.0f;
+                sycl::half * __restrict S_row = S_f16
+                    + jc_rel * (int64_t)chunk_size;
+
+                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)];
+                    }
+                    float val = sycl::native::exp(s - new_max);
+                    S_row[i] = sycl::half(val);
+                    local_sum += val;
+                }
+
+                KQ_sum[jc_abs] = KQ_sum[jc_abs] * rescale + local_sum;
+                KQ_max[jc_abs] = 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, 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;
+
+    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;
+                // 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
+                    + ((int64_t)jc * n_q_heads + iqh) * (int64_t)DV;
+
+                for (int v = 0; v < DV; v++) {
+                    dst_row[v] = src[v] * inv_sum;
+                }
+            });
+    });
+}
+
+// ---------------------------------------------------------------------------
+// Per-chunk dequant
+//
+// Rather than dequantizing all of K/V up front (footprint scales with
+// context), we dequant one KV-head chunk at a time into a dense
+// [this_chunk x D] fp16 buffer (row-major, lda = D). The source address of
+// element (head=ikvh, row=chunk_start+r, col=c) decomposes into independent
+// linear terms head_off(ikvh) + row_off(chunk_start) + (r,c), so slicing a
+// chunk is a clean pointer offset in every layout case. The true-Gemma-
+// interleave vs padded-seq-view distinction is resolved once when the
+// descriptor is built; slicing does not reintroduce it.
+// ---------------------------------------------------------------------------
+enum mkl_fa_kv_desc_mode {
+    MKL_FA_KV_MODE_F16_DENSE       = 0,
+    MKL_FA_KV_MODE_F16_INTERLEAVED = 1,
+    MKL_FA_KV_MODE_QUANT_CONTIG    = 2,
+    MKL_FA_KV_MODE_QUANT_NC        = 3,
+};
+
+struct mkl_fa_kv_desc {
+    const char *         data = nullptr;
+    ggml_type            type = GGML_TYPE_F16;
+    int64_t              D    = 0;      // ne[0]
+    int64_t              nb1  = 0;      // byte stride, seq dim
+    int64_t              nb2  = 0;      // byte stride, head dim
+    mkl_fa_kv_desc_mode  mode = MKL_FA_KV_MODE_F16_DENSE;
+    int64_t              ts   = 0;      // type size (mode 3 base offset)
+    int64_t              s01  = 0;      // nc row stride in blocks (mode 3)
+    int64_t              s02  = 0;      // nc head stride in blocks (mode 3)
+};
+
+static mkl_fa_kv_desc mkl_fa_make_desc(const ggml_tensor * T, bool interleaved, int n_kv_heads) {
+    mkl_fa_kv_desc d;
+    d.data = (const char *)T->data;
+    d.type = T->type;
+    d.D    = T->ne[0];
+    d.nb1  = (int64_t)T->nb[1];
+    d.nb2  = (int64_t)T->nb[2];
+    d.ts   = (int64_t)ggml_type_size(T->type);
+
+    if (T->type == GGML_TYPE_F16) {
+        d.mode = interleaved ? MKL_FA_KV_MODE_F16_INTERLEAVED
+                             : MKL_FA_KV_MODE_F16_DENSE;
+    } else if (ggml_is_contiguously_allocated(T) && !interleaved) {
+        d.mode = MKL_FA_KV_MODE_QUANT_CONTIG;
+    } else {
+        d.mode = MKL_FA_KV_MODE_QUANT_NC;
+        const int64_t bs          = (int64_t)ggml_blck_size(T->type);
+        const int64_t blk_per_row = T->ne[0] / bs;
+        // True Gemma interleave packs heads within a row (nb[2] < ne[1]*nb[1])
+        // → reconstruct physical strides. Padded seq-views (nb[2] > ne[1]*nb[1])
+        // already have correct physical strides.
+        const bool gemma = interleaved &&
+            ((int64_t)T->nb[2] < (int64_t)T->ne[1] * (int64_t)T->nb[1]);
+        if (gemma) {
+            d.s01 = (int64_t)n_kv_heads * blk_per_row;
+            d.s02 = blk_per_row;
+        } else {
+            d.s01 = d.nb1 / d.ts;
+            d.s02 = d.nb2 / d.ts;
+        }
+    }
+    return d;
+}
+
+// Dequant one KV-head chunk into a dense [this_chunk x D] fp16 buffer.
+static void mkl_fa_dequant_chunk(
+    dpct::queue_ptr stream, const mkl_fa_kv_desc & d, ggml_tensor * dst_ctx,
+    sycl::half * out, int ikvh, int chunk_start, int this_chunk) {
+
+    const int64_t D = d.D;
+    switch (d.mode) {
+        case MKL_FA_KV_MODE_F16_DENSE: {
+            const char * base = d.data + (int64_t)ikvh * d.nb2
+                + (int64_t)chunk_start * d.nb1;
+            stream->memcpy(out, base, (size_t)this_chunk * D * sizeof(sycl::half));
+            break;
+        }
+        case MKL_FA_KV_MODE_F16_INTERLEAVED: {
+            const char * base = d.data + (int64_t)ikvh * d.nb2
+                + (int64_t)chunk_start * d.nb1;
+            const int64_t row_halfs = d.nb1 / (int64_t)sizeof(sycl::half);
+            const sycl::half * src = (const sycl::half *)base;
+            stream->parallel_for(
+                sycl::range<2>((size_t)this_chunk, (size_t)D),
+                [=](sycl::item<2> it) {
+                    int64_t r = it.get_id(0);
+                    int64_t c = it.get_id(1);
+                    out[r * D + c] = src[r * row_halfs + c];
+                });
+            break;
+        }
+        case MKL_FA_KV_MODE_QUANT_CONTIG: {
+            const char * base = d.data + (int64_t)ikvh * d.nb2
+                + (int64_t)chunk_start * d.nb1;
+            to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(d.type, dst_ctx);
+            to_fp16(base, out, (int64_t)this_chunk * D, stream);
+            break;
+        }
+        default: {  // MKL_FA_KV_MODE_QUANT_NC
+            to_fp16_nc_sycl_t to_fp16 = ggml_get_to_fp16_nc_sycl(d.type);
+            const int64_t base_blocks = (int64_t)ikvh * d.s02
+                + (int64_t)chunk_start * d.s01;
+            const char * base = d.data + base_blocks * d.ts;
+            // ne02 = ne03 = 1 → s02/s03 inert; head+chunk offset carried by base.
+            to_fp16(base, out, D, this_chunk, 1, 1, d.s01, d.s02, d.s02, stream);
+            break;
+        }
+    }
+}
+
+// ---------------------------------------------------------------------------
+// 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];
+    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 float q_scale = 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);
+
+    // Query rows are processed in tiles of q_tile_rows so the score buffers
+    // (KQ_f32/S_f16 = q_tile_rows * chunk_size) stay bounded regardless of
+    // batch size. n_query_rows <= Q_TILE is a single tile (no extra work).
+    static int q_tile_env = ggml_sycl_get_env("GGML_SYCL_MKL_FA_Q_TILE", MKL_FA_Q_TILE);
+    const int q_tile_rows = std::max(1, std::min(q_tile_env, n_query_rows));
+
+    const int64_t wg_size = MKL_FA_WG_SIZE;
+
+    // --- Debug output (gated by GGML_SYCL_MKL_FA_DEBUG=1) ---
+    static int mkl_call_count = 0;
+    mkl_call_count++;
+    static int mkl_debug = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DEBUG", 0);
+    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) {
+        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,
+                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),
+                k_early_interleaved ? " K_ILV" : "",
+                v_early_interleaved ? " V_ILV" : "");
+        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)));
+    }
+
+    // --- Stream and allocators ---
+    dpct::queue_ptr stream = ctx.stream();
+
+#define MKL_TAKE_TIME(t0)  auto t0 = std::chrono::steady_clock::now()
+#define MKL_ACCUM(acc, t0) do { if (do_print) { \
+    acc += (int64_t)std::chrono::duration_cast \
+    <std::chrono::microseconds>(std::chrono::steady_clock::now() - (t0)).count(); \
+} } while(0)
+
+    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);
+
+    // --- K/V dequant descriptors ---
+    // Dequant is done per-chunk inside the KV loop (footprint independent of
+    // context). Output is always dense row-major fp16 [this_chunk x D], lda=D.
+    // Interleaved detection: ne[1]*nb[1] != nb[2] means heads are interleaved.
+    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;
+
+    const mkl_fa_kv_desc K_desc = mkl_fa_make_desc(K, k_interleaved, n_kv_heads);
+    const mkl_fa_kv_desc V_desc = V_is_K_view
+        ? K_desc : mkl_fa_make_desc(V, v_interleaved, n_kv_heads);
+
+    MKL_ACCUM(dequant_time_us, t_deq);
+
+    // --- 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) {
+        // 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 from pool ---
+    ggml_sycl_pool & pool = ctx.pool();
+
+    ggml_sycl_pool_alloc<float>      KQ_f32(pool);      // [q_tile_rows x chunk]
+    ggml_sycl_pool_alloc<sycl::half> S_f16(pool);       // [q_tile_rows x chunk]
+    ggml_sycl_pool_alloc<float>      VKQ_chunk(pool);   // [q_tile_rows x DV]
+    ggml_sycl_pool_alloc<float>      VKQ_accum(pool);   // [n_query_rows x DV] (full)
+    ggml_sycl_pool_alloc<float>      KQ_max(pool);      // [n_query_rows] (full)
+    ggml_sycl_pool_alloc<float>      KQ_sum(pool);      // [n_query_rows] (full)
+    ggml_sycl_pool_alloc<sycl::half> Q_head_f16(pool);  // [n_query_rows x DKQ] (full)
+    ggml_sycl_pool_alloc<sycl::half> K_chunk_f16(pool); // [chunk x DKQ] (per-chunk dequant)
+    ggml_sycl_pool_alloc<sycl::half> V_chunk_f16(pool); // [chunk x DV] (per-chunk dequant)
+
+    KQ_f32.alloc((size_t)q_tile_rows * chunk_size);
+    S_f16.alloc((size_t)q_tile_rows * chunk_size);
+    VKQ_chunk.alloc((size_t)q_tile_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);
+    K_chunk_f16.alloc((size_t)chunk_size * DKQ);
+
+    sycl::half * V_chunk_f16_ptr;
+    if (V_is_K_view) {
+        V_chunk_f16_ptr = K_chunk_f16.ptr;   // V aliases K (DV == DKQ)
+    } else {
+        V_chunk_f16.alloc((size_t)chunk_size * DV);
+        V_chunk_f16_ptr = V_chunk_f16.ptr;
+    }
+
+    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;
+    sycl::half * K_chunk_f16_ptr = K_chunk_f16.ptr;
+
+    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));
+
+        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] / 2);  // 2 = actual fp16 device size
+        }
+
+        for (int ikvh = 0; ikvh < n_kv_heads; ikvh++) {
+            int kvh_base_head = ikvh * gqa_ratio;
+
+            // 1. Pack all GQA Q heads into fp16 (full n_query_rows)
+            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, q_row_stride, q_head_stride, wg_size);
+
+            // 2. Initialize softmax state (full n_query_rows)
+            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();
+
+            // 3. KV chunk loop (OUTER): dequant each chunk once, then tile queries.
+            for (int chunk_start = 0; chunk_start < n_kv; chunk_start += chunk_size) {
+                int this_chunk = std::min(chunk_size, n_kv - chunk_start);
+
+                // 3a. Dequant this KV chunk to dense fp16 (once per chunk)
+                {
+                    MKL_TAKE_TIME(t0);
+                    mkl_fa_dequant_chunk(stream, K_desc, KQV,
+                        K_chunk_f16_ptr, ikvh, chunk_start, this_chunk);
+                    if (!V_is_K_view) {
+                        mkl_fa_dequant_chunk(stream, V_desc, KQV,
+                            V_chunk_f16_ptr, ikvh, chunk_start, this_chunk);
+                    }
+                    stream->wait();  // dequant must be ready before MKL GEMM
+                    MKL_ACCUM(dequant_time_us, t0);
+                }
+
+                // 3b. Query tile loop (INNER) — bounds KQ_f32/S_f16 footprint.
+                for (int q0 = 0; q0 < n_query_rows; q0 += q_tile_rows) {
+                    int q_rows = std::min(q_tile_rows, n_query_rows - q0);
+
+                    // GEMM: KQ = Q_tile × K_chunk^T
+                    {
+                        MKL_TAKE_TIME(t0);
+                        sycl::event ev = gemm(*stream,
+                            transpose::trans, transpose::nontrans,
+                            this_chunk, q_rows, DKQ,
+                            alpha,
+                            K_chunk_f16_ptr, DKQ,
+                            Q_head_f16_ptr + (int64_t)q0 * DKQ, DKQ,
+                            beta,
+                            KQ_f32_ptr, this_chunk);
+                        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");
+                        }
+                        MKL_ACCUM(gemm_kq_time_us, t0);
+                    }
+                    // Online softmax over this chunk for this query tile
+                    {
+                        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,
+                            q0, q_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);
+                    }
+
+                    // GEMM: VKQ_chunk = S × V_chunk
+                    {
+                        MKL_TAKE_TIME(t0);
+                        sycl::event ev = gemm(*stream,
+                            transpose::nontrans, transpose::nontrans,
+                            DV, q_rows, this_chunk,
+                            alpha,
+                            V_chunk_f16_ptr, DV,
+                            S_f16_ptr, this_chunk,
+                            beta,
+                            VKQ_chunk_ptr, DV);
+                        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");
+                        }
+                        MKL_ACCUM(gemm_vkq_time_us, t0);
+                    }
+                    // VKQ_accum[q0..] += VKQ_chunk
+                    {
+                        const int64_t n_total = (int64_t)q_rows * DV;
+                        const int64_t wg = ((n_total + wg_size - 1) / wg_size)
+                            * wg_size;
+                        float * accum = VKQ_accum_ptr + (int64_t)q0 * DV;
+                        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) {
+                                        accum[i] += VKQ_chunk_ptr[i];
+                                    }
+                                });
+                        });
+                    }
+                }
+            }
+
+            // 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, n_q_heads,
+                    src_offset, wg_size);
+            }
+        }
+    }
+
+#undef MKL_TAKE_TIME
+#undef MKL_ACCUM
+
+    if (do_print) {
+        const int64_t v_chunk_elems = V_is_K_view ? 0 : (int64_t)chunk_size * DV;
+        double total_mb = (double)(
+            (int64_t)q_tile_rows * chunk_size * sizeof(float)      // KQ_f32
+          + (int64_t)q_tile_rows * chunk_size * sizeof(sycl::half) // S_f16
+          + (int64_t)q_tile_rows * DV * sizeof(float)              // VKQ_chunk
+          + (int64_t)n_query_rows * DV * sizeof(float)             // VKQ_accum
+          + (int64_t)n_query_rows * sizeof(float)                  // KQ_max
+          + (int64_t)n_query_rows * sizeof(float)                  // KQ_sum
+          + (int64_t)n_query_rows * DKQ * sizeof(sycl::half)       // Q_head_f16
+          + (int64_t)chunk_size * DKQ * sizeof(sycl::half)         // K_chunk_f16
+          + v_chunk_elems * (int64_t)sizeof(sycl::half)            // V_chunk_f16
+        ) / (1024.0 * 1024.0);
+        GGML_LOG_INFO("[MKL-FA] #%d n_kv=%d n_q=%d q_tile=%d time_us: "
+                "dequant=%lld GEMM_KQ=%lld softmax=%lld GEMM_VKQ=%lld "
+                "buf_mb=%.1f\n",
+                mkl_call_count, n_kv, n_queries, q_tile_rows,
+                (long long)dequant_time_us,
+                (long long)gemm_kq_time_us,
+                (long long)softmax_time_us,
+                (long long)gemm_vkq_time_us,
+                total_mb);
+    }
+}
@@ -99,8 +99,10 @@ enum best_fattn_kernel {
    BEST_FATTN_KERNEL_VEC      = 100,
    BEST_FATTN_KERNEL_ONEDNN   = 150, // added enum for onednn==150
    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) {
    GGML_UNUSED(device);
 #ifndef SYCL_FLASH_ATTN
@@ -115,6 +117,7 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const
    const ggml_tensor * K     = dst->src[1];
    const ggml_tensor * V     = dst->src[2];
    const ggml_tensor * mask  = dst->src[3];
+    const ggml_tensor * sinks = dst->src[4];

    const int gqa_ratio = Q->ne[2] / K->ne[2];
    GGML_ASSERT(Q->ne[2] % K->ne[2] == 0);
@@ -122,7 +125,49 @@ 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. 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.
+    static int mkl_enable = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1);
+    // MKL is validated for the mainstream GQA envelope: grouped-query
+    // (gqa_ratio >= 2), head_dim a multiple of 64 in [64,512] with matching
+    // K/V head size, mask, no sinks/ALiBi/softcap. Gemma's global layers use
+    // head_dim 512, so the cap must include it. Head sizes not a multiple of
+    // 64 (72/80/96), MHA (gqa_ratio == 1), and MLA (DKQ != DV, e.g. 576/512)
+    // fall through to TILE/VEC; see follow-up work.
+    if (mkl_enable == 1 && mask && !sinks && gqa_ratio >= 2 &&
+        Q->ne[0] >= 64 && Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 &&
+        Q->ne[0] == V->ne[0] &&
+        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}) {
        if (t == nullptr || ggml_is_quantized(t->type)) {
            continue;
@@ -216,6 +261,37 @@ 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 = 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]);
+    }
+
    switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) {
        case BEST_FATTN_KERNEL_NONE:
            GGML_ABORT("Not support Flash-Attention");
@@ -232,6 +308,51 @@ 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;
+    }
+
+    // --- 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 = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DIAG", 0);
+    static int fa_diag_count = 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]);
+            }
+        }
    }
 }

@@ -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
@@ -998,6 +998,7 @@ struct vk_device_struct {
    vk_pipeline pipeline_snake_f32;
    vk_pipeline pipeline_snake_f16;
    vk_pipeline pipeline_snake_bf16;
+    vk_pipeline pipeline_pool1d_f32;
    vk_pipeline pipeline_pool2d_f32;
    vk_pipeline pipeline_rwkv_wkv6_f32;
    vk_pipeline pipeline_rwkv_wkv7_f32;
@@ -1687,6 +1688,17 @@ struct vk_op_snake_push_constants {
    uint32_t ne1;
 };

+struct vk_op_pool1d_push_constants {
+    uint32_t IL;
+    uint32_t OL;
+    uint32_t OC;
+    uint32_t pelements;
+    uint32_t op;
+    int32_t k0;
+    int32_t s0;
+    int32_t p0;
+};
+
 struct vk_op_pool2d_push_constants {
    uint32_t IW; uint32_t IH;
    uint32_t OW; uint32_t OH;
@@ -1934,6 +1946,7 @@ struct ggml_vk_garbage_collector {
 static void ggml_vk_preallocate_buffers(ggml_backend_vk_context * ctx, vk_context subctx);
 static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested = nullptr);
 static void ggml_pipeline_allocate_descriptor_sets(ggml_backend_vk_context * ctx);
+static bool ggml_vk_intel_windows_driver_equals_or_newer_than(uint32_t driver_version, uint32_t threshold_major, uint32_t threshold_minor);

 static bool vk_memory_logger_enabled = false;

@@ -5552,8 +5565,11 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
    ggml_vk_create_pipeline(device, device->pipeline_argmax_f32, "argmax_f32", argmax_f32_len, argmax_f32_data, "main", 2, sizeof(vk_op_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);

    ggml_vk_create_pipeline(device, device->pipeline_sum_rows_f32, "sum_rows_f32", sum_rows_f32_len, sum_rows_f32_data, "main", 2, sizeof(vk_op_sum_rows_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);
-    // Intel Arc B390 was observed segfaulting with this shader.
-    if (device->subgroup_basic && device->subgroup_shuffle && device->vendor_id != VK_VENDOR_ID_INTEL) {
+    // Intel Windows driver older than 32.0.101.8860 will crash when using fwht kernels on Xe2+ GPUS so we gate that here
+    const bool can_use_fwht = device->driver_id != vk::DriverId::eIntelProprietaryWindows ||
+        device->architecture != vk_device_architecture::INTEL_XE2 ||
+        (device->architecture == vk_device_architecture::INTEL_XE2 && ggml_vk_intel_windows_driver_equals_or_newer_than(device->properties.driverVersion, 101, 8860));
+    if (can_use_fwht && device->subgroup_basic && device->subgroup_shuffle) {
        int idx = 0;
        for (uint32_t n : {64, 128, 256, 512}) {
            if (device->subgroup_size <= n) {
@@ -5561,8 +5577,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
            }
            ++idx;
        }
-    } else if (device->driver_id != vk::DriverId::eIntelProprietaryWindows) {
-        // Disabled on Intel Windows due to a driver bug: https://github.com/ggml-org/llama.cpp/pull/23964#issuecomment-4598226147
+    } else if (can_use_fwht) {
        int idx = 0;
        for (uint32_t n : {64, 128, 256, 512}) {
            const uint32_t block_size = std::min(device->subgroup_size, n);
@@ -5619,6 +5634,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
    ggml_vk_create_pipeline(device, device->pipeline_snake_f16,  "snake_f16",  snake_f16_len,  snake_f16_data,  "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1);
    ggml_vk_create_pipeline(device, device->pipeline_snake_bf16, "snake_bf16", snake_bf16_len, snake_bf16_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1);

+    ggml_vk_create_pipeline(device, device->pipeline_pool1d_f32, "pool1d_f32", pool1d_f32_len, pool1d_f32_data, "main", 2, sizeof(vk_op_pool1d_push_constants), {512, 1, 1}, {}, 1);
    ggml_vk_create_pipeline(device, device->pipeline_pool2d_f32, "pool2d_f32", pool2d_f32_len, pool2d_f32_data, "main", 2, sizeof(vk_op_pool2d_push_constants), {512, 1, 1}, {}, 1);

    ggml_vk_create_pipeline(device, device->pipeline_rwkv_wkv6_f32, "rwkv_wkv6_f32", rwkv_wkv6_f32_len, rwkv_wkv6_f32_data, "main", 7, sizeof(vk_op_rwkv_wkv6_push_constants), {1, 1, 1}, {device->subgroup_size}, 1);
@@ -11332,6 +11348,11 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
            case GGML_TYPE_BF16: return ctx->device->pipeline_col2im_1d_bf16;
            default:             return nullptr;
        }
+    case GGML_OP_POOL_1D:
+        if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+            return ctx->device->pipeline_pool1d_f32;
+        }
+        return nullptr;
    case GGML_OP_POOL_2D:
        if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
            return ctx->device->pipeline_pool2d_f32;
@@ -11854,6 +11875,13 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
        {
            elements = { uint32_t(dst->ne[0]), uint32_t(dst->ne[1]), 1 };
        } break;
+    case GGML_OP_POOL_1D:
+        {
+            const uint32_t N = dst->ne[3] * dst->ne[2];
+            const uint32_t OC = dst->ne[1];
+            const uint32_t OL = dst->ne[0];
+            elements = { N * OC * OL, 1, 1};
+        } break;
    case GGML_OP_POOL_2D:
        {
            const uint32_t N = dst->ne[3];
@@ -13773,6 +13801,29 @@ static void ggml_vk_snake_dispatch_fused(ggml_backend_vk_context * ctx, vk_conte
    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, a_buf, inv_b_buf, dst_buf }, pc, elements);
 }

+static void ggml_vk_pool_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
+    uint32_t op = static_cast<uint32_t>(dst->op_params[0]);
+    const int32_t k0 = dst->op_params[1];
+    const int32_t s0 = dst->op_params[2];
+    const int32_t p0 = dst->op_params[3];
+
+    const uint32_t IL = src0->ne[0];
+
+    const uint32_t N = dst->ne[3] * dst->ne[2];
+
+    const uint32_t OC = dst->ne[1];
+    const uint32_t OL = dst->ne[0];
+
+    const uint32_t parallel_elements = N * OC * OL;
+
+    ggml_vk_op_f32<vk_op_pool1d_push_constants>(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_POOL_1D, {
+        IL, OL, OC,
+        parallel_elements,
+        op,
+        k0, s0, p0,
+    });
+}
+
 static void ggml_vk_pool_2d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
    uint32_t op = static_cast<uint32_t>(dst->op_params[0]);
    const int32_t k1 = dst->op_params[1];
@@ -15284,6 +15335,10 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr
    case GGML_OP_CONV_TRANSPOSE_1D:
        ggml_vk_conv_transpose_1d(ctx, compute_ctx, src0, src1, node);

+        break;
+    case GGML_OP_POOL_1D:
+        ggml_vk_pool_1d(ctx, compute_ctx, src0, node);
+
        break;
    case GGML_OP_POOL_2D:
        ggml_vk_pool_2d(ctx, compute_ctx, src0, node);
@@ -18001,6 +18056,8 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
        case GGML_OP_CONV_2D_DW:
            return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16)
                && op->src[1]->type == GGML_TYPE_F32;
+        case GGML_OP_POOL_1D:
+            return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
        case GGML_OP_POOL_2D:
            return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
        case GGML_OP_RWKV_WKV6:
@@ -18466,6 +18523,22 @@ static uint32_t ggml_vk_intel_shader_core_count(const vk::PhysicalDevice& vkdev)
    }
 }

+static bool ggml_vk_intel_windows_driver_equals_or_newer_than(uint32_t driver_version, uint32_t threshold_major, uint32_t threshold_minor) {
+#if defined(_WIN32)
+    // Intel Windows encodes xxx.yyyy as [31:14].[13:0].
+    const uint32_t major = driver_version >> 14;
+    const uint32_t minor = driver_version & 0x3fff;
+
+    return major > threshold_major || (major == threshold_major && minor >= threshold_minor);
+#else
+    GGML_UNUSED(driver_version);
+    GGML_UNUSED(threshold_major);
+    GGML_UNUSED(threshold_minor);
+    return true;
+#endif
+}
+
+
 // checks

 #ifdef GGML_VULKAN_CHECK_RESULTS
@@ -18920,6 +18993,13 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
            const int32_t oc     = tensor->op_params[1];
            const int32_t p0     = tensor->op_params[2];
            tensor_clone = ggml_col2im_1d(ggml_ctx, src_clone[0], stride, oc, p0);
+        } else if (tensor->op == GGML_OP_POOL_1D) {
+            enum ggml_op_pool op = static_cast<ggml_op_pool>(tensor->op_params[0]);
+            const int32_t k0 = tensor->op_params[1];
+            const int32_t s0 = tensor->op_params[2];
+            const int32_t p0 = tensor->op_params[3];
+
+            tensor_clone = ggml_pool_1d(ggml_ctx, src_clone[0], op, k0, s0, p0);
        } else if (tensor->op == GGML_OP_POOL_2D) {
            enum ggml_op_pool op = static_cast<ggml_op_pool>(tensor->op_params[0]);
            const int32_t k0 = tensor->op_params[1];
@@ -0,0 +1,65 @@
+#version 450
+
+#include "types.glsl"
+
+#extension GL_EXT_shader_16bit_storage : require
+
+layout(push_constant) uniform parameter {
+    uint IL;
+    uint OL;
+    uint OC;
+    uint pelements;
+    uint op;
+    int k0;
+    int s0;
+    int p0;
+} p;
+
+#define BLOCK_SIZE 512
+#define FLT_MAX 3.402823466e+38F
+#define OP_POOL_MAX 0u
+#define OP_POOL_AVG 1u
+
+layout (local_size_x = BLOCK_SIZE, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0) readonly buffer X {A_TYPE data_a[];};
+layout(binding = 1) writeonly buffer D {D_TYPE data_d[];};
+
+void main() {
+    const uint idx = gl_GlobalInvocationID.x;
+    if (idx >= p.pelements) {
+        return;
+    }
+
+    const uint nc = idx / p.OL;
+    const uint cur_ol = idx % p.OL;
+
+    const int start = int(cur_ol) * p.s0 - p.p0;
+    const int bl = max(start, 0);
+    const int el = min(max(start + p.k0, 0), int(p.IL));
+
+    const int window_size = el - bl;
+    const float scale = window_size > 0 ? 1.0 / float(window_size) : 0.0;
+    float res;
+
+    if (p.op == OP_POOL_AVG) {
+        res = 0.0;
+    } else if (p.op == OP_POOL_MAX) {
+        res = -FLT_MAX;
+    } else {
+        return;
+    }
+
+    #pragma unroll
+    for (uint i = bl; i < el; i++) {
+        const float cur = D_TYPE(data_a[nc * p.IL + i]);
+
+        if (p.op == OP_POOL_AVG) {
+            res += cur * scale;
+        } else if (p.op == OP_POOL_MAX) {
+            res = max(res, cur);
+        }
+    }
+
+    data_d[nc * p.OL + cur_ol] = res;
+}
@@ -1052,6 +1052,7 @@ void process_shaders() {
    string_to_spv("snake_f16",  "snake.comp", {{"DATA_A_F16", "1"},  {"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}});
    string_to_spv("snake_bf16", "snake.comp", {{"DATA_A_BF16", "1"}, {"DATA_D_BF16", "1"}, {"A_TYPE", "uint16_t"},  {"D_TYPE", "uint16_t"}});

+    string_to_spv("pool1d_f32", "pool1d.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
    string_to_spv("pool2d_f32", "pool2d.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));

    string_to_spv("rwkv_wkv6_f32", "wkv6.comp", merge_maps(base_dict, {{"A_TYPE", "float"}}));
@@ -353,6 +353,7 @@ class Keys:
        class Attention:
            HEAD_COUNT      = "clip.vision.attention.head_count"
            HEAD_COUNT_KV   = "clip.vision.attention.head_count_kv" # used by mimovl (GQA)
+            HEAD_DIM        = "clip.vision.attention.head_dim" # set when qkv width != n_embd
            LAYERNORM_EPS   = "clip.vision.attention.layer_norm_epsilon"

        class Projector:
@@ -1226,6 +1226,9 @@ class GGUFWriter:
    def add_vision_head_count_kv(self, value: int) -> None:
        self.add_uint32(Keys.ClipVision.Attention.HEAD_COUNT_KV, value)

+    def add_vision_head_dim(self, value: int) -> None:
+        self.add_uint32(Keys.ClipVision.Attention.HEAD_DIM, value)
+
    def add_vision_attention_layernorm_eps(self, value: float) -> None:
        self.add_float32(Keys.ClipVision.Attention.LAYERNORM_EPS, value)

@@ -337,6 +337,7 @@ extern "C" {
        bool use_extra_bufts; // use extra buffer types (used for weight repacking)
        bool no_host;         // bypass host buffer allowing extra buffers to be used
        bool no_alloc;        // only load metadata and simulate memory allocations
+        bool load_mtp;        // whether to load MTP layers
    };

    struct llama_sampler_seq_config {
@@ -526,6 +526,7 @@ llama_model_loader::llama_model_loader(
        llama_load_mode load_mode,
        bool check_tensors,
        bool no_alloc,
+        bool load_mtp,
        const llama_model_kv_override * param_overrides_p,
        const llama_model_tensor_buft_override * param_tensor_buft_overrides_p)
        : metadata(meta), set_tensor_data(set_tensor_data), set_tensor_data_ud(set_tensor_data_ud) {
@@ -812,6 +813,7 @@ llama_model_loader::llama_model_loader(

    this->check_tensors = check_tensors;
    this->no_alloc = no_alloc;
+    this->load_mtp = load_mtp;
 }

 std::string llama_model_loader::get_arch_name() const {
@@ -79,6 +79,7 @@ struct llama_model_loader {
    bool use_direct_io = false;
    bool check_tensors;
    bool no_alloc;
+    bool load_mtp;

    llama_files files;
    llama_ftype ftype;
@@ -129,6 +130,7 @@ struct llama_model_loader {
        llama_load_mode load_mode,
        bool check_tensors,
        bool no_alloc,
+        bool load_mtp,
        const llama_model_kv_override * param_overrides_p,
        const llama_model_tensor_buft_override * param_tensor_buft_overrides_p);

@@ -2390,6 +2390,7 @@ llama_model_params llama_model_default_params() {
        /*.use_extra_bufts             =*/ true,
        /*.no_host                     =*/ false,
        /*.no_alloc                    =*/ false,
+        /*.load_mtp                    =*/ false,
    };

    return result;
@@ -893,7 +893,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
    const llama_model_kv_override * kv_overrides = params->kv_overrides;
    std::vector<std::string> splits = {};
    llama_model_loader ml(/*metadata*/ nullptr, /*set_tensor_data*/ nullptr, /*set_tensor_data_ud*/ nullptr,
-        fname_inp, splits, /*file*/ nullptr, /*load_mode*/ load_mode, /*check_tensors*/ true, /*no_alloc*/ false, kv_overrides, nullptr);
+        fname_inp, splits, /*file*/ nullptr, /*load_mode*/ load_mode, /*check_tensors*/ true, /*no_alloc*/ false, /*load_mtp*/ true, kv_overrides, nullptr);
    ml.init_mappings(false); // no prefetching

    auto mparams = llama_model_default_params();
@@ -305,7 +305,7 @@ static std::pair<int, llama_model *> llama_model_load(struct gguf_context * meta
        const std::string & fname, std::vector<std::string> & splits, FILE * file, llama_model_params & params) {
    try {
        llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, file, params.load_mode,
-            params.check_tensors, params.no_alloc, params.kv_overrides, params.tensor_buft_overrides);
+            params.check_tensors, params.no_alloc, params.load_mtp, params.kv_overrides, params.tensor_buft_overrides);

        ml.print_info();
        std::unique_ptr<llama_model> model_ptr(llama_model_create(ml, params));
@@ -55,7 +55,11 @@ void llama_model_cohere2moe::load_arch_tensors(llama_model_loader & ml) {
    const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
    const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
    const int trunk_flags = mtp_only  ? TENSOR_NOT_REQUIRED : 0;
-    const int mtp_flags   = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags         = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+
+    if (!ml.load_mtp) {
+        mtp_flags |= TENSOR_SKIP;
+    }

    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);

@@ -91,7 +91,11 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader & ml) {
    const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
    const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
    const int trunk_flags = mtp_only   ? TENSOR_NOT_REQUIRED : 0;
-    const int mtp_flags   = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags         = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+
+    if (!ml.load_mtp) {
+        mtp_flags |= TENSOR_SKIP;
+    }

    const bool is_mla = hparams.is_mla();
    if (!is_mla) {
@@ -33,7 +33,11 @@ void llama_model_hy_v3::load_arch_tensors(llama_model_loader & ml) {
    const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
    const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
    const int trunk_flags = mtp_only   ? TENSOR_NOT_REQUIRED : 0;
-    const int mtp_flags   = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags         = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+
+    if (!ml.load_mtp) {
+        mtp_flags |= TENSOR_SKIP;
+    }

    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);

@@ -271,8 +271,6 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
            } else {
                const int64_t n_idx_dim = hparams.indexer_head_size;   // 128

-                GGML_ASSERT(!inp_attn->self_k_rot && !inp_attn->self_v_rot && "MSA: attn-rot not supported");
-
                // Index Branch, project, norm, partial RoPE, cache
                ggml_tensor * iq = build_lora_mm(model.layers[il].index_q_proj, cur);
                ggml_tensor * ik = build_lora_mm(model.layers[il].index_k_proj, cur);
@@ -289,6 +287,14 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
                ggml_build_forward_expand(gf, mctx_cur->cpy_k_idx(ctx0, ik, inp_attn->get_k_idxs(), il));
                ggml_tensor * ik_kv = mctx_cur->get_k_idx(ctx0, il);

+                if (inp_attn->self_k_rot) {
+                    Qcur = llama_mul_mat_hadamard(ctx0, Qcur, inp_attn->self_k_rot);
+                    Kcur = llama_mul_mat_hadamard(ctx0, Kcur, inp_attn->self_k_rot);
+                }
+                if (inp_attn->self_v_rot) {
+                    Vcur = llama_mul_mat_hadamard(ctx0, Vcur, inp_attn->self_v_rot);
+                }
+
                // Main branch: store K/V, take cache views
                ggml_build_forward_expand(gf, Qcur);
                ggml_build_forward_expand(gf, Kcur);
@@ -431,7 +437,9 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
                        cur = ggml_concat(ctx0, cur, outs[st], 1);
                    }
                }
-
+                if (inp_attn->self_v_rot) {
+                    cur = llama_mul_mat_hadamard(ctx0, cur, inp_attn->self_v_rot);
+                }
                cb(cur, "kqv_out", il);
                if (model.layers[il].wo) {
                    cur = build_lora_mm(model.layers[il].wo, cur, model.layers[il].wo_s);
@@ -39,6 +39,7 @@ void llama_model_qwen35::load_arch_tensors(llama_model_loader & ml) {

    const bool mtp_only = (hparams.n_layer_nextn > 0) && (ml.get_weight("blk.0.attn_norm.weight") == nullptr);
    const int trunk_flags = mtp_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags = !ml.load_mtp ? TENSOR_SKIP : 0;

    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);

@@ -97,25 +98,25 @@ void llama_model_qwen35::load_arch_tensors(llama_model_loader & ml) {
        auto & layer = layers[il];

        // MTP block looks like a full-attention Qwen3.5 decoder block.
-        layer.attn_norm      = create_tensor(tn(LLM_TENSOR_ATTN_NORM,      "weight", il), { n_embd }, 0);
-        layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, 0);
+        layer.attn_norm      = create_tensor(tn(LLM_TENSOR_ATTN_NORM,      "weight", il), { n_embd }, mtp_flags);
+        layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, mtp_flags);

-        create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, 0);
-        layer.wo          = create_tensor(tn(LLM_TENSOR_ATTN_OUT,    "weight", il), { n_embd_head_k * n_head, n_embd }, 0);
-        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", il), { n_embd_head_k }, 0);
-        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", il), { n_embd_head_k }, 0);
+        create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, mtp_flags);
+        layer.wo          = create_tensor(tn(LLM_TENSOR_ATTN_OUT,    "weight", il), { n_embd_head_k * n_head, n_embd }, mtp_flags);
+        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", il), { n_embd_head_k }, mtp_flags);
+        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", il), { n_embd_head_k }, mtp_flags);

-        layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", il), {n_embd,   n_ff}, 0);
-        layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", il), {  n_ff, n_embd}, 0);
-        layer.ffn_up   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", il), {n_embd,   n_ff}, 0);
+        layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", il), {n_embd,   n_ff}, mtp_flags);
+        layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", il), {  n_ff, n_embd}, mtp_flags);
+        layer.ffn_up   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", il), {n_embd,   n_ff}, mtp_flags);

        // NextN-specific tensors that define the MTP block.
-        layer.nextn.eh_proj          = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ,          "weight", il), { 2 * n_embd, n_embd }, 0);
-        layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", il), { n_embd },              0);
-        layer.nextn.hnorm            = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM,            "weight", il), { n_embd },              0);
-        layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS,     "weight", il), { n_embd, n_vocab },     TENSOR_NOT_REQUIRED);
-        layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", il), { n_embd, n_vocab },     TENSOR_NOT_REQUIRED);
-        layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", il), { n_embd },              TENSOR_NOT_REQUIRED);
+        layer.nextn.eh_proj          = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ,          "weight", il), { 2 * n_embd, n_embd }, mtp_flags);
+        layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", il), { n_embd },              mtp_flags);
+        layer.nextn.hnorm            = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM,            "weight", il), { n_embd },              mtp_flags);
+        layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS,     "weight", il), { n_embd, n_vocab },     mtp_flags|TENSOR_NOT_REQUIRED);
+        layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", il), { n_embd, n_vocab },     mtp_flags|TENSOR_NOT_REQUIRED);
+        layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", il), { n_embd },              mtp_flags|TENSOR_NOT_REQUIRED);
    };

    for (int i = 0; i < n_layer; ++i) {
@@ -42,6 +42,7 @@ void llama_model_qwen35moe::load_arch_tensors(llama_model_loader & ml) {

    const bool mtp_only = (hparams.n_layer_nextn > 0) && (ml.get_weight("blk.0.attn_norm.weight") == nullptr);
    const int trunk_flags = mtp_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags = !ml.load_mtp ? TENSOR_SKIP : 0;

    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);

@@ -113,32 +114,32 @@ void llama_model_qwen35moe::load_arch_tensors(llama_model_loader & ml) {
        const int64_t n_ff_shexp = hparams.n_ff_shexp ? hparams.n_ff_shexp : n_ff;

        // MTP block looks like a full-attention Qwen3.5 decoder block with MoE FFN.
-        layer.attn_norm      = create_tensor(tn(LLM_TENSOR_ATTN_NORM,      "weight", il), { n_embd }, 0);
-        layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, 0);
+        layer.attn_norm      = create_tensor(tn(LLM_TENSOR_ATTN_NORM,      "weight", il), { n_embd }, mtp_flags);
+        layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, mtp_flags);

-        create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, 0);
-        layer.wo          = create_tensor(tn(LLM_TENSOR_ATTN_OUT,    "weight", il), { n_embd_head_k * n_head, n_embd }, 0);
-        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", il), { n_embd_head_k }, 0);
-        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", il), { n_embd_head_k }, 0);
+        create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, mtp_flags);
+        layer.wo          = create_tensor(tn(LLM_TENSOR_ATTN_OUT,    "weight", il), { n_embd_head_k * n_head, n_embd }, mtp_flags);
+        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", il), { n_embd_head_k }, mtp_flags);
+        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", il), { n_embd_head_k }, mtp_flags);

        // Routed experts
-        layer.ffn_gate_inp  = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,  "weight", il), { n_embd, n_expert }, 0);
-        layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { n_ff_exp, n_embd, n_expert }, 0);
-        create_tensor_gate_up_exps(layer, il, n_embd, n_ff_exp, n_expert, 0);
+        layer.ffn_gate_inp  = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,  "weight", il), { n_embd, n_expert }, mtp_flags);
+        layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { n_ff_exp, n_embd, n_expert }, mtp_flags);
+        create_tensor_gate_up_exps(layer, il, n_embd, n_ff_exp, n_expert, mtp_flags);

        // Shared experts
-        layer.ffn_gate_inp_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP_SHEXP, "weight", il), { n_embd }, 0);
-        layer.ffn_gate_shexp     = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP,     "weight", il), { n_embd, n_ff_shexp }, 0);
-        layer.ffn_up_shexp       = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP,       "weight", il), { n_embd, n_ff_shexp }, 0);
-        layer.ffn_down_shexp     = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP,     "weight", il), { n_ff_shexp, n_embd }, 0);
+        layer.ffn_gate_inp_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP_SHEXP, "weight", il), { n_embd }, mtp_flags);
+        layer.ffn_gate_shexp     = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP,     "weight", il), { n_embd, n_ff_shexp }, mtp_flags);
+        layer.ffn_up_shexp       = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP,       "weight", il), { n_embd, n_ff_shexp }, mtp_flags);
+        layer.ffn_down_shexp     = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP,     "weight", il), { n_ff_shexp, n_embd }, mtp_flags);

        // NextN-specific tensors that define the MTP block.
-        layer.nextn.eh_proj          = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ,          "weight", il), { 2 * n_embd, n_embd }, 0);
-        layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", il), { n_embd },              0);
-        layer.nextn.hnorm            = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM,            "weight", il), { n_embd },              0);
-        layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS,     "weight", il), { n_embd, n_vocab },     TENSOR_NOT_REQUIRED);
-        layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", il), { n_embd, n_vocab },     TENSOR_NOT_REQUIRED);
-        layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", il), { n_embd },              TENSOR_NOT_REQUIRED);
+        layer.nextn.eh_proj          = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ,          "weight", il), { 2 * n_embd, n_embd }, mtp_flags);
+        layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", il), { n_embd },              mtp_flags);
+        layer.nextn.hnorm            = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM,            "weight", il), { n_embd },              mtp_flags);
+        layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS,     "weight", il), { n_embd, n_vocab },     mtp_flags|TENSOR_NOT_REQUIRED);
+        layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", il), { n_embd, n_vocab },     mtp_flags|TENSOR_NOT_REQUIRED);
+        layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", il), { n_embd },              mtp_flags|TENSOR_NOT_REQUIRED);
    };

    for (int i = 0; i < n_layer; ++i) {
@@ -48,7 +48,11 @@ void llama_model_step35::load_arch_tensors(llama_model_loader & ml) {
    const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
    const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
    const int trunk_flags = mtp_only  ? TENSOR_NOT_REQUIRED : 0;
-    const int mtp_flags   = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+    int mtp_flags         = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+
+    if (!ml.load_mtp) {
+        mtp_flags |= TENSOR_SKIP;
+    }

    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);

@@ -8192,9 +8192,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {

    for (ggml_type type_input : {GGML_TYPE_F32}) {
        for (ggml_op_pool pool_type : {GGML_OP_POOL_AVG, GGML_OP_POOL_MAX}) {
-            for (int k0 : {1, 3}) {
-                for (int s0 : {1, 2}) {
-                    for (int p0 : {0, 1}) {
+            for (int k0 : {1, 2, 3}) {
+                for (int s0 : {1, 2, 3}) {
+                    for (int p0 : {0, 1, 2, 3}) {
                        test_cases.emplace_back(new test_pool1d(pool_type, type_input, { 10,  3, 2, 1 }, k0, s0, p0));
                        test_cases.emplace_back(new test_pool1d(pool_type, type_input, { 11,  1, 3, 2 }, k0, s0, p0));
                        test_cases.emplace_back(new test_pool1d(pool_type, type_input, { 128, 2, 1, 3 }, k0, s0, p0));
@@ -9513,6 +9513,18 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
        }
    }

+    // prefill-shaped cases with long KV (nb >= 32, kv >= 1024): covers the
+    // XMX/GEMM-accelerated SYCL FA path which only activates for these shapes.
+    for (int kv : { 1024, 2048, }) {
+        for (int hs : { 64, 128, 256, }) {
+            for (int nb : { 32, 64, }) {
+                for (ggml_type type_KV : { GGML_TYPE_F16, GGML_TYPE_Q8_0, GGML_TYPE_Q4_0, }) {
+                    test_cases.emplace_back(new test_flash_attn_ext(hs, hs, 8, {4, 1}, kv, nb, true, false, 0, 0, GGML_PREC_F32, type_KV, type_KV));
+                }
+            }
+        }
+    }
+
    for (int hsk : { 40, 64, 72, 80, 96, 128, 192, 256, 320, 512, 576 }) {
        for (int hsv : { 40, 64, 72, 80, 96, 128, 192, 256, 512 }) {
            if (hsk != 192 && hsk != 320 && hsk != 576 && hsk != hsv) continue;
@@ -41,6 +41,7 @@
 #define KEY_PROJ_DIM            "clip.%s.projection_dim"
 #define KEY_N_HEAD              "clip.%s.attention.head_count"
 #define KEY_N_HEAD_KV           "clip.%s.attention.head_count_kv"
+#define KEY_N_EMBD_HEAD         "clip.%s.attention.head_dim"
 #define KEY_LAYER_NORM_EPS      "clip.%s.attention.layer_norm_epsilon"
 #define KEY_FEATURE_LAYERS      "clip.%s.feature_layer"

@@ -54,6 +54,8 @@ struct clip_hparams {
    int32_t projection_dim = 0;
    int32_t n_head = 0;
    int32_t n_head_kv = 0;
+    // 0 = derive from n_embd; set when qkv width != n_embd
+    int32_t n_embd_head = 0;
    int32_t n_layer = 0;
    int32_t n_merge = 1; // number of patch merges **per-side**

@@ -253,7 +253,7 @@ clip_graph::clip_graph(clip_ctx * ctx, const clip_image_f32 & img) :
        n_embd(hparams.n_embd),
        n_head(hparams.n_head),
        n_head_kv(hparams.n_head_kv),
-        d_head(n_head > 0 ? n_embd / n_head : 0),
+        d_head(hparams.n_embd_head > 0 ? hparams.n_embd_head : (n_head > 0 ? n_embd / n_head : 0)),
        n_layer(hparams.n_layer),
        n_mmproj_embd(clip_n_mmproj_embd(ctx)),
        eps(hparams.eps),
@@ -372,13 +372,13 @@ ggml_tensor * clip_graph::build_vit(
                /* nb1    */ ggml_row_size(cur->type, d_head),
                /* nb2    */ cur->nb[1],
                /* nb3    */ cur->nb[1] * n_pos,
-                /* offset */ ggml_row_size(cur->type, n_embd));
+                /* offset */ ggml_row_size(cur->type, n_head * d_head));

                Vcur = ggml_view_4d(ctx0, cur, d_head, n_head, n_pos, B,
                /* nb1    */ ggml_row_size(cur->type, d_head),
                /* nb2    */ cur->nb[1],
                /* nb3    */ cur->nb[1] * n_pos,
-                /* offset */ ggml_row_size(cur->type, 2 * n_embd));
+                /* offset */ ggml_row_size(cur->type, 2 * n_head * d_head));

                if (layer.q_norm) {
                    GGML_ASSERT(layer.q_norm->ne[0] == Qcur->ne[0]);
@@ -1190,6 +1190,7 @@ struct clip_model_loader {
            const char * prefix = is_vision ? "vision" : "audio";
            get_u32(string_format(KEY_N_EMBD,         prefix), hparams.n_embd);
            get_u32(string_format(KEY_N_HEAD,         prefix), hparams.n_head);
+            get_u32(string_format(KEY_N_EMBD_HEAD,    prefix), hparams.n_embd_head, false);
            get_u32(string_format(KEY_N_FF,           prefix), hparams.n_ff);
            get_u32(string_format(KEY_N_BLOCK,        prefix), hparams.n_layer);
            get_u32(string_format(KEY_PROJ_DIM,       prefix), hparams.projection_dim);
@@ -212,6 +212,7 @@ struct server_slot {
    llama_tokens spec_prompt;
    std::vector<int32_t> spec_i_batch;
    common_prompt_checkpoint spec_ckpt;
+    bool spec_is_replay = false;

    // TODO: move members that belong to the task (such as `generated_text`, `has_new_line`) to task_results_state
    //       see https://github.com/ggml-org/llama.cpp/pull/18283#issuecomment-3710175837
@@ -332,6 +333,8 @@ struct server_slot {
    void reset() {
        SLT_DBG(*this, "%s", "\n");

+        spec_is_replay = false;
+
        n_prompt_tokens_cache = 0;

        last_nl_pos    = 0;
@@ -3875,6 +3878,7 @@ private:
                        }

                        // partial acceptance is not supported by the context -> truncate the draft and restore the state
+                        slot.spec_is_replay = true;
                        slot.spec_draft = std::move(accepted);

                        const auto & ckpt = slot.spec_ckpt;
@@ -3909,16 +3913,22 @@ private:

            const auto ids = std::move(slot.spec_draft);

+            size_t n_accepted = ids.size() - 1;
+            if (slot.spec_is_replay && n_accepted > 0) {
+                n_accepted--;
+            }
+            slot.spec_is_replay = false;
+
            slot.t_token_generation = std::max<int64_t>(1, t_now - slot.t_start_generation) / 1e3;

            // update how many tokens out of those tested were accepted
-            slot.n_draft_accepted += ids.size() - 1;
+            slot.n_draft_accepted += n_accepted;
            slot.n_draft_verif_steps += 1;

            if (slot.n_accepted_per_pos.empty()) {
                slot.n_accepted_per_pos.resize(common_speculative_n_max(&params_base.speculative), 0);
            }
-            for (size_t i = 0; i < ids.size() - 1 && i < slot.n_accepted_per_pos.size(); ++i) {
+            for (size_t i = 0; i < n_accepted && i < slot.n_accepted_per_pos.size(); ++i) {
                slot.n_accepted_per_pos[i]++;
            }

@@ -3954,7 +3964,7 @@ private:

            slot.print_timings_tg();

-            SLT_DBG(slot, "accepted %d/%d draft tokens, new n_tokens = %d\n", (int) ids.size() - 1, (int) n_draft, slot.prompt.n_tokens());
+            SLT_DBG(slot, "accepted %d/%d draft tokens, new n_tokens = %d\n", (int) n_accepted, (int) n_draft, slot.prompt.n_tokens());
        });
    }
Author	SHA1	Message	Date
Anand Patil	876a432116	vulkan: add POOL_1D op (#25431 ) * vulkan : add pool1d push constants and pipeline field Declared data structures needed for POOL1D OP, which are the vk_op_pool1d_push_constants struct and pipeline_pool1d_f32 field. * vulkan : add pool1d compute shader Added pool1d.comp for Vulkan backend mirroring the existing pool2d shader. * vulkan : add full GGML_OP_POOL_1D support Added pipeline creation and op dispatch for 1D pooling in the Vulkan backend. * vulkan : fix pool1d shader logic Registered pool1d_f32 in vulkan-shaders-gen.cpp and fixed tensor dimension indices and avg pool scale. * vulkan : fix pool1d end boundary crash and expand test coverage Fixed an issue where the shader crashed when the end boundary was negative when k0 < p0. Also, added more test cases related to this fix.	2026-07-31 16:48:58 +02:00
Masato Nakasaka	eb41d503ba	vulkan: Introduce driver version check for Windows Intel GPU to mitigate crashing (#25192 ) * Removed crash guard for Intel Crash fixed from driver 32.0.101.8860 * Added driver version check for windows * Change to convert from driverVersion rather than string * No need to use signed * Refactor * allow GPU other than Xe2+ * adjusted function body position	2026-07-31 16:26:37 +02:00
Xuan-Son Nguyen	db7d8b24b5	mtmd: add n_embd_head (#26342 ) Co-authored-by: Daniel Han <unslothai@gmail.com>	2026-07-31 15:30:19 +02:00
timkhronos	a09d8abf8c	Support rotated kv cache quant (#26180 )	2026-07-31 21:06:40 +08:00
fairydreaming	82dbc4f017	llama : load MTP tensors only if they are really used (#26296 ) * llama : load MTP tensors only if they are really used * llama : skip loading MTP (if not used) in remaining models that support MTP --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>	2026-07-31 14:57:02 +02:00
Jeff Bolz	6f3c0a790b	vulkan: update vulkan sdk to 1.4.357.0 (#26303 )	2026-07-31 07:27:03 -05:00
Ruixiang Wang	000547513f	server: correct accepted tokens when need draft token replay (#26320 ) * spec: correct accepted tokens when need draft token replay * cont : naming --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-07-31 11:16:17 +03:00
David Friehs	15e755f30d	cuda: extract Q2_0 elements via __byte_perm (#25603 )	2026-07-31 11:15:44 +03:00
Ozymandias_EBON	9d9a6d29f6	SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc… (#25025 ) * SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing * fattn-mkl: fix interleaved dst layout in normalize kernel - 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 * Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md 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 * Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros - 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 <cfloat> - 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 * fattn-mkl: enable MKL FA for all KV cache types 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. * fattn-mkl: rename mkl_disable -> mkl_enable for clarity * 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-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards 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 == ne02) and interleaved (nb1 == H * ne02) 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 <noreply@anthropic.com> fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests Adding K>=1024 flash-attn test cases surfaced several MKL bugs: - Quant K/V with a padded seq-view (real KV cache) used the wrong strides in the dequant path... only the true Gemma interleave layout should reconstruct strides. nb[2] vs ne[1]nb[1] - Gate was firing on shapes the kernel doesn't handle: head_dim < 64 or not a multiple of 64, MHA, attention sinks, and bf16 decode... fell through to vec which no bf16 case. Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512 (has to be a multiple of 64) with matching K/V head size, mask, no sinks/alibi/softcap... everything else falls back to tile. Covers Qwen Dense/MoE and Gemma4 Dense/MoE Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass. Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using Qwen 27b q5_k_xl Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review. * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners --------- Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>	2026-07-31 10:43:16 +03:00