* Add DMMV Q4_K and Q6_K ESIMD kernels
Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Refactor ESIMD kernels to share common code
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Move control of ESIMD from compile to runtime
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Use ESIMD by default when available
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Fix possible error when using ESIMD by default
While not an issue in the current version, this will become an
issue when additional QK ESIMD kernels are added (such as Q2_K).
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add explicit unroll to ESIMD kernels
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Tidy up ESIMD kernels a bit
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add DMMV Q5_K ESIMD kernel
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Remove redundant copyright notice
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
---------
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* add params
* cpu kernel
* metal kernel
* add test backend ops
* gate other backends
* ggml: (cuda) support ggml_rope_set_offset (#27121)
* rm cuda supports_op guard, fix webgpu clang-format
* ggml: support ggml_rope_set_offset on vulkan (#27344)
* ggml: support ggml_rope_set_offset on vulkan
* remove inplace optimization
Kernel is a port of `ggml-cuda/fwht.cu`
(us/run, median):
```
m x n x k GEMM FWHT speedup
64 x 1 x 64 10.20 2.93 3.48x
64 x 2048 x 64 10.75 2.71 3.97x
128 x 1 x 128 10.33 2.88 3.59x
128 x 32 x 128 9.20 2.77 3.33x
128 x 2048 x 128 16.46 2.76 5.95x
256 x 1 x 256 10.19 2.77 3.68x
256 x 2048 x 256 16.69 3.41 4.89x
512 x 2048 x 512 54.16 12.89 4.20x
```
Adjusts the thread/block count to be proportional to the size
of the quant, reducing under/over subscription.
Largest perf improvement is the q4_0 -> f32 path, with, on
a Arc 70, throughput goes from 20.21 GB/s to 158.19 GB/s
The rest of the quants are flat in performance uplift.
* Initial changes for Recurrent state rollback for nemotron for cpu and cuda
* Removing CPU RS rollback. Will enable it in subsequent PRs
* addition of test case
* Removing assert and calling runtime API to check if op is supported
* removing extra API and updating the call sites for K
* replace static cuda detection to runtime fused_op api
* address review comments and fallback when SSM rollback not supprted
* Adding changes for supporting RS-rollback in CPU. Also added test-backend-ops for cpu and cuda
* removing memory manipulation as rs rollback is now supported in CPU
* removing the static probe which is not needed now
* correcting the format
* address review comments
* enabling test for all the backends, unsupported backends will fallback to CPU
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* choose different graph based on the result of fused_ssm_op is supported or not and also handled memory->n_rs_seq >1 case incase of op is not supported
* Support K > 1 in ssm_scan for all backends
* Fix CI Issues
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
* Add DMMV Q4_K and Q6_K ESIMD kernels
Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Refactor ESIMD kernels to share common code
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Move control of ESIMD from compile to runtime
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Use ESIMD by default when available
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Fix possible error when using ESIMD by default
While not an issue in the current version, this will become an
issue when additional QK ESIMD kernels are added (such as Q2_K).
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add explicit unroll to ESIMD kernels
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Tidy up ESIMD kernels a bit
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add DMMV Q3_K ESIMD kernel
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
---------
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* llama: add new default load-mode auto which picks mmap unless a non-Metal iGPU is used
* Update ggml/src/ggml-hexagon/ggml-hexagon.cpp
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* set mmap_support to false on OpenCL backend
* fix order of load modes
* use -1 for auto
* resolve load mode auto earlier to correctly pick gpu host or cpu memory
* add load mode auto to llama-bench
* bump virtgpu api version, regenerate docs
---------
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
The NVFP4 quantization format stores a scaling factor for every group of
16 weights, packed into a single UE4M3 byte.
The SYCL GPU code was converting these scale values using the E4M3 path,
but that's *signed*, and these are unsigned values.
* tests: add SWIGLU perf cases
perf mode had no GLU coverage. Adds SWIGLU at 17408 columns, 512 and
2048 tokens, f16 and f32, with the operands both fused and split.
* sycl: consolidate fused-GLU kernels
They differed only in which op_* they called, so take the op as an argument and share a common launcher.
Their block sizes were all 256, so launch geometry is unchanged;
SYCL_GELU_BLOCK_SIZE and SYCL_SILU_BLOCK_SIZE lose their last users so are dropped.
* sycl: contiguous fast path for the fused GLU ops
o0 == n and o1 == n collapse the de-interleave index math to the
identity, so dispatch a flat kernel in that case. It fires for
ggml_glu_split with packed operands; a fused [gate|up] tensor keeps the
strided path. test-backend-ops perf -o SWIGLU on an Arc Pro B70: split
+14% f16 and +4% f32, fused unchanged.
* sycl: parallelize the non-contiguous concat kernel
Launch geometry only: the non-contiguous concat kernel launched a single-lane
work-group (1, 1, 1), now it will launch a (1, 1, SYCL_CONCAT_BLOCK_SIZE) one.
SYCL_CONCAT_BLOCK_SIZE is defined in `ggml/src/ggml-sycl/presets.hpp`.
llama-bench (Arc Pro B70, Qwen3.6-27B-UD-Q4_K_XL, -fa on, q8_0 KV),
on top of upstream master: pp2048 920 -> 1006 t/s (+9.4%)
* sycl: cap non-contiguous concat block at ne0
* sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_CONCAT_BLOCK_SIZE)
* Revert "sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_CONCAT_BLOCK_SIZE)"
This reverts commit 2709909e790d298cc0746950caff16ad8ecf9d6f.
* sycl: extend oneDNN SDPA to Q4_0-Q8_0 and F32 KV caches
Extends the oneDNN SDPA path (PR #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 #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>
* docs: drop GGML_SYCL_FA_DEBUG from SYCL.md (not shipped in this PR)
Co-Authored-By: Claude <noreply@anthropic.com>
---------
Co-authored-by: Claude <noreply@anthropic.com>
* SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing
* fattn-mkl: fix interleaved dst layout in normalize kernel
- Fix mkl_fa_normalize_head: use interleaved dst layout
((query * n_q_heads + head) * DV) matching TILE's
flash_attn_combine_results. Previously used dense head-major
layout which wrote head outputs to wrong addresses, corrupting
attention for all models except Qwen3.6-27B (where GQA=6 heads
were sparse enough to avoid visible overlap).
- Remove 7 redundant stream->wait() calls — SYCL in-order queue
already serializes pure SYCL kernel dependencies. Retain only
the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its
own internal queue that does not respect SYCL in-order).
- Remove unused dst_row_stride, diagnostic clutter, and dead
K/V hex dump (fa_diag block in fattn-mkl.cpp).
- Add MKL_FA_DISABLE=1 env var for A/B testing.
- Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output
fingerprint (MKL_FA_DIAG=1) in fattn.cpp.
Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B
Perf (B70/Battlemage, 32K, q8_0 KV):
Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x)
Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x)
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md
Completed the following:
- Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable)
- Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG
- Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG
- Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro
- Document all three env vars in docs/backend/SYCL.md under Runtime
- Add comment explaining MKL FA activation trigger (flash-attn + quantized
KV cache + batch-size >= 1024 + n_kv >= 1024)
Resolves review feedback from arthw.
Again, thank you!!!
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros
- Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks
(fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG,
GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG)
- Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM
KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion
- Gate MKL_ACCUM macro behind do_print so timing accumulators are
no-ops in normal operation
- Remove redundant MIT/Intel copyright header from fattn-mkl.cpp
- Remove unused #include <cfloat>
- Expand SYCL.md MKL FA docs with step-by-step activation trigger
and example llama-cli command
Again, thank you!!!
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* fattn-mkl: enable MKL FA for all KV cache types
Remove the quantized-only restriction on MKL activation — the MKL
kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM,
so F16 (default), BF16, and F32 caches all benefit from XMX hardware
acceleration. The type restriction was an unnecessary gate.
Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path)
After: ~670 t/s (MKL path, matching quantized-cache baseline)
Minimal change: two conditions removed, one comment updated in fattn.cpp.
No kernel or conversion code changes — the dequant pipeline already
covers all types.
* fattn-mkl: rename mkl_disable -> mkl_enable for clarity
* fattn-mkl: refine MKL FA dispatch gates
Three changes:
1. Remove quantized-only restriction - MKL FA activates for all
KV cache types (F16 default, BF16, F32, quantized). The MKL
kernel converts non-F16 K/V via to_fp16_sycl before GEMM.
2. Rename mkl_disable -> mkl_enable to match env var
(GGML_SYCL_ENABLE_MKL_FA).
3. Replace batch-size threshold with Q->ne[1] >= 32 gate.
Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where
fused kernel beats MKL launch overhead. Routes all
multi-token prefill through XMX-accelerated GEMM.
Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse,
128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s.
* ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards
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 (#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>
* fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests
Adding K>=1024 flash-attn test cases surfaced several MKL bugs:
- Quant K/V with a padded seq-view (real KV cache) used the wrong
strides in the dequant path... only the true Gemma interleave
layout should reconstruct strides. nb[2] vs ne[1]*nb[1]
- Gate was firing on shapes the kernel doesn't handle: head_dim < 64
or not a multiple of 64, MHA, attention sinks, and
bf16 decode... fell through to vec which no bf16 case.
Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512
(has to be a multiple of 64) with matching K/V head size, mask,
no sinks/alibi/softcap... everything else falls back to tile.
Covers Qwen Dense/MoE and Gemma4 Dense/MoE
Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass.
Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using
Qwen 27b q5_k_xl
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review.
* Update ggml/src/ggml-sycl/fattn-mkl.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn-mkl.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners
---------
Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path
The scale was uploaded with an async memcpy sourced from a stack local. On the
in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is
large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host
stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale.
Output then collapses to a single repeated token and the KV cache is poisoned
for the rest of the session.
Short contexts win the race by accident, and test-backend-ops caps
FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous
device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the
symptom only by keeping the frame alive across the copy at the cost of a host
sync on every FA call.
Fix: cache one device scalar per (device, value) -- the scale is constant per
model -- and upload it synchronously once. The single-device fast path (no
per-call host sync) is then safe: every device-side hazard already serializes
on the in-order queue. The multi-GPU conservative wait is kept unchanged.
Also:
- GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that
routes very long sequences to the native FA kernel.
- test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B
geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind
spot. Note the race itself needs a live multi-op pipeline to reproduce;
single-op runs pass even on broken builds.
Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output
byte-identical at temp 0 to the native FA path through 32k-deep prefill, with
prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth).
Assisted-by: Claude Fable 5
* sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it
Review feedback on #25880:
- read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via
ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it)
- document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md
runtime table
Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536
shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels
(long-sequence fp16 accumulation drift, present before this PR) and would fail
CI for an unrelated reason.
Assisted-by: Claude Fable 5
* sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled
Assisted-by: Claude Fable 5
* sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly
Assisted-by: Claude Fable 5
* Update ggml/src/ggml-sycl/fattn-onednn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* sycl: write the SDPA scale from a kernel instead of caching it
The per-(device, value) scale cache was a function-local static
unordered_map with no synchronization, so concurrent backend instances
could access and rehash it at the same time.
Write the scalar with a single_task instead. The value is captured into
the command, so no host memory has to outlive the call -- which is what
the use-after-return fix needed in the first place. That removes the
shared container, the leaked device allocation and the string key, and
it also closes the remaining async-memcpy-from-a-stack-local on the
first flash-attention call.
Ordering does not rely on timing: the queue is created with
sycl::property::queue::in_order and the dnnl stream wraps that same
queue, so the write completes before the SDPA reads the scalar. The
multi-GPU wait_and_throw() branch is unchanged.
Also drop the <cstdlib> include, which is unused.
Assisted-by: Claude Opus 5
---------
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
Raise the threshold for minimum buffer size from 1 GiB to 4 GiB, based
on real-world experiments of overcommitting device memory with model
weights larger than available VRAM, for example Qwen3.5-35B-A3B-Q8
running on a B70.
Also add a debug message to better track USM system allocations.
Signed-off-by: Francois Dugast <francois.dugast@intel.com>
* [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>
* support op col2im_1d
* update ops.md
* rm unused words
* update for bf16
* optimize 1%-11% as the review comments
* fix the format issue
* update as the review comments
* sycl: add supported types to ggml_sycl_supports_reorder_dmmv
The reordered feature is implemented in ggml_sycl_op_dequantize_mul_mat_vec,
but gated by ggml_sycl_supports_reorder_dmmv. This commit fixes the gate.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* sycl: set K_QUANTS_PER_ITERATION=1 to improve utilization
When combined with opening the reorder gate, this improves GPU
utilization on B70, giving a significant boost to tg t/s.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* sycl: replace QK_WARP_SIZE with WARP_SIZE for QK_5
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* sycl: add missing types to ggml_backend_sycl_buffer_init_tensor
Without this, the extra field is not allocated and the reorder path
will not take effect.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
---------
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
Todd Malsbary