sycl: Update gate logic for Alchemist GPUs regarding OneDNN features. by HumerousGorgon · Pull Request #26635 · ggml-org/llama.cpp

HumerousGorgon · 2026-08-05T13:25:27Z

Overview

There have been a few recent OneDNN specific features that have aimed to increase the performance of prompt processing on Intel GPUs. Unfortunately, in pull #25222 it was found that SPDA routes threw incorrect results on Alchemist GPUs. As a result, Alchemist GPUs were gated from the SPDA paths.
The PR also left Alchemist users a testing script to verify whether the SPDA paths were correct or not by comparing it to a known good CPU result. When running this, I found that it was only the SPDA shapes with a head size of d=64 that failed.

Instead of gating all SPDA routes, I rewrote the gate to block Alchemist users from SPDA shapes with a head size of d=64, instead having them fall back to the regular SYCL FA routes. This allowed all other routes to run as OneDNN SPDAs.

Additional information

Prompt processing for models with head sizes !=64 is greatly increased. Qwen3.6-27B sees speeds of around 590-680t/s PP.
Prompt processing for models with head sizes = 64 is unchanged, but very performant given that most models with smaller head sizes are small.

The OneDNN issue with regards to head sizes of 64 is tracked in their repo and is reportedly in "triage" according to Intel. When the script shows that the failed shapes now pass, this gate can be completely removed.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: Claude helped me to understand what the testing script did, I wrote the new gate.

arthw

It's good idea!

Could you provide the test script or cmd for the abnormal output with head=64?
I will test this solution on other GPUs.

Thank you!

Accepted recommendations to add bmg_g31 arch. Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

HumerousGorgon · 2026-08-06T03:03:59Z

It's good idea!

Could you provide the test script or cmd for the abnormal output with head=64? I will test this solution on other GPUs.

Thank you!

The program to run it was from #25222:

// 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;
}

Running with OneDNN debugging enabled to show complete shapes, it gives this output:

onednn_verbose,v1,graph,compile:cache_miss,gpu,100002,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x40:5120s1280s1280s40s1 in1_f16:1:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 out0_f16:13:strided:undef:1x4x1x32x40:5120s1280s1280s40s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,12.2781
onednn_verbose,v1,graph,exec,gpu,100002,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x40:5120s1280s1280s40s1 in1_f16:1:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 out0_f16:13:strided:undef:1x4x1x32x40:5120s1280s1280s40s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,2.07202
  d=40  rep=1  q=32  seq=512   |  NMSE=1.18e-07  relmax=5.08e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100008,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x4x32x40:20480s5120s1280s40s1 in1_f16:1:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 out0_f16:13:strided:undef:1x4x4x32x40:20480s5120s1280s40s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,10.0129
onednn_verbose,v1,graph,exec,gpu,100008,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x4x32x40:20480s5120s1280s40s1 in1_f16:1:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x40:81920s20480s20480s40s1 out0_f16:13:strided:undef:1x4x4x32x40:20480s5120s1280s40s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.900879
  d=40  rep=4  q=32  seq=512   |  NMSE=1.59e-07  relmax=5.81e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100014,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x64:8192s2048s2048s64s1 in1_f16:1:strided:undef:1x4x1x512x64:131072s32768s32768s64s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x64:131072s32768s32768s64s1 out0_f16:13:strided:undef:1x4x1x32x64:8192s2048s2048s64s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,9.83691
onednn_verbose,v1,graph,exec,gpu,100014,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x64:8192s2048s2048s64s1 in1_f16:1:strided:undef:1x4x1x512x64:131072s32768s32768s64s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x64:131072s32768s32768s64s1 out0_f16:13:strided:undef:1x4x1x32x64:8192s2048s2048s64s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.885986
  d=64  rep=1  q=32  seq=512   |  NMSE=1.71e-01  relmax=5.67e-01  |  FAIL
onednn_verbose,v1,graph,compile:cache_miss,gpu,100020,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x64:8192s2048s2048s64s1 in1_f16:1:strided:undef:1x4x1x1024x64:262144s65536s65536s64s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x1024:32768s32768s32768s1024s1 in4_f16:12:strided:undef:1x4x1x1024x64:262144s65536s65536s64s1 out0_f16:13:strided:undef:1x4x1x32x64:8192s2048s2048s64s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.182129
onednn_verbose,v1,graph,exec,gpu,100020,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x64:8192s2048s2048s64s1 in1_f16:1:strided:undef:1x4x1x1024x64:262144s65536s65536s64s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x1024:32768s32768s32768s1024s1 in4_f16:12:strided:undef:1x4x1x1024x64:262144s65536s65536s64s1 out0_f16:13:strided:undef:1x4x1x32x64:8192s2048s2048s64s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.62915
  d=64  rep=1  q=32  seq=1024  |  NMSE=1.83e-01  relmax=5.04e-01  |  FAIL
onednn_verbose,v1,graph,compile:cache_miss,gpu,100026,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x72:9216s2304s2304s72s1 in1_f16:1:strided:undef:1x4x1x512x72:147456s36864s36864s72s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x72:147456s36864s36864s72s1 out0_f16:13:strided:undef:1x4x1x32x72:9216s2304s2304s72s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.37695
onednn_verbose,v1,graph,exec,gpu,100026,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x72:9216s2304s2304s72s1 in1_f16:1:strided:undef:1x4x1x512x72:147456s36864s36864s72s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x72:147456s36864s36864s72s1 out0_f16:13:strided:undef:1x4x1x32x72:9216s2304s2304s72s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.870117
  d=72  rep=1  q=32  seq=512   |  NMSE=1.23e-07  relmax=5.06e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100032,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x80:10240s2560s2560s80s1 in1_f16:1:strided:undef:1x4x1x512x80:163840s40960s40960s80s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x80:163840s40960s40960s80s1 out0_f16:13:strided:undef:1x4x1x32x80:10240s2560s2560s80s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.3291
onednn_verbose,v1,graph,exec,gpu,100032,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x80:10240s2560s2560s80s1 in1_f16:1:strided:undef:1x4x1x512x80:163840s40960s40960s80s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x80:163840s40960s40960s80s1 out0_f16:13:strided:undef:1x4x1x32x80:10240s2560s2560s80s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.802002
  d=80  rep=1  q=32  seq=512   |  NMSE=1.29e-07  relmax=4.23e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100038,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x96:12288s3072s3072s96s1 in1_f16:1:strided:undef:1x4x1x512x96:196608s49152s49152s96s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x96:196608s49152s49152s96s1 out0_f16:13:strided:undef:1x4x1x32x96:12288s3072s3072s96s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.16992
onednn_verbose,v1,graph,exec,gpu,100038,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x96:12288s3072s3072s96s1 in1_f16:1:strided:undef:1x4x1x512x96:196608s49152s49152s96s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x96:196608s49152s49152s96s1 out0_f16:13:strided:undef:1x4x1x32x96:12288s3072s3072s96s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.78418
  d=96  rep=1  q=32  seq=512   |  NMSE=1.12e-07  relmax=3.86e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100044,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x128:16384s4096s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x1x32x128:16384s4096s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.26709
onednn_verbose,v1,graph,exec,gpu,100044,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x128:16384s4096s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x1x32x128:16384s4096s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,1.1521
  d=128 rep=1  q=32  seq=512   |  NMSE=1.27e-07  relmax=4.22e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100050,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x4x32x128:65536s16384s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x4x32x128:65536s16384s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.271
onednn_verbose,v1,graph,exec,gpu,100050,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x4x32x128:65536s16384s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x4x32x128:65536s16384s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.546875
  d=128 rep=4  q=32  seq=512   |  NMSE=1.23e-07  relmax=4.07e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100056,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x12x32x128:196608s49152s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x12x32x128:196608s49152s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,8.60498
onednn_verbose,v1,graph,exec,gpu,100056,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x12x32x128:196608s49152s4096s128s1 in1_f16:1:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x128:262144s65536s65536s128s1 out0_f16:13:strided:undef:1x4x12x32x128:196608s49152s4096s128s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.521973
  d=128 rep=12 q=32  seq=512   |  NMSE=1.26e-07  relmax=4.82e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100062,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x256:32768s8192s8192s256s1 in1_f16:1:strided:undef:1x4x1x512x256:524288s131072s131072s256s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x256:524288s131072s131072s256s1 out0_f16:13:strided:undef:1x4x1x32x256:32768s8192s8192s256s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,9.83716
onednn_verbose,v1,graph,exec,gpu,100062,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x256:32768s8192s8192s256s1 in1_f16:1:strided:undef:1x4x1x512x256:524288s131072s131072s256s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x256:524288s131072s131072s256s1 out0_f16:13:strided:undef:1x4x1x32x256:32768s8192s8192s256s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,0.836914
  d=256 rep=1  q=32  seq=512   |  NMSE=1.26e-07  relmax=3.87e-04  |  PASS
onednn_verbose,v1,graph,compile:cache_miss,gpu,100068,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x512:65536s16384s16384s512s1 in1_f16:1:strided:undef:1x4x1x512x512:1048576s262144s262144s512s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x512:1048576s262144s262144s512s1 out0_f16:13:strided:undef:1x4x1x32x512:65536s16384s16384s512s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,11.3828
onednn_verbose,v1,graph,exec,gpu,100068,sdp,bmm1;scale_div;mask_add;softmax;bmm2,,in0_f16:0:strided:undef:1x4x1x32x512:65536s16384s16384s512s1 in1_f16:1:strided:undef:1x4x1x512x512:1048576s262144s262144s512s1 in2_f16:4:strided:undef:1x1x1x1x1:1s1s1s1s1 in3_f16:7:strided:undef:1x1x1x32x512:16384s16384s16384s512s1 in4_f16:12:strided:undef:1x4x1x512x512:1048576s262144s262144s512s1 out0_f16:13:strided:undef:1x4x1x32x512:65536s16384s16384s512s1,fpm:strict,sdp_primitive_kernel_t,dnnl_backend,1.42188
  d=512 rep=1  q=32  seq=512   |  NMSE=1.28e-07  relmax=4.90e-04  |  PASS

10/12 shapes PASS.
>> SDPA path is WRONG on this device (expected on Alchemist/DG2) -- do NOT enable here.

You can see that it's only the 2/12 shapes with head sizes of 64 that fail. This could likely be rectified on Intel's end as last I saw it's a driver issue.

arthw

It's good job!

Thank you!

arthw · 2026-08-08T01:52:39Z

@HumerousGorgon
After the new update, there is compiling error.
Please fix them!

Thank you!

HumerousGorgon · 2026-08-09T04:45:52Z

@HumerousGorgon After the new update, there is compiling error. Please fix them!

Thank you!

Forgot to add back in the arch var after I reworked the gate logic.. whoops!
Fixed now :)

arthw · 2026-08-12T05:23:54Z

@HumerousGorgon
Could you fix the code format issue in: EditorConfig Checker / editorconfig (pull_request)?

Thank you!

HumerousGorgon · 2026-08-12T16:49:32Z

@HumerousGorgon Could you fix the code format issue in: EditorConfig Checker / editorconfig (pull_request)?

Thank you!

Fixed!

arthw · 2026-08-13T06:58:35Z

I re-triger CI to check it.

Thank you!

HumerousGorgon · 2026-08-19T04:14:11Z

Is there anything else that needs to be done here prior to merge?

HumerousGorgon added 2 commits August 5, 2026 12:46

feat: updated gating logic of fattn-onednn.cpp

161626d

verified device types

17421e2

HumerousGorgon requested a review from a team as a code owner August 5, 2026 13:25

github-actions Bot added ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels Aug 5, 2026

HumerousGorgon changed the title ~~SYCL: Update gate logic for Alchemist GPUs regarding OneDNN features.~~ sycl: Update gate logic for Alchemist GPUs regarding OneDNN features. Aug 5, 2026

arthw reviewed Aug 6, 2026

View reviewed changes

Comment thread ggml/src/ggml-sycl/fattn-onednn.cpp Outdated

Update ggml/src/ggml-sycl/fattn-onednn.cpp

b291b92

Accepted recommendations to add bmg_g31 arch. Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

arthw reviewed Aug 6, 2026

View reviewed changes

Comment thread ggml/src/ggml-sycl/fattn-onednn.cpp Outdated

Improved SPDA gate, added documentation.

d388f31

arthw approved these changes Aug 6, 2026

View reviewed changes

Added arch var to reworked gate, fixing build errors.

044fb7b

Fix trailing whitespaces.

8e55842

arthw added the merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. label Aug 14, 2026

ggerganov merged commit 1cb3f5e into ggml-org:master Aug 21, 2026
29 of 31 checks passed

Navigation Menu

Uh oh!

sycl: Update gate logic for Alchemist GPUs regarding OneDNN features. - #26635

sycl: Update gate logic for Alchemist GPUs regarding OneDNN features.#26635
ggerganov merged 6 commits into
ggml-org:masterfrom
HumerousGorgon:onednn/alchemist-head-gate

HumerousGorgon commented Aug 5, 2026

Uh oh!

arthw left a comment

Uh oh!

Uh oh!

HumerousGorgon commented Aug 6, 2026 •

edited

Loading

Uh oh!

Uh oh!

arthw left a comment

Uh oh!

arthw commented Aug 8, 2026

Uh oh!

HumerousGorgon commented Aug 9, 2026

Uh oh!

arthw commented Aug 12, 2026

Uh oh!

HumerousGorgon commented Aug 12, 2026

Uh oh!

arthw commented Aug 13, 2026

Uh oh!

HumerousGorgon commented Aug 19, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

HumerousGorgon commented Aug 5, 2026

Overview

Additional information

Requirements

Uh oh!

arthw left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

HumerousGorgon commented Aug 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

arthw left a comment

Choose a reason for hiding this comment

Uh oh!

arthw commented Aug 8, 2026

Uh oh!

HumerousGorgon commented Aug 9, 2026

Uh oh!

arthw commented Aug 12, 2026

Uh oh!

HumerousGorgon commented Aug 12, 2026

Uh oh!

arthw commented Aug 13, 2026

Uh oh!

HumerousGorgon commented Aug 19, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

HumerousGorgon commented Aug 6, 2026 •

edited

Loading