SYCL: add TurboQuant KV cache support for Intel GPUs

[gagarinlg]

Overview

Add complete TurboQuant KV cache support to the SYCL backend for Intel GPUs: quantized SET_ROWS kernels (turbo2/3/4), Walsh-Hadamard rotation, VEC flash attention with turbo K/V dequant, and all cross-type FA template instantiations (21 files). Includes GGML_SYCL_TURBO_QUANT cmake gate, GGML_SYCL_SCALAR_MMQ for Battlemage tile sizes, and supports_op guards for safe scheduler fallback.

Tested on Intel Arc Pro B70 (Battlemage G31, 32GB VRAM, oneAPI 2025.3). Lloyd-Max optimal centroids match CUDA reference.

Benchmarks (Gemma 4 12B QAT Q4_0, tg100, ngl 99, Intel Arc Pro B70):

KV cache (K / V)	Decode (t/s)	vs f16
f16 / f16	51.49	baseline
q8_0 / q8_0	49.12	-4.6%
turbo3 / turbo3	46.37	-9.9%
turbo2 / turbo2	43.77	-15.0%
turbo4 / turbo4	40.53	-21.3%

Perplexity (wikitext-2-raw, 50 chunks, n_ctx=512): turbo3/turbo3 PPL 318.29 vs q8_0 baseline 280.78 (+13.4%, consistent with turbo3 compression ratio on 12B model).

Coherence: -ctk turbo3 -ctv turbo3 produces correct reasoning output.

Additional information

The CUDA upstream refactored FA type validation to check both K and V types (a395d83) — this SYCL implementation uses the same pattern. Vulkan turbo4 centroids were also updated upstream to matching Lloyd-Max values (175a652).

A proper code review is required.

Build config:

cmake -B build \
  -DGGML_SYCL=ON -DGGML_SYCL_F16=ON -DGGML_SYCL_DNN=ON \
  -DGGML_SYCL_GRAPH=ON -DGGML_SYCL_SCALAR_MMQ=ON \
  -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES

[Copilot]

Pull request overview

Adds TurboQuant KV-cache support to the SYCL backend (Intel GPUs), including SYCL-side TurboQuant constants/dequant helpers, TurboQuant SET_ROWS quantization kernels, a SYCL Walsh-Hadamard transform op, and extensive FlashAttention VEC template instantiations/dispatch for turbo K/V type combinations.

Changes:

• Implement SYCL Walsh-Hadamard transform (GGML_OP_TURBO_WHT) and wire it into SYCL op dispatch / supports-op checks.
• Add TurboQuant centroids/WHT sign tables and turbo2/3/4 dequant helpers, and use them in SYCL SET_ROWS + FlashAttention VEC paths.
• Add SYCL build/config switches (GGML_SYCL_TURBO_QUANT, GGML_SYCL_SCALAR_MMQ, optional SG32 warp size) and document SYCL support in the README.

Reviewed changes

Copilot reviewed 34 out of 34 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
README.md	Documents SYCL TurboQuant coverage and a oneAPI build example
ggml/src/ggml-sycl/turbo-wht.hpp	Declares SYCL Turbo WHT entry point
ggml/src/ggml-sycl/turbo-wht.cpp	Implements SYCL Turbo WHT kernels + dispatch
ggml/src/ggml-sycl/turbo-quant.hpp	Adds TurboQuant constants/sign tables + inline centroid/dequant helpers
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo4_0-turbo4_0.cpp	Explicit VEC FA instantiations for turbo4/turbo4
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo4_0-turbo3_0.cpp	Explicit VEC FA instantiations for turbo4/turbo3
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo4_0-turbo2_0.cpp	Explicit VEC FA instantiations for turbo4/turbo2
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo4_0-q8_0.cpp	Explicit VEC FA instantiations for turbo4/q8_0
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo4_0-f16.cpp	Explicit VEC FA instantiations for turbo4/f16
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo3_0-turbo4_0.cpp	Explicit VEC FA instantiations for turbo3/turbo4
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cpp	Explicit VEC FA instantiations for turbo3/turbo3
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo3_0-turbo2_0.cpp	Explicit VEC FA instantiations for turbo3/turbo2
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo3_0-q8_0.cpp	Explicit VEC FA instantiations for turbo3/q8_0
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo3_0-f16.cpp	Explicit VEC FA instantiations for turbo3/f16
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo2_0-turbo4_0.cpp	Explicit VEC FA instantiations for turbo2/turbo4
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo2_0-turbo3_0.cpp	Explicit VEC FA instantiations for turbo2/turbo3
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo2_0-turbo2_0.cpp	Explicit VEC FA instantiations for turbo2/turbo2
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo2_0-q8_0.cpp	Explicit VEC FA instantiations for turbo2/q8_0
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-turbo2_0-f16.cpp	Explicit VEC FA instantiations for turbo2/f16
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-q8_0-turbo4_0.cpp	Explicit VEC FA instantiations for q8_0/turbo4
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-q8_0-turbo3_0.cpp	Explicit VEC FA instantiations for q8_0/turbo3
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-q8_0-turbo2_0.cpp	Explicit VEC FA instantiations for q8_0/turbo2
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-f16-turbo4_0.cpp	Explicit VEC FA instantiations for f16/turbo4
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-f16-turbo3_0.cpp	Explicit VEC FA instantiations for f16/turbo3
ggml/src/ggml-sycl/template-instances/fattn-vec-instance-f16-turbo2_0.cpp	Explicit VEC FA instantiations for f16/turbo2
ggml/src/ggml-sycl/set_rows.cpp	Adds TurboQuant SET_ROWS cooperative quantization kernels and dispatch cases
ggml/src/ggml-sycl/mmvq.cpp	Adjusts workgroup/subgroup sizing for varying WARP_SIZE (Battlemage)
ggml/src/ggml-sycl/ggml-sycl.cpp	Registers Turbo WHT op and extends supports-op validation for turbo paths
ggml/src/ggml-sycl/fattn.cpp	Adds turbo K/V dispatch cases and selection guards for VEC FA
ggml/src/ggml-sycl/fattn-vec.hpp	Adds turbo-aware VEC FA logic (turbo LUT path + type handling)
ggml/src/ggml-sycl/fattn-common.hpp	Adds turbo KQ dot + turbo V dequant specializations for VEC FA
ggml/src/ggml-sycl/convert.cpp	Adds turbo2/3/4 dequant-to-fp16/fp32 conversion kernels
ggml/src/ggml-sycl/common.hpp	Disables XMX path when scalar MMQ is requested
ggml/src/ggml-sycl/CMakeLists.txt	Adds SYCL TurboQuant/scalar MMQ/SG32 options and WARP_SIZE selection

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[gagarinlg]

@cclecle maybe you want to test this SYCL port.

[Copilot]

Pull request overview

Copilot reviewed 34 out of 34 changed files in this pull request and generated 4 comments.

[Copilot]

Pull request overview

Copilot reviewed 34 out of 34 changed files in this pull request and generated 7 comments.

Comments suppressed due to low confidence (1)

ggml/src/ggml-sycl/mmvq.cpp:1595

• This reorder kernel asserts block_num_y % num_subgroups == 0, which effectively requires nrows to be a multiple of the workgroup-row granularity (since GGML_SYCL_MMV_Y==1). The underlying mul_mat_vec_q_reorder kernel already checks if (row >= nrows) return;, so it's safe (and more robust for split tails) to round block_num_y up instead of aborting.

    constexpr size_t num_subgroups = 256 / WARP_SIZE;
    GGML_ASSERT(block_num_y % num_subgroups == 0);

    const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, block_num_y * WARP_SIZE);
    const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);

[Copilot]

Pull request overview

Copilot reviewed 34 out of 34 changed files in this pull request and generated no new comments.

Comments suppressed due to low confidence (1)

ggml/src/ggml-sycl/fattn-common.hpp:13

• fattn-common.hpp uses std::is_same_v in the newly added turbo dequant helpers, but the header does not include <type_traits>. This is undefined to rely on transitive includes from SYCL/DPCT headers and can break builds with different toolchains or header versions.

#include "turbo-quant.hpp"

#include "ggml.h"

#include <cstdint>

[Copilot]

Pull request overview

Copilot reviewed 34 out of 34 changed files in this pull request and generated 1 comment.

[cclecle]

@cclecle maybe you want to test this SYCL port.

To be honest I gave up on my A380 because it is too limited and bought 2x R9700 :).
But I can definitly test this to help validating the changes. Will try soon.

[TheTom]

Thanks @gagarinlg! Merged to feature/turboquant-kv-cache as dd9b27c60.

Isolation check before merge:

• Diff is only ggml/src/ggml-sycl/* plus README backend docs
• GGML_SYCL stays OFF by default, so Metal / CUDA / HIP / Vulkan trees do not compile this path
• Clean auto-merge onto tip; Metal llama-server rebuild on macOS still green after the merge

Appreciate the Arc Pro B70 benches, Lloyd-Max centroid alignment with CUDA, and the follow-up validation commits.

[TheTom]

Thanks again @gagarinlg. Really solid SYCL TurboQuant work. Happy to have it in.

[gagarinlg]

Don't hesitate to ping me, in case an issue in that code comes up.