[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#25222
Conversation
github-actions
Bot
added
ggml
|
Hi @hmscider, thanks for your contribution! Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:
Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below. |
arthw
left a comment
arthw
left a comment
There was a problem hiding this comment.
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% |
| #else // !GGML_SYCL_DNNL | ||
|
|
||
| void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { | ||
| ggml_sycl_flash_attn_ext_tile(ctx, dst); |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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
There was a problem hiding this comment.
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!
Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (ggml-org#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>
|
Very good Pull Request. I tried that with |
|
@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. |
|
The PR gave me a big boost in pre-fill but OpenCode starts spitting out gibberish with the It may be due to my dual CUDA+SYCL setup CUDA0,SYCL0 |
arthw
left a comment
arthw
left a comment
There was a problem hiding this comment.
It's good job!
It shows good performance increasing!
Thank you!
|
Thanks for the PR. Prefill boost is impressive. I can confirm the gibberish output too. On a dual B70 setup, with In OpenCode/Pi: Environment: |
|
The following UTs seems to fail on Alchemist (DG2) with this PR: |
|
@hmscider |
|
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. |
|
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 |
|
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. |
…corruption fix) Fixes the multi-turn 'GGGGG...' repetition collapse on dual B60 with PR ggml-org#25222 oneDNN flash-attn on. Root cause: cont_to_f16 -> oneDNN execute -> permute_sdpa_out_sycl is fully async on the SYCL stream, but the ggml_sycl_pool_alloc<Qf/Kf/Vf/ scbuf/outf> RAII destructors fire on the host at function return, returning all temp buffers to the host-side pool. The next scheduler-dispatched op (e.g. the very next layer's flash-attn) can re-acquire those same device bytes from the pool while the GPU is still doing the oneDNN SDPA / permute, so the in-flight op reads garbage -- the softmax NaNs/infs and the model collapses to emitting a single repeated token ('GGGGG...'). Fix: stream->wait_and_throw() at the end of the try block, before the RAII destructors fire. This is the same fix pattern used in the XMX path (commit da66dd3: 'Add stream->wait() before sycl::free() in XMX KV-split cleanup'), ported to the oneDNN graph path. Verified head-to-head (dual B60, Meta-Llama-3-8B-Instruct Q4_0, oneDNN FA, --split-mode layer --tensor-split 1,1 --cache-type-k/v f16): Turn 1 (q=26 < MIN_Q=32 -> TILE): correct (unchanged, baseline) Turn 2 (q=53 >= 32 -> oneDNN XMX, GGGGG... -> correct Q/K/V explanation 105.66 t/s prompt eval confirms XMX systolic kernel still active) A/B isolation: 'wait()' alone (no host->device mask copy) is sufficient; removing the mask copy does not regress Turn 2. So the minimal fix is this single stream flush -- no host-mask migration, no graph recompilation. Note: this serialises layers through this op and may slightly reduce prefill throughput, but it is correctness-only additive -- the only behavioural change is forcing the async chain to retire before freeing.
| tensor to(E.out, eng, outf.get()); | ||
| E.cp.execute(strm, ti, {to}); | ||
|
|
||
| permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream); |
There was a problem hiding this comment.
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.
| 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(); |
|
One-line reproducibility fix posted as a review suggestion at This PR's oneDNN SDPA path silently breaks multi-turn prefill on multi-GPU — the 2nd chat turn collapses to a single repeated token Reproduced on dual Arc B60 with
Root cause: the 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 — |
@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: |
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. |
|
|
I hope this helps. Arc 140T (arrow lake H, xe-lpg+), linux Build: ... 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: 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) run: -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
changed the title
|
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 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: |
Add four-tier XMX-accelerated flash attention dispatch: 1. ONEDNN: native F16 oneDNN SDPA (from PR ggml-org#25222 by @hmscider) 2. HYBRID: dequant K/V to F16, then oneDNN SDPA (new) 3. MKL: oneMKL GEMM fallback for SDPA-incompatible shapes 4. TILE/VEC: existing scalar fallback for decode Fixes: permute formula (h/t swap), pool-alloc stream sync, MKL gate guards (mask, sinks, head_dim >= 64). Co-Authored-By: Claude Code on DeepSeek-v4-Pro
Good that it works for Q4_K. It may be the gguf version that cause trouble (Q5_k_xl ud unsloth). |
|
Same result of Qwen3-Coder-30B-A3B-Instruct-UD-Q4_K_XL.gguf. 25.3.3
2026.0 ./build/bin/llama-bench --model ../models/Qwen3-Coder-30B-A3B-Instruct-UD-Q4_K_XL.gguf -fa 1 -n 0 --n-cpu-moe 99
|
|
@hmscider |
… API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k
… 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf
f6d9183 to
2f4e8e1
Compare
|
@arthw Did the rebase, built llama.cpp, and confirm it working. Can you please check whether the rebase look good to go? thank you. |
…vice sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious. Co-authored-by: maxious <81432+maxious@users.noreply.github.com>
2f4e8e1 to
77a7b5b
Compare
arthw
left a comment
arthw
left a comment
There was a problem hiding this comment.
It's good job!
The related UT cases are passed.
Thank you!
arthw
added
the
merge ready
|
I confirmed that ddfa616#diff-56e3ee0e3b82833d2b02cafc29f63b2df7e4c0a8cc7c6dd98c57e9a730e35f88R254 fixed the issue if I remove the condition on it. The current code only counts for >1 SYCL devices. My CUDA+SYCL setup was still working off the single GPU branch and producing gibberish. |
* [SYCL] F16 (default) Flash Attention with XMX engine via oneDNN graph API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k * [SYCL] Address review on FA oneDNN path. Result: llama-bench---pp512; 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf * PR-25222 revision v2: addressed audits * [SYCL] flash-attn oneDNN SDPA KV F16 rev 3.0: add BMG gate + multi-device sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious. Co-authored-by: maxious <81432+maxious@users.noreply.github.com> * updated comment on bmg gate, noted the issue --------- Co-authored-by: scientist3 <scientist.3@users.noreply.github.com> Co-authored-by: hmscider <hmscider@users.noreply.github.com> Co-authored-by: maxious <81432+maxious@users.noreply.github.com>
Extends the oneDNN SDPA path (PR ggml-org#25222) to handle non-F16 KV caches by dequantizing or converting K/V to dense FP16 on-device before feeding them into the SDPA graph. The fused systolic kernel then runs identically to the native FP16 path. Supported KV types: - Q4_0, Q4_1, Q5_0, Q5_1, Q8_0: to_fp16_sycl / to_fp16_nc_sycl - F32: cont_to_f16_sycl<float> - BF16 and IQ types are excluded (no conversion kernel available) Gate: non-F16 requires K >= 1024 and Q >= 32 (prefill only). F16 KV runs at any length (existing behavior). Also includes the stream sync fix (stream->wait_and_throw() unconditional, PR ggml-org#25741 by @malsbat) and removal of V_is_K_view aliasing (K and V are always dequantized to separate buffers). Co-Authored-By: Claude <noreply@anthropic.com>
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). |
|
* [SYCL] F16 (default) Flash Attention with XMX engine via oneDNN graph API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k * [SYCL] Address review on FA oneDNN path. Result: llama-bench---pp512; 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf * PR-25222 revision v2: addressed audits * [SYCL] flash-attn oneDNN SDPA KV F16 rev 3.0: add BMG gate + multi-device sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious. Co-authored-by: maxious <81432+maxious@users.noreply.github.com> * updated comment on bmg gate, noted the issue --------- Co-authored-by: scientist3 <scientist.3@users.noreply.github.com> Co-authored-by: hmscider <hmscider@users.noreply.github.com> Co-authored-by: maxious <81432+maxious@users.noreply.github.com>
|
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! |
13 participants
Overview
flash-attention-xmx-fp16ggml-org/llama.cppProblem: 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
ggml/src/ggml-sycl/fattn-onednn.cppggml/src/ggml-sycl/fattn-onednn.hpp---XMX FA implementation.ggml/src/ggml-sycl/fattn.cpp---3 changes: include the header; addBEST_FATTN_KERNEL_ONEDNN = 150; call the support gate and dispatch to it.Spec
b9856, oneDNN 3.11.2, Level-Zero V1,GGML_SYCL_FA_ONEDNN=1Qwen3.6-27B-Q8_0.gguf&&gemma-4-12b-it-Q8_0.ggufMinimal build command used for testing, for replication (just llama-bench and 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
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
How it works
5 ops in attention
Matmul->Divide->Add(mask)->SoftMax->Matmulwere fused byoneDNN 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)
Q->ne[3]==1)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