SYCL: add TurboQuant KV cache support for Intel GPUs

README.md

@@ -62,6 +62,7 @@ All turbo formats use Walsh-Hadamard rotation followed by polar codebook quantiz
 | **CUDA** (NVIDIA) | `dp4a` for `TQ4_1S`, warp-cooperative dequant (16× less compute per block), multi-token / multi-GPU | Yes — turbo VEC FA (+9% decode); mixed `f16/bf16 + q8_0` without `GGML_CUDA_FA_ALL_QUANTS` | Load-time `TQ4_1S → q8_0` conversion path |
 | **HIP / ROCm** (AMD) | Portable `ggml_cuda_dp4a`; scalar half path for `TQ4_1S` on AMD | Yes — VEC FA forced for quantized KV; pool bypass for FA f16 temp buffers | RDNA3 (gfx1100), RDNA4, CDNA3 (MI300X / gfx942), CDNA4 (MI355X / gfx950). [cross-engine-mi300x](https://github.com/TheTom/turboquant_plus/blob/main/docs/papers/cross-engine-mi300x.md) |
 | **Vulkan** | `TQ4_1S` weights, `SET_ROWS` for `turbo2`/`turbo4` | coopmat flash attention with `turbo3` KV | Compute-shader path; nix-buildable |
+| **SYCL** (Intel) | `SET_ROWS` for `turbo2`/`turbo3`/`turbo4`, WHT rotation (group 64+128) | VEC FA for all turbo K/V combos (D=64-512) | Intel Arc (A380, B70), oneAPI 2025.2+; cross-turbo K/V combos supported |
 
 ### Model-family support
 
@@ -111,6 +112,11 @@ cmake -B build -DGGML_HIP=ON -DCMAKE_HIP_ARCHITECTURES="gfx1100;gfx942;gfx950" &
 
 # Vulkan
 cmake -B build -DGGML_VULKAN=ON && cmake --build build -j
+
+# Intel SYCL (oneAPI)
+source /opt/intel/oneapi/setvars.sh
+cmake -B build -DGGML_SYCL=ON -DGGML_SYCL_F16=ON \
+  -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx && cmake --build build -j
 ```
 
 ## Usage

ggml/src/ggml-sycl/CMakeLists.txt

@@ -162,8 +162,24 @@ if (GGML_SYCL_F16)
     add_compile_definitions(GGML_SYCL_F16)
 endif()
 
+option(GGML_SYCL_SCALAR_MMQ "Force scalar MMQ tile sizes (for B70 without XMX kernels)" OFF)
+if (GGML_SYCL_SCALAR_MMQ)
+    add_compile_definitions(GGML_SYCL_SCALAR_MMQ)
+endif()
+
+option(GGML_SYCL_TURBO_QUANT "Enable dispatch to TurboQuant KV cache FA kernels" ON)
+if (GGML_SYCL_TURBO_QUANT)
+    add_compile_definitions(GGML_SYCL_TURBO_QUANT)
+endif()
+
+option(GGML_SYCL_SG32 "Use sub-group size 32 on Intel GPU (experimental)" OFF)
+
 if (GGML_SYCL_TARGET STREQUAL "INTEL")
-    add_compile_definitions(GGML_SYCL_WARP_SIZE=16)
+    if (GGML_SYCL_SG32)
+        add_compile_definitions(GGML_SYCL_WARP_SIZE=32)
+    else()
+        add_compile_definitions(GGML_SYCL_WARP_SIZE=16)
+    endif()
     if (NOT GGML_SYCL_DEVICE_ARCH)
         target_link_options(ggml-sycl PRIVATE -Xs -ze-intel-greater-than-4GB-buffer-required)
     else()

ggml/src/ggml-sycl/common.hpp

@@ -103,7 +103,7 @@ extern int g_ggml_sycl_enable_flash_attention;
 
 // define for XMX in Intel GPU
 // TODO: currently, it's not used for XMX really.
-#if !defined(GGML_SYCL_FORCE_MMQ)
+#if !defined(GGML_SYCL_FORCE_MMQ) && !defined(GGML_SYCL_SCALAR_MMQ)
     #define SYCL_USE_XMX
 #endif
 

ggml/src/ggml-sycl/convert.cpp

@@ -1,6 +1,7 @@
 #include "convert.hpp"
 #include "dequantize.hpp"
 #include "presets.hpp"
+#include "turbo-quant.hpp"
 
 template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
 static void dequantize_block(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t k,
@@ -640,6 +641,74 @@ static void convert_unary_sycl(const void * vx, dst_t * y, const int64_t k, dpct
 }
 
 
+// ---- TurboQuant dequantization kernels ----
+
+template <typename dst_t>
+static void dequantize_block_turbo2_0_kernel(const void * __restrict__ vx, dst_t * __restrict__ y,
+                                             const int64_t k, const sycl::nd_item<3> & item) {
+    const int64_t ib = item.get_group(2);
+    const int j = item.get_local_id(2);
+    const auto * x = static_cast<const block_turbo2_0 *>(vx) + ib;
+    const int64_t y_offset = ib * QK_TURBO2;
+    if (y_offset + j >= k) return;
+    const float norm = static_cast<float>(x->norm);
+    y[y_offset + j] = static_cast<dst_t>(turbo2_dequant_element(x, j, norm));
+}
+
+template <typename dst_t>
+static void dequantize_block_turbo3_0_kernel(const void * __restrict__ vx, dst_t * __restrict__ y,
+                                             const int64_t k, const sycl::nd_item<3> & item) {
+    const int64_t ib = item.get_group(2);
+    const int j = item.get_local_id(2);
+    const auto * x = static_cast<const block_turbo3_0 *>(vx) + ib;
+    const int64_t y_offset = ib * QK_TURBO3;
+    if (y_offset + j >= k) return;
+    const float norm = static_cast<float>(x->norm);
+    y[y_offset + j] = static_cast<dst_t>(turbo3_dequant_element(x, j, norm));
+}
+
+template <typename dst_t>
+static void dequantize_block_turbo4_0_kernel(const void * __restrict__ vx, dst_t * __restrict__ y,
+                                             const int64_t k, const sycl::nd_item<3> & item) {
+    const int64_t ib = item.get_group(2);
+    const int j = item.get_local_id(2);
+    const auto * x = static_cast<const block_turbo4_0 *>(vx) + ib;
+    const int64_t y_offset = ib * QK_TURBO4;
+    if (y_offset + j >= k) return;
+    const float norm = static_cast<float>(x->norm);
+    y[y_offset + j] = static_cast<dst_t>(turbo4_dequant_element(x, j, norm));
+}
+
+template <typename dst_t>
+static void dequantize_row_turbo2_0_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
+    GGML_ASSERT(k > 0 && k % QK_TURBO2 == 0);
+    const int64_t nb = k / QK_TURBO2;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, QK_TURBO2),
+                          sycl::range<3>(1, 1, QK_TURBO2)),
+        [=](sycl::nd_item<3> item) { dequantize_block_turbo2_0_kernel(vx, y, k, item); });
+}
+
+template <typename dst_t>
+static void dequantize_row_turbo3_0_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
+    GGML_ASSERT(k > 0 && k % QK_TURBO3 == 0);
+    const int64_t nb = k / QK_TURBO3;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, QK_TURBO3),
+                          sycl::range<3>(1, 1, QK_TURBO3)),
+        [=](sycl::nd_item<3> item) { dequantize_block_turbo3_0_kernel(vx, y, k, item); });
+}
+
+template <typename dst_t>
+static void dequantize_row_turbo4_0_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
+    GGML_ASSERT(k > 0 && k % QK_TURBO4 == 0);
+    const int64_t nb = k / QK_TURBO4;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, QK_TURBO4),
+                          sycl::range<3>(1, 1, QK_TURBO4)),
+        [=](sycl::nd_item<3> item) { dequantize_block_turbo4_0_kernel(vx, y, k, item); });
+}
+
 to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
     switch (type) {
         case GGML_TYPE_Q4_0:
@@ -710,6 +779,12 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
             return dequantize_row_mxfp4_sycl;
         case GGML_TYPE_NVFP4:
             return dequantize_row_nvfp4_sycl;
+        case GGML_TYPE_TURBO2_0:
+            return dequantize_row_turbo2_0_sycl;
+        case GGML_TYPE_TURBO3_0:
+            return dequantize_row_turbo3_0_sycl;
+        case GGML_TYPE_TURBO4_0:
+            return dequantize_row_turbo4_0_sycl;
         case GGML_TYPE_F32:
             return convert_unary_sycl<float>;
 #ifdef GGML_SYCL_HAS_BF16
@@ -793,6 +868,12 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
             return dequantize_row_mxfp4_sycl;
         case GGML_TYPE_NVFP4:
             return dequantize_row_nvfp4_sycl;
+        case GGML_TYPE_TURBO2_0:
+            return dequantize_row_turbo2_0_sycl;
+        case GGML_TYPE_TURBO3_0:
+            return dequantize_row_turbo3_0_sycl;
+        case GGML_TYPE_TURBO4_0:
+            return dequantize_row_turbo4_0_sycl;
         case GGML_TYPE_F16:
             return convert_unary_sycl<sycl::half>;
 #ifdef GGML_SYCL_HAS_BF16