UPSTREAM PR #20571: sycl: add GGML_OP_GATED_DELTA_NET fused kernel

[loci-dev]

Note

Source pull request: ggml-org/llama.cpp#20571

Summary

• Port the Gated Delta Net (GDN) recurrence kernel from the Vulkan compute shader (gated_delta_net.comp) to the SYCL backend
• Enables Qwen3.5 and other delta-net architecture models to run efficiently on Intel GPUs via oneAPI
• Previously, the SYCL backend had no GDN support, causing fallback to CPU for these operations

Implementation

New files:

• ggml/src/ggml-sycl/gdn.cpp — fused kernel implementation
• ggml/src/ggml-sycl/gdn.hpp — header

Modified files:

• ggml/src/ggml-sycl/backend.hpp — add include
• ggml/src/ggml-sycl/ggml-sycl.cpp — add dispatch case and supports_op entry

Kernel features:

• Both GDA (scalar gate) and KDA (key-dependent / vector gate) variants
• Head sizes 32, 64, 128 via compile-time templates
• GQA/MQA support through stride-based tensor access (matching Vulkan push constants layout)
• sycl::float4 vectorized inner loops (same pattern as existing gla.cpp)
• One workgroup per (head, seq), S_V threads per workgroup, state held in registers

Benchmark

Tested on Intel Arc 140V (Lunar Lake iGPU) with Qwen3.5-0.8B-Q4_K_M, -ngl 99:

Metric	Before (no GDN kernel)	After (this PR)	Change
Decode (tok/s)	22.0	54.0	+145%
Prompt (tok/s)	8.6	23.1	+169%

The decode improvement comes from GDN layers now running as a fused kernel on GPU instead of falling back to per-op CPU execution.

Test plan

• Builds successfully with oneAPI 2025.3 + Ninja on Windows
• End-to-end inference with Qwen3.5-0.8B produces coherent output
• test-backend-ops passes GATED_DELTA_NET tests (test cases already exist in upstream)
• CI with SYCL build

🤖 Generated with Claude Code