Commit d2a79e604 for llama.cpp
commit d2a79e60464903f9832fbecdba48bd4cc6c47e75
Author: Anant Shrivastava <anant@anantshri.info>
Date: Wed Oct 7 12:54:07 2026 +0530
sycl: accelerate GLM MLA prefill with MKL flash attention (#29171)
* sycl: accelerate GLM MLA prefill with MKL flash attention
GLM-4.7 Flash uses an MLA shape with 576-wide Q/K heads, a 512-wide
V head, GQA 20, and F16 KV. The SYCL dispatcher rejects this shape
because the normal MKL flash-attention gate requires matching K/V
widths and caps the head dimension at 512, so prompt processing falls
back to the substantially slower TILE kernel.
Admit only the validated 576/576/512, GQA-20 F16 shape to the existing
MKL pipeline. Keep all other mismatched K/V shapes on their current
fallback paths.
Handle GLM's V cache as a narrower strided view of K rows. Select the
strided F16 descriptor when row stride is padded, and alias K/V
dequantization buffers only when their logical widths match. Restrict
the stride exception to a real V view sharing K's row stride.
Add the exact 576/512, GQA-20 prompt-path backend test.
On an Intel Arc Pro B70 at master e613ef2, pp8192 improves from
432.80 to 1292.29 tok/s (2.99x, +198.6%). tg256 remains unchanged
within noise at 45.67 versus 45.65 tok/s. The exact MLA test passes
and debug output confirms MKL dispatch.
* sycl: store MKL flash attention scores in F16
Keep the QK GEMM output in F16 instead of F32. The online softmax still
converts each score to F32 for its max, exponent, and sum, so the
per-element math is unchanged apart from score rounding, and the F32
matrix was being written only to be consumed as F16 probabilities.
The F32 score matrix is the largest flash-attention intermediate on this
path; storing it as F16 halves its size and traffic. This builds on the
coalesced softmax loads from 1aa2954bd, which read each score row
cooperatively, so the smaller dtype pays off.
Measured on an Intel Arc Pro B70 with the dispatch from the previous
commit, -ngl 999 -b 4096 -ub 1024 -ctk f16 -ctv f16 -fa on:
pp8192 1583.9 -> 1657.1 tok/s (+4.6%)
pp64000 610.0 -> 684.5 tok/s (+12.2%)
pp131072 ~354 -> 402.1 tok/s (+13.5%)
tg256 at 8k context is unchanged (32.64), and the FLASH_ATTN_EXT suite
shows no new failures. The exact GLM MLA backend cases pass against CPU.
Adjust the ~354 baseline figure if you prefer citing only measured pairs (the 131k dispatch-only point came from the equivalent maintained build). Optionally add Assisted-by: <tool name> per the contribution guidelines since AI contributed to the change.
* Revert "sycl: store MKL flash attention scores in F16"
This reverts commit 265f9748164dc88e02678425b75f520ccd528b32.
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;
}