[SYCL] Flash Attention with XMX engine via oneDNN graph API (SDPA) on KV f16 for Xe2 ; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k

[hmscider]

Overview

• Branch: flash-attention-xmx-fp16
• Target: ggml-org/llama.cpp
• Hardware tested: Intel Arc Pro B70 (Xe2, 256 XMX engines), confirmed working under oneAPI toolkit 2025.3 && 2026.0.0

Problem: The current SYCL flash attention degrades the prefill worse than fa=0, because it uses shader unit (XVE engine). Prefill decays with context and doesn't benefit by flash attention's O(N) memory efficiency.

Solution: Implement flash attention that utilizes XMX engines, as the CUDA backend already does. Routes eligible prefill flash-attention through oneDNN Graph API fused Scaled Dot Product Attention (sdp_primitive_kernel_t, XMX kernel), instead of the current XVE kernels.

What changed

• New: ggml/src/ggml-sycl/fattn-onednn.cpp ggml/src/ggml-sycl/fattn-onednn.hpp---XMX FA implementation.
• Modified: ggml/src/ggml-sycl/fattn.cpp---3 changes: include the header; add BEST_FATTN_KERNEL_ONEDNN = 150; call the support gate and dispatch to it.

Spec

GPU	Arc Pro B70 (BMG-G31, Xe2), 32 Xe2 cores, 256 XMX engines, 4096 shading units, 32 GB GDDR6, 608 GB/s
Build	b9856, oneDNN 3.11.2, Level-Zero V1, GGML_SYCL_FA_ONEDNN=1
Model	Qwen3.6-27B-Q8_0.gguf && gemma-4-12b-it-Q8_0.gguf

Minimal build command used for testing, for replication (just llama-bench and llama-perplexity)

cmake -S . -B build -G Ninja \
    -DGGML_SYCL=ON -DGGML_SYCL_F16=ON -DGGML_SYCL_DNN=ON \
    -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx \
    -DGGML_SYCL_DEVICE_ARCH=xe2 -DCMAKE_BUILD_TYPE=Release \
    -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_EXAMPLES=OFF -DLLAMA_BUILD_SERVER=OFF

cmake --build build -j"$(nproc)" --target llama-bench llama-perplexity

1. Throughput Result llama-bench Qwen3.6 27b Q8_0 gguf

Command llama-bench -m Qwen3.6-27B-Q8_0.gguf -n 0 -r 2 -p 512, 1024, ...

Table

ctx	fa=0 default (t/s)	fa=1 default (t/s)	fa=1 XMX (t/s)	fa1 XMX/fa0 default	fa1 XMX/ fa1 default
pp512	813.96 ± 0.24	783.01 ± 3.09	948.53 ± 2.04	1.165	1.21×
pp1024	810.11 ± 1.49	773.44 ± 2.32	946.09 ± 0.63	1.168	1.22×
pp2048	800.95 ± 0.43	744.12 ± 1.20	940.92 ± 0.41	1.175	1.26×
pp4096	790.40 ± 1.07	690.28 ± 0.02	930.81 ± 0.65	1.178	1.35×
pp8192	767.15 ± 1.38	608.52 ± 0.19	914.38 ± 1.40	1.192	1.50×
pp12288	743.55 ± 0.08	—	901.01 ± 0.40	1.212	—
pp16384	718.14 ± 0.01	482.77 ± 0.00	886.94 ± 0.12	1.235	1.84×
pp24576	676.54 ± 0.06	389.64 ± 0.23	864.45 ± 0.04	1.278	2.22×
pp32768	638.14 ± 0.25	326.33 ± 0.11	843.38 ± 0.13	1.322	2.58×
pp36864	621.72 ± 0.17	—	826.25 ± 0.15	1.329	—
pp40960	605.87 ± 0.18	281.36 ± 0.24	818.80 ± 0.20	1.351	2.91×
pp45056	OOM	263.54 ± 0.32	810.26 ± 0.00	—	3.07×
pp49152	—	247.68 ± 0.10	799.24 ± 0.27	—	3.23×
pp65536	—	200.42 ± 0.04	762.00 ± 0.02	—	3.80×
pp72000	—	186.10 ± 0.28	745.94 ± 0.72	—	4.01×
pp80000	—	171.22 ± 0.20	730.20 ± 0.12	—	4.26×
pp82000	—	—	OOM	—	—

Figure

2. Correctness---llama-perplexity

Command: llama-perplexity -m gemma-4-12b-it-Q8_0.gguf -f wiki.test.raw -c 512 -ngl 999 -fa {0,1}
Dataset: wikitext-2-raw-v1 / wiki.test.raw

run	PPL	error	s/pass	wall-clock
fa=0 (default path)	698.19	± 10.44	4.14	9 min 31.6 sec
fa=1 (FA XMX fix path)	695.79	± 10.40	2.67	6 min 7.0 sec
delta	2.39	—	-35%	-35% (1.56x faster)

How it works

5 ops in attention Matmul->Divide->Add(mask)->SoftMax->Matmul were fused by oneDNN SDPA. oneDNN's graph pattern-matcher recognizes the shape of the operations and substitutes the flash-attention-style fused kernel for it, which is what Graph API is for.

Additional information

When it engages (else falls through unchanged)

• f16 K/V, mask present (f16, single head/batch broadcast), attention sinks & softcap fall back to default.
• head_dim ∈ {40,64,72,80,96,112,128,256,512} with V==K
• GQA divides evenly
• single sequence (Q->ne[3]==1)
• prefill (Q->ne[1] ≥ 32).

Safety — strictly additive

Every non-eligible shape falls through to the existing default FA path (no behavior change). Any oneDNN exception at runtime is caught and falls back to the default. Compile is guarded by GGML_SYCL_DNNL; oneDNN must be presented. Implemented runtime kill-switch: GGML_SYCL_FA_ONEDNN=0.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: Code was written with AI assistance (Claude), reviewed and tested by the author on intel GPU.

[ggml-gh-bot]

Hi @hmscider, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• AI-generated content: This project does not accept PRs, descriptions or commit messages that are fully or predominantly AI-generated. If you have used AI to assist you in writing code, please make sure to disclose that explicitly.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[arthw]

It's good job!

With Qwen3.5-35B-A3B-Q4_K_M.gguf on B60:

Test	fa	Base t/s	Primary t/s	Increase Rate (Primary vs Base)
pp512	0	159.31	190.44	19.54%
pp512	1	244.46	273.09	11.71%
tg128	0	11.44	12.32	7.69%
tg128	1	11.56	15.18	31.31%

[arthw]

If there is no GGML_SYCL_DNNL, ggml_sycl_flash_attn_ext_onednn() can't be called.
So, no need define ggml_sycl_flash_attn_ext_onednn() too.

[hmscider]

Done---removed defining it, and guarded the dispatch call site in fattn.cpp with #if GGML_SYCL_DNNL.
Since the defined switch references ggml_sycl_flash_attn_ext_onednn always, removing it breaks the DNNL=0 link (undefined reference). The new guard placed in fattn.cpp allows the linking to properly work. This behavior is confirmed with icpx (oneAPI 2026.0), running the following code (no llama.cpp rebuild required).

// linktest.cpp -- the fix: call wrapped like #if GGML_SYCL_DNNL, built with it OFF
extern void onednn_fn();            // declared, never defined (== stub removed)
int main(int argc, char **) {
    switch (argc) {
        case 999:
#if GGML_SYCL_DNNL                  // the guard we added in fattn.cpp; here it's off -> reference removed
            onednn_fn();
#endif
            break;
        default:
            break;
    }
    return 0;
}

Error case, when just removing the definition:

icpx -fsycl -O2 linktest.cpp -o linktest      # Intel oneAPI DPC++/C++ 2026.0.0

/usr/bin/ld: linktest-*.o: in function main': linktest.cpp:(.text+0x1c): undefined reference to onednn_fn()'
icpx: error: linker command failed with exit code 1

Success case, with the guard in place:

icpx -fsycl -O2 linktest.cpp -o linktest # Intel oneAPI DPC++/C++ 2026.0.0

(no linker error) # build exit 0
$ ./linktest ; echo $?
0

[hmscider]

Results---PR #25222 review-fix build (DNN=ON), Llama-3.1-8B-Instruct-Q8_0, B70

Here are some quick tests with 8B model:

Test 1---PP512 throughput

fa	t/s
0 (XVE-SIMD)	3009.38 ± 6.57
1 (XMX oneDNN SDPA)	3959.89 ± 0.78

Works as intended, 1.32× (+31.6%)

Test 2---perplexity (wikitext-2-raw, 12 chunks)

fa	PPL ± SE
0 (XVE-SIMD)	8.8780 ± 0.40061
1 (XMX oneDNN SDPA)	8.8769 ± 0.40051

delta=0.11%, accurate

Overall, it works well!

[WizardlyBump17]

Very good Pull Request. I tried that with ./llama-server --host 0.0.0.0 --port 8080 --model /models/Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf.ignore --jinja --threads 8 --ctx-size 262144 --cache-ram 0 --parallel 1 --temperature 0.0 --top-p 0.2 --top-k 20 --no-mmap --spec-type draft-mtp --spec-draft-n-max 3 --batch-size 2700 --ubatch-size 2700 --n-gpu-layers 99 --n-cpu-moe 99 --reasoning-preserve --flash-attn on and my 116k prompt went from 245t/s to 462t/s. I noticed that it only works for F16 KV. Do you plan to bring it to the other KV quants?

[hmscider]

@WizardlyBump17 Thanks for reporting, 245 tk/s to 462tk/s increase is phenomenal. Your right, it currently only works for F16 KV. And yes, after getting this PR landed, the next goal would be to implement other kv quants.

[AshotN]

The PR gave me a big boost in pre-fill but OpenCode starts spitting out gibberish with the GGML_SYCL_FA_ONEDNN=1 env set.

It may be due to my dual CUDA+SYCL setup

CUDA0,SYCL0
split-mode layer
tensor-split 1,3
Model: Qwen3.6-27B-MTP-GGUF Q8_0
KV: f16/f16
FA: on
draft-mtp, n=2

[arthw]

It's good job!
It shows good performance increasing!

Thank you!

[saurabh-deochake]

Thanks for the PR. Prefill boost is impressive.

I can confirm the gibberish output too. On a dual B70 setup, with GGML_SYCL_FA_ONEDNN=1 seeing garbled output in the thinking block on first response when using OpenCode and Pi coding agent. However, a direct curl to /v1/chat/completions seems to be working just fine.

In OpenCode/Pi:

The model produced output that does not match the expected peg-native format.

Environment:
2x Intel Arc Pro B70 (xe2), oneAPI DPC++/C++ Compiler 2026.0.0, Level-Zero 1.28.2, oneDNN 3.11.2.
Build b9863, Qwen3.6-27B-Q4_K_M with MTP, draft-mtp, n=2, fa=1, single GPU for this test (--device SYCL0), f16 KV cache.

[qnixsynapse]

The following UTs seems to fail on Alchemist (DG2) with this PR:

Failing tests:
  FLASH_ATTN_EXT(hsk=40,hsv=40,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=40,hsv=40,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=40,hsv=40,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=40,hsv=40,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=1024,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=1024,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=72,hsv=72,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=72,hsv=72,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=72,hsv=72,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=72,hsv=72,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=80,hsv=80,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=80,hsv=80,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=80,hsv=80,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=80,hsv=80,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=96,hsv=96,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=96,hsv=96,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=96,hsv=96,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=113,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=113,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=1024,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=1024,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[12,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=def,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=256,hsv=256,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=256,hsv=256,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=256,hsv=256,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=256,hsv=256,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[1,1],kv=1024,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[1,1],kv=1024,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=512,hsv=512,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  Backend SYCL0: FAIL
Backend 2/2: CPU
  Skipping
1/2 backends passed
FAIL

[arthw]

@hmscider
Could you check the UT failed issue above?

[SleepinDevil]

I am on Windows 11 with an Intel Arc Pro B70.

Have built both FP16 SYCL and Vulkan for my testing and the below are interesting results from the Intel dev's Vulkan PRs over at #24408 vs this (#25222) SYCL PR here.

This SYCL build is very interesting though, definitely a huge improvement over not having DNN active, very well done on getting the SYCL speeds up! So excited to see if you can sort out the bugs mentioned in the thread above and after that seeing you move this onto quantised KV because that's where the real benefits are imo!

Nice work, keen to see this get sorted and merged :)

The Vulkan build with #24406, & #24407 is my daily driver at the moment with Qwen3.6-27b-Q6_K (MTP) and for my use case it seems to still be the faster build compared to this SYCL build on my environment at least.

|----- SYCL FP16 /w PR#25222 - dnnl & GGML_SYCL_FA_ONEDNN=1 -----|

call "C:\Program Files (x86)\Intel\oneAPI\setvars.bat" intel64
set "ONEAPI_DEVICE_SELECTOR=level_zero:gpu"
set "GGML_SYCL_FA_ONEDNN=1"
llama-bench  -p 8192  -n 128  -d 0,8192,65536  -r 3  -fa 1  --delay 10  -ngl 99  -m ..\LLM-models\Qwen3.6-27B-MTP-Q6_K.gguf
| model                          |       size |     params | backend    | ngl |  fa |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --: | --------------: | -------------------: |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |          pp8192 |        933.74 ± 0.27 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |           tg128 |         20.55 ± 0.02 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |  pp8192 @ d8192 |        732.70 ± 0.29 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |   tg128 @ d8192 |         19.19 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 | pp8192 @ d65536 |        570.49 ± 0.07 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |  tg128 @ d65536 |         12.70 ± 0.00 |

build: 8241e1a83 (9863)


|----- SYCL FP16 /w PR#25222 - GGML_SYCL_FA_ONEDNN=0 -----|

call "C:\Program Files (x86)\Intel\oneAPI\setvars.bat" intel64
set "ONEAPI_DEVICE_SELECTOR=level_zero:gpu"
set "GGML_SYCL_FA_ONEDNN=0"
llama-bench  -p 8192  -n 128  -d 0,8192,65536  -r 3  -fa 1  --delay 10  -ngl 99  -m ..\LLM-models\Qwen3.6-27B-MTP-Q6_K.gguf
| model                          |       size |     params | backend    | ngl |  fa |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --: | --------------: | -------------------: |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |          pp8192 |        653.19 ± 1.20 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |           tg128 |         20.55 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |  pp8192 @ d8192 |        384.98 ± 0.13 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |   tg128 @ d8192 |         18.99 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 | pp8192 @ d65536 |        104.46 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | SYCL       |  99 |   1 |  tg128 @ d65536 |         12.50 ± 0.02 |

build: 8241e1a83 (9863)


|----- Vulkan /w PR#24406, & PR#24407 -----|

llama-bench  -p 8192  -n 128  -d 0,8192,65536  -r 3  -fa 1  --delay 10  -ngl 99  --device Vulkan1  -m ..\LLM-models\Qwen3.6-27B-MTP-Q6_K.gguf
ggml_vulkan: Found 2 Vulkan devices:
ggml_vulkan: 0 = NVIDIA GeForce RTX 3080 Ti (NVIDIA) | uma: 0 | fp16: 1 | bf16: 1 | warp size: 32 | shared memory: 49152 | int dot: 1 | matrix cores: NV_coopmat2
ggml_vulkan: 1 = Intel(R) Arc(TM) Pro B70 Graphics (Intel Corporation) | uma: 0 | fp16: 1 | bf16: 0 | warp size: 32 | shared memory: 49152 | int dot: 1 | matrix cores: KHR_coopmat
| model                          |       size |     params | backend    | ngl |  fa | dev          |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --: | ------------ | --------------: | -------------------: |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      |          pp8192 |       1013.48 ± 1.13 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      |           tg128 |         21.19 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      |  pp8192 @ d8192 |        910.56 ± 8.21 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      |   tg128 @ d8192 |         20.27 ± 0.01 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      | pp8192 @ d65536 |       554.22 ± 18.99 |
| qwen35 27B Q6_K                |  21.30 GiB |    27.32 B | Vulkan     |  99 |   1 | Vulkan1      |  tg128 @ d65536 |         15.73 ± 0.00 |

build: 9da5214fd (9595)

[hmscider]

@hmscider
Could you check the UT failed issue above?

Hi, @arthw, it appears that all checks are passed on my end. Could you please point me to the exact failure issue?

Edit: got it. Since my PR was based on bmg, perhaps the issue may be specific to alchemist. Let me look into it.

[hmscider]

@saurabh-deochake

Thanks for the report. Its interesting that querying server directly works but going through harness cause issues. I don't have dual GPU so can't test its settings. Could you run llama-perplexity test with dual GPU and share the results here? Trying to see if untested dual GPU output setting is causing accuracy issue.

[hmscider]

@qnixsynapse @arthw

The error comes from onednn graph api spda. Can you check oneDNN & icpx version? It should be icpx 2026.0.0 & oneDNN 3.11.2 or older. If version matches then it's likely GPU architecture issue.

edit: confirmed the build works for icpx 2025.3 (oneAPI toolkit 2025.3). @arthw confirmed that this is a bug specific to alchemist and opened issue on the intel side.

[maxious]

The oneDNN SDPA path silently breaks multi-turn prefill on multi-GPU — the second turn collapses to a repeated token (GGGGG…).

cont_to_f16 → E.cp.execute → permute_sdpa_out_sycl are async on this stream, but the ggml_sycl_pool_allocs below return their device buffers to the host pool now. Before the GPU retires, the next layer's flash-attn can re-acquire the same bytes → in-flight SDPA reads garbage → softmax NaN/inf → token collapse. Invisible on single-turn because q < GGML_SYCL_FA_ONEDNN_MIN_Q falls back to TILE.

Reproduced on dual Arc B60, --flash-attn on --cache-type-k f16 --cache-type-v f16 --split-mode layer --tensor-split 1,1: Turn 2 (q >= 32 → oneDNN XMX, ~105 tok/s prompt eval) → GGGGG…; Turn 1 (q < 32 → TILE) correct. With this flush applied, Turn 2 produces coherent output and keeps the XMX throughput.

Suggested change

	permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream);
	permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream);
	// DO NOT REMOVE: cont_to_f16 -> oneDNN execute -> permute is async on this stream, but the
	// pool_alloc*s above free their device buffers at host return. Without this wait the next
	// scheduler op re-acquires those bytes while the GPU is still computing the SDPA, turning
	// it into garbage and collapsing multi-turn output to a single repeated token ("GGGGG...").
	stream->wait_and_throw();

[maxious]

One-line reproducibility fix posted as a review suggestion at ggml/src/ggml-sycl/fattn-onednn.cpp:239.

This PR's oneDNN SDPA path silently breaks multi-turn prefill on multi-GPU — the 2nd chat turn collapses to a single repeated token GGGGG….

Reproduced on dual Arc B60 with --flash-attn on --cache-type-k f16 --cache-type-v f16 --split-mode layer --tensor-split 1,1:

• Turn 1 (q < 32 → TILE): correct.
• Turn 2 (q ≥ 32 → oneDNN XMX, ~105 tok/s prompt-eval): collapses to GGGGG….

Root cause: the ggml_sycl_pool_alloc temps Qf/Kf/Vf/scbuf/outf return to the host-side pool at function-return, before the GPU has retired the async cont_to_f16 → E.cp.execute → permute_sdpa_out_sycl chain. The next layer's flash-attn then re-acquires those same device bytes while the in-flight SDPA is still writing → NaN/inf scores → token collapse.

Single-turn perplexity/benchmarks miss this entirely because the prefill prompt is short enough to fall back to the TILE kernel.

Fix is one line — stream->wait_and_throw() — please apply from the inline suggestion at #discussion_r3523162505. End-to-end verified on the same setup: multi-turn output renders coherent; XMX prompt-eval throughput (~105 tok/s) preserved.

[qnixsynapse]

@qnixsynapse @arthw

The error comes from onednn graph api spda. Can you check oneDNN & icpx version? It should be icpx 2026.0.0 & oneDNN 3.11.2 or older. If version matches then it's likely GPU architecture issue.

@hmscider Please disable fa-onednn in older version if new version is a requirement. I am guessing garbled output might be the result of this.

Update: I upgraded oneapi base toolkit to 2026.1.0. Still UTs are failing:

Failing tests:
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=1024,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[1,1],kv=1024,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=32,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,1,2,3])
  FLASH_ATTN_EXT(hsk=64,hsv=64,nh=4,nr23=[4,1],kv=512,nb=75,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=f16,type_V=f16,permute=[0,2,1,3])
  Backend SYCL0: FAIL

[hmscider]

@qnixsynapse @arthw

The error comes from onednn graph api spda. Can you check oneDNN & icpx version? It should be icpx 2026.0.0 & oneDNN 3.11.2 or older. If version matches then it's likely GPU architecture issue.

@hmscider Please disable fa-onednn in older version if new version is a requirement. I am guessing garbled output might be the result of this.

We can do that, but it'd be a temporary fix. I believe we need to access the root cause—whether this is from onednn & icpx version or if this is alchemist-specific issue. Please provide onednn & icpx version first, and upgrade onednn & icpx if yours is older version. It may work fine on alchemist.

[qnixsynapse]

@qnixsynapse @arthw
The error comes from onednn graph api spda. Can you check oneDNN & icpx version? It should be icpx 2026.0.0 & oneDNN 3.11.2 or older. If version matches then it's likely GPU architecture issue.

@hmscider Please disable fa-onednn in older version if new version is a requirement. I am guessing garbled output might be the result of this.

We can do that, but it'd be a temporary fix. I believe we need to access the root cause—whether this is from onednn & icpx version or if this is alchemist-specific issue. Please provide onednn & icpx version first, and upgrade onednn & icpx if yours is older version. It may work fine on alchemist.

• oneDNN: 2026.0
• icpx: Intel oneAPI DPC++/C++ Compiler 2026.1.0 (build 2026.1.0.20260617)

[jlionhan]

I hope this helps.

Arc 140T (arrow lake H, xe-lpg+), linux

Build:
cmake --fresh -B build -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DCMAKE_BUILD_TYPE=Release -DGGML_SYCL=1 -DBUILD_SHARED_LIBS=0 -DGGML_SYCL_F16=1

...
-- The CXX compiler identification is IntelLLVM 2026.1.0
...
-- Using oneAPI Release SYCL compiler (icpx).
-- SYCL found
-- SYCL Compiler version: 20260100
-- SYCL_INCLUDE_DIR: /opt/intel/oneapi/compiler/2026.1/include
-- SYCL_LIBRARY=/opt/intel/oneapi/compiler/2026.1/lib/libsycl.so
-- Found IntelSYCL: /opt/intel/oneapi/compiler/2026.1/include (found version "202012")
-- GGML_SYCL_SUPPORT_LEVEL_ZERO_API ON
-- Level Zero loader found: /lib/libze_loader.so
-- Level Zero headers found: /usr/include
-- Found oneDNN: /opt/intel/oneapi/dnnl/2026.0/lib/libdnnl.so.3.11
...

cmake --build build --config Release -j 12

llama-cli --temperature 1.0 --top-k 64 --top-p 0.95 --gpu-layers 999 --model ~/models/google_gemma-4-26B-A4B-it-Q5_K_M.gguf --ctx-size 32000 --threads 6 --parallel 1 --reasoning off --offline -fit off --cache-ram 0 --ctx-checkpoints 2

run:
> Where is the capital of the UK?

The capital of the United Kingdom is London. It is located in the southeast of England.

[ Prompt: 21.5 t/s | Generation: 13.3 t/s ]

> Who is Edison?

-- own own 싶 than--c--눔---e-ع-ngen----父亲-ingarg own ownер-n---ات-loglevel-epagination own ownな---s-on-べき-this-en로-----s-yق-服用--s--s----ات- own-ط-to- own--ಗ than-t-ed로ะ thaning-e-y---**--o---ات ownprinting-4-en-y-- neutral0ة-

(gibberish response to the second prompt)
> /exit

run:
> What are the most important environmental, climatic, economic, and geopolitical roles of the Arctic Ocean, and how do its unique ecosystems, sea ice, and marine biodiversity influence global climate stability, international shipping routes, and scientific research efforts?

-less冠-------м-- way Photographs---ر----ه-员𒋯𒑍--------मset--- own-y-phalط-----//-3-t-fade-tي-set----精度-enなclassifiers----よう-姓---s--с------ own-e-ر----な斬-r---際s--en`/-اً----遇到的--い---ا--口な-----اً----and---- own own----ing와-ers-------ad- own--कतों-ق-ed-z-y-து----ဍ-ه- Cot-서--imide 技術set own-6-jos-ca-ing---ه-y-es-ر-u- wa

(gibberish response to the first prompt)

Thank you

[hmscider]

@arthw @maxious

Placed a Battlemage-specific gate && applied @maxious suggestion for multi-gpu settings. Please refer to my comment as well.

It appears that for alchemist && arc-140t (xe-hpg && xe-lpg ) the oneDNN graph API isn't currently working. I realized that since I only have Xe2 (battlemage), there is no practical way for myself to dig into the oneDNN codebase and come up with working solution. Narrowing the scope of the PR to Xe2 (Battlemage)-specific is the best I can do at this rate. I hope other contributors with other architectures @qnixsynapse @jlionhan can step up in the future, after this PR, to work on architecture-specific builds, eventually.

@qnixsynapse @jlionhan For anyone interested in digging into oneDNN graph API on their hardware, I wrote a probe that you could run without extra build. This probe rebuilds the exact build_sdpa graph from my fattn-onednn.cpp, runs it on the failing shapes, and checks NMSE < 5e-4 vs a CPU reference. Also prints dnnl_version() + device. Add ONEDNN_VERBOSE=2 to see which kernel dispatched.

Please note, that upgrading oneDNN to the latest version (3.13.0 or newer) may fix the current issues on other architectures, as I checked their source code, and confirmed that xe_hpg (alchemist) support is there.

On my B70, built under oneDNN 3.11.2, running below code granted NMSE = 1e-7, so there shouldn't be bug in the code. Please use the below code for probing NMSE for other architectures.

Code:

// onednn_sdpa_alchemist_probe.cpp
// ============================================================
// PURPOSE: Show that the PR #25222 oneDNN Graph SDPA path is numerically
// CORRECT on Battlemage (BMG/Xe2, e.g. Arc Pro B70) but WRONG on
// Alchemist (DG2/XeHPG, e.g. Arc A-series) -- with NO llama.cpp rebuild.
//
// It rebuilds the EXACT SDPA graph that fattn-onednn.cpp::build_sdpa()
// compiles (MatMul->Divide->Add->SoftMax->MatMul, f16 in/out, f32 score),
// runs it on the exact shapes that FAIL test-backend-ops FLASH_ATTN_EXT on
// Alchemist, and compares to a CPU reference using the SAME metric and
// threshold the unit test uses: NMSE = sum((got-ref)^2)/sum(ref^2) < 5e-4
// (ggml test_flash_attn_ext::max_nmse_err).
//
// EXPECTED:
//   B70 (BMG):        all shapes PASS   (systolic sdp_primitive_kernel_t)
//   Arc A / DG2:      shapes FAIL       (different kernel / f16 precision)
//
// To also see WHICH oneDNN kernel is picked (systolic vs larger_partition):
//   ONEDNN_VERBOSE=2 ./alchemist_probe
//
// COMPILE (in the SYCL container, oneDNN 3.11.2):
//   icpx -fsycl -O2 onednn_sdpa_alchemist_probe.cpp \
//        -I${DNNLROOT}/include -L${DNNLROOT}/lib -ldnnl \
//        -o alchemist_probe && ./alchemist_probe
// ============================================================

#include <sycl/sycl.hpp>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstring>
#include <random>
#include <string>
#include <vector>
#include "oneapi/dnnl/dnnl.hpp"
#include "oneapi/dnnl/dnnl_graph.hpp"
#include "oneapi/dnnl/dnnl_sycl.hpp"

using namespace dnnl;
using namespace dnnl::graph;
using half = sycl::half;

// ---- CPU reference: MHA + GQA scaled-dot-product attention, causal mask ----
// Layout is head-major contiguous, matching the 5D logical tensors:
//   Q [H, q, d]   (head h uses kv-head h/rep),  K,V [Hkv, seq, d],  out [H, q, d]
static std::vector<float> cpu_attention(
        const std::vector<float>& Q, const std::vector<float>& K, const std::vector<float>& V,
        int H, int Hkv, int q, int seq, int d, float scale) {
    const int rep = H / Hkv;
    const int causal_offset = seq - q;            // query i sees keys 0..(causal_offset+i)
    std::vector<float> out((size_t) H * q * d, 0.0f);
    std::vector<float> scores(seq);
    for (int h = 0; h < H; ++h) {
        const int hk = h / rep;                   // kv head for this query head
        const float* Qh = &Q[(size_t) h  * q   * d];
        const float* Kh = &K[(size_t) hk * seq * d];
        const float* Vh = &V[(size_t) hk * seq * d];
        float* Oh = &out[(size_t) h * q * d];
        for (int i = 0; i < q; ++i) {
            float mx = -1e30f;
            for (int j = 0; j < seq; ++j) {
                float s = 0.0f;
                for (int k = 0; k < d; ++k) s += Qh[i*d+k] * Kh[j*d+k];
                s *= scale;
                if (j > causal_offset + i) s = -INFINITY;   // causal mask
                scores[j] = s;
                if (s > mx) mx = s;
            }
            float sum = 0.0f;
            for (int j = 0; j < seq; ++j) {
                scores[j] = (scores[j] == -INFINITY) ? 0.0f : std::exp(scores[j] - mx);
                sum += scores[j];
            }
            if (sum == 0.0f) sum = 1.0f;
            for (int k = 0; k < d; ++k) {
                float acc = 0.0f;
                for (int j = 0; j < seq; ++j) acc += (scores[j] / sum) * Vh[j*d+k];
                Oh[i*d+k] = acc;
            }
        }
    }
    return out;
}

struct sdpa_partition {
    compiled_partition cp;
    std::vector<logical_tensor> ins;
    logical_tensor out;
    size_t id_q=0, id_k=0, id_v=0, id_scale=0, id_mask=0;
    bool ok=false;
};

// Identical graph to fattn-onednn.cpp::build_sdpa().
static sdpa_partition build_sdpa(const engine& eng, int mb,int H,int Hkv,int q,int seq,int d) {
    using ltype = logical_tensor::layout_type;
    using dt    = logical_tensor::data_type;
    using ldims = logical_tensor::dims;
    const int rep = H / Hkv;
    int64_t id = 0;
    ldims q_sz={mb,Hkv,rep,q,d}, kv_sz={mb,Hkv,1,seq,d}, s_sz={mb,Hkv,rep,q,seq};
    sdpa_partition E;

    auto query  = logical_tensor(id++, dt::f16, q_sz,  ltype::strided);
    auto key    = logical_tensor(id++, dt::f16, kv_sz, ltype::strided);
    auto score  = logical_tensor(id++, dt::f32, s_sz,  ltype::strided);
    op   bmm1(id++, op::kind::MatMul, "bmm1");
    bmm1.set_attr<bool>(op::attr::transpose_b, true);
    bmm1.add_inputs({query, key}); bmm1.add_outputs({score});

    auto scale  = logical_tensor(id++, dt::f16, {1,1,1,1,1}, ltype::strided);
    auto scaled = logical_tensor(id++, dt::f32, s_sz,  ltype::strided);
    op   sdiv(id++, op::kind::Divide, "scale_div");
    sdiv.add_inputs({score, scale}); sdiv.add_outputs({scaled});

    auto mask   = logical_tensor(id++, dt::f16, {mb,1,1,q,seq}, ltype::strided);
    auto masked = logical_tensor(id++, dt::f32, s_sz,  ltype::strided);
    op   madd(id++, op::kind::Add, "mask_add");
    madd.add_inputs({scaled, mask}); madd.add_outputs({masked});

    auto probs  = logical_tensor(id++, dt::f16, s_sz,  ltype::strided);
    op   smax(id++, op::kind::SoftMax, "softmax");
    smax.set_attr<int64_t>(op::attr::axis, -1);
    smax.set_attr<std::string>(op::attr::mode, "inf_as_zero");
    smax.add_inputs({masked}); smax.add_outputs({probs});

    auto value  = logical_tensor(id++, dt::f16, kv_sz, ltype::strided);
    auto output = logical_tensor(id++, dt::f16, q_sz,  ltype::strided);   // f16 out (systolic path)
    op   bmm2(id++, op::kind::MatMul, "bmm2");
    bmm2.add_inputs({probs, value}); bmm2.add_outputs({output});

    dnnl::graph::graph g(eng.get_kind());
    g.add_op(bmm1); g.add_op(sdiv); g.add_op(madd); g.add_op(smax); g.add_op(bmm2);
    g.finalize();
    auto parts = g.get_partitions();
    if (parts.size() != 1 || !parts[0].is_supported()) return E;

    E.ins = parts[0].get_input_ports();
    E.out = parts[0].get_output_ports()[0];
    E.cp  = parts[0].compile(E.ins, {E.out}, eng);
    E.out = E.cp.query_logical_tensor(E.out.get_id());
    E.id_q=query.get_id(); E.id_k=key.get_id(); E.id_v=value.get_id();
    E.id_scale=scale.get_id(); E.id_mask=mask.get_id();
    E.ok = true;
    return E;
}

// Run one shape, return true on PASS (NMSE < 5e-4).
static bool run_shape(sycl::queue& sq, const engine& eng, dnnl::stream& strm,
                      int H, int Hkv, int q, int seq, int d) {
    const int mb = 1;
    const float scale = 1.0f / std::sqrt((float) d);

    std::mt19937 rng(1234 + d*131 + H*17 + q);
    std::uniform_real_distribution<float> dist(-1.0f, 1.0f);
    std::vector<float> hQ((size_t)H*q*d), hK((size_t)Hkv*seq*d), hV((size_t)Hkv*seq*d);
    for (auto& x : hQ) x = dist(rng);
    for (auto& x : hK) x = dist(rng);
    for (auto& x : hV) x = dist(rng);

    auto ref = cpu_attention(hQ, hK, hV, H, Hkv, q, seq, d, scale);

    half* dQ    = sycl::malloc_device<half>((size_t)H*q*d, sq);
    half* dK    = sycl::malloc_device<half>((size_t)Hkv*seq*d, sq);
    half* dV    = sycl::malloc_device<half>((size_t)Hkv*seq*d, sq);
    half* dScale= sycl::malloc_device<half>(1, sq);
    half* dMask = sycl::malloc_device<half>((size_t)q*seq, sq);
    half* dOut  = sycl::malloc_device<half>((size_t)H*q*d, sq);

    std::vector<half> tQ(hQ.begin(),hQ.end()), tK(hK.begin(),hK.end()), tV(hV.begin(),hV.end());
    sq.memcpy(dQ, tQ.data(), tQ.size()*sizeof(half));
    sq.memcpy(dK, tK.data(), tK.size()*sizeof(half));
    sq.memcpy(dV, tV.data(), tV.size()*sizeof(half));
    half hscale = (half)(1.0f / scale);
    sq.memcpy(dScale, &hscale, sizeof(half));

    std::vector<half> hMask((size_t)q*seq);
    const int causal_offset = seq - q;
    for (int i = 0; i < q; ++i)
        for (int j = 0; j < seq; ++j)
            hMask[(size_t)i*seq+j] = (j > causal_offset + i) ? (half)(-65504.0f) : (half)0.0f;
    sq.memcpy(dMask, hMask.data(), hMask.size()*sizeof(half));
    sq.wait();

    auto E = build_sdpa(eng, mb, H, Hkv, q, seq, d);
    if (!E.ok) {
        printf("  d=%-3d rep=%-2d q=%-3d seq=%-4d  |  partition NOT supported (falls to TILE)\n",
               d, H/Hkv, q, seq);
        sycl::free(dQ,sq); sycl::free(dK,sq); sycl::free(dV,sq);
        sycl::free(dScale,sq); sycl::free(dMask,sq); sycl::free(dOut,sq);
        return true;   // not the failure mode we're hunting; TILE would handle it
    }

    auto id2ptr = [&](size_t r) -> void* {
        if (r == E.id_q)     return dQ;
        if (r == E.id_k)     return dK;
        if (r == E.id_v)     return dV;
        if (r == E.id_scale) return dScale;
        if (r == E.id_mask)  return dMask;
        return nullptr;
    };
    std::vector<tensor> ti; ti.reserve(E.ins.size());
    for (auto& lt : E.ins) ti.emplace_back(lt, eng, id2ptr(lt.get_id()));
    tensor to(E.out, eng, dOut);
    E.cp.execute(strm, ti, {to});
    strm.wait();

    std::vector<half> hOut((size_t)H*q*d);
    sq.memcpy(hOut.data(), dOut, hOut.size()*sizeof(half)).wait();

    double se = 0.0, sref = 0.0;      // NMSE = sum(err^2)/sum(ref^2)  (ggml UT metric)
    float max_err = 0.0f, max_ref = 0.0f;
    for (size_t i = 0; i < ref.size(); ++i) {
        float got = (float) hOut[i];
        float e = got - ref[i];
        se += (double) e * e;
        sref += (double) ref[i] * ref[i];
        max_err = std::max(max_err, std::fabs(e));
        max_ref = std::max(max_ref, std::fabs(ref[i]));
    }
    double nmse = sref > 0 ? se / sref : se;
    float  relmax = max_err / (max_ref > 0 ? max_ref : 1.0f);
    bool pass = nmse < 5e-4;
    printf("  d=%-3d rep=%-2d q=%-3d seq=%-4d  |  NMSE=%.2e  relmax=%.2e  |  %s\n",
           d, H/Hkv, q, seq, nmse, relmax, pass ? "PASS" : "FAIL");

    sycl::free(dQ,sq); sycl::free(dK,sq); sycl::free(dV,sq);
    sycl::free(dScale,sq); sycl::free(dMask,sq); sycl::free(dOut,sq);
    return pass;
}

int main() {
    sycl::queue sq{sycl::gpu_selector_v};
    auto eng  = dnnl::sycl_interop::make_engine(sq.get_device(), sq.get_context());
    auto strm = dnnl::sycl_interop::make_stream(eng, sq);

    const dnnl_version_t* v = dnnl_version();
    printf("=== oneDNN SDPA correctness on: %s ===\n",
           sq.get_device().get_info<sycl::info::device::name>().c_str());
    printf("oneDNN version: %d.%d.%d\n", v->major, v->minor, v->patch);
    printf("Threshold: NMSE < 5e-4 (same as test-backend-ops FLASH_ATTN_EXT)\n");
    printf("Shapes below are the exact ones that FAIL on Alchemist (nh=4, kv f16, causal).\n\n");

    // Exact failing shapes from PR #25222 UT report (Hkv = nh = 4).
    // {d, rep, q, seq}
    struct Shape { int d, rep, q, seq; };
    const std::vector<Shape> shapes = {
        { 40, 1, 32, 512},
        { 40, 4, 32, 512},
        { 64, 1, 32, 512},
        { 64, 1, 32, 1024},
        { 72, 1, 32, 512},
        { 80, 1, 32, 512},
        { 96, 1, 32, 512},
        {128, 1, 32, 512},
        {128, 4, 32, 512},
        {128,12, 32, 512},
        {256, 1, 32, 512},
        {512, 1, 32, 512},
    };

    const int Hkv = 4;
    int passed = 0, total = 0;
    for (const auto& s : shapes) {
        ++total;
        passed += run_shape(sq, eng, strm, Hkv * s.rep, Hkv, s.q, s.seq, s.d) ? 1 : 0;
    }
    printf("\n%d/%d shapes PASS.\n", passed, total);
    printf("%s\n", passed == total
           ? ">> SDPA path is correct on this device (expected on BMG/Xe2)."
           : ">> SDPA path is WRONG on this device (expected on Alchemist/DG2) -- do NOT enable here.");
    return passed == total ? 0 : 1;
}

Usage:

source /opt/intel/oneapi/setvars.sh
icpx -fsycl -O2 onednn_sdpa_alchemist_probe.cpp \
    -I${DNNLROOT}/include -L${DNNLROOT}/lib -ldnnl \
    -o alchemist_probe && ./alchemist_probe  

[F3zz1k]

@WizardlyBump17 Thanks for reporting, 245 tk/s to 462tk/s increase is phenomenal. Your right, it currently only works for F16 KV. And yes, after getting this PR landed, the next goal would be to implement other kv quants.

Are we tracking this anywhere? I've made an initial attempt at supporting q8_0 KV by converting to FP16 on the fly though this was a no-go with 30% throughput at low context and worsening down to 15% at 40k context on qwen3.6-27b. It was numerically correct and even saw a tiny improvement to perplexity however the overhead killed any gains. The scalar tile was faster at every depth on Intel Arc Pro B70 with D=256.

Things I haven't tried yet are native s8×s8 to s32 DPAS (skip the dequant) or feeding q8 as s8+group-scales into the oneDNN int8 SDPA ukernel.

Figured this might help save others from wasted effort. Can share more info if needed.

Edit: oneDNN is also a no-go for useful context sizes (unless you happen to use ~8k). Results compared to baseline: 100% @ 512, 115% @ 8k but 44% @ 40k. Quantizing down Q and softmax matrix could improve throughput but I assume the model quality hit isn't worth it (unless someone specifically trained a model for this setup).

[johnkarlhill]

@WizardlyBump17 Thanks for reporting, 245 tk/s to 462tk/s increase is phenomenal. Your right, it currently only works for F16 KV. And yes, after getting this PR landed, the next goal would be to implement other kv quants.

Are we tracking this anywhere? I've made an initial attempt at supporting q8_0 KV by converting to FP16 on the fly though this was a no-go with 30% throughput at low context and worsening down to 15% at 40k context on qwen3.6-27b. It was numerically correct and even saw a tiny improvement to perplexity however the overhead killed any gains. The scalar tile was faster at every depth on Intel Arc Pro B70 with D=256.

Things I haven't tried yet are native s8×s8 to s32 DPAS (skip the dequant) or feeding q8 as s8+group-scales into the oneDNN int8 SDPA ukernel.

Figured this might help save others from wasted effort. Can share more info if needed.

Edit: oneDNN is also a no-go for useful context sizes (unless you happen to use ~8k). Results compared to baseline: 100% @ 512, 115% @ 8k but 44% @ 40k. Quantizing down Q and softmax matrix could improve throughput but I assume the model quality hit isn't worth it (unless someone specifically trained a model for this setup).

#25874

[sheigl]

I just wanted to say thank you for this! Prompt processing is soo much faster in high contexts! I have a B70 and B580 serving up qwen 3.6 27b and it’s working great!