diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index 27a8bce8b..5a5cf69fd 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -282,8 +282,10 @@ static mkl_fa_kv_desc mkl_fa_make_desc(const ggml_tensor * T, bool interleaved, 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; + // MLA's V cache is a 512-wide view of 576-wide K rows. Treat any + // padded row stride as strided even when there is only one KV head. + d.mode = interleaved || d.nb1 != d.D * (int64_t)sizeof(sycl::half) + ? 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 { @@ -412,7 +414,9 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * 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 + // Alias the dequantized buffers only when K and V expose the same values. + // MLA V is a narrower view of K and needs its own strided dequantization. + const bool V_is_K_view = V->ne[0] == K->ne[0] && V->view_src && (V->view_src == K || (V->view_src == K->view_src && V->view_offs == K->view_offs)); diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index 541ae8a82..8bd946e09 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -146,15 +146,17 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing. // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ... // Note: MKL GEMM calls are incompatible with SYCL graph capture replay. - // 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. + const bool standard_shape = Q->ne[0] >= 64 && Q->ne[0] <= 512 && + Q->ne[0] % 64 == 0 && Q->ne[0] == V->ne[0]; + // GLM-4.7 Flash's MLA shape is already expressible by the MKL pipeline: + // KQ GEMM uses DKQ=576 while VKQ and output use DV=512. Keep this narrow + // until other mismatched K/V shapes have independent correctness data. + const bool glm_mla_shape = Q->ne[0] == 576 && K->ne[0] == 576 && + V->ne[0] == 512 && gqa_ratio == 20 && + K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16; + if (g_ggml_sycl_enable_mkl_fa == 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] && + (standard_shape || glm_mla_shape) && 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)) { @@ -164,7 +166,10 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const // 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) { + const bool glm_v_stride = glm_mla_shape && t == V && + V->view_src && V->nb[1] == K->nb[1]; + if (!glm_v_stride && t->type == GGML_TYPE_F16 && + t->nb[1] % (t->ne[0] * 2) != 0) { kv_strides_ok = false; break; }