Principle:Ggml org Ggml SYCL GPU Computation

Field	Value
sources	GGML SYCL Specification Intel oneAPI oneDNN
domains	GPU, Intel, SYCL
last_updated	2026-02-10

Overview

SYCL GPU Computation is the principle of running tensor operations on Intel GPUs (and other SYCL-compatible devices) via the SYCL/oneAPI programming model, with optional integration of oneDNN for optimized primitives and oneMKL for BLAS operations.

Description

SYCL is an open standard for heterogeneous parallel programming built on top of standard C++. It provides a single-source programming model where host code and device kernels are written in the same C++ source file, using lambda expressions to define parallel regions. Intel's oneAPI toolkit provides the primary SYCL implementation (DPC++, based on LLVM/Clang) targeting Intel GPUs, CPUs, and FPGAs.

The GGML SYCL backend enables inference on Intel discrete GPUs (Arc series), integrated GPUs, and potentially other SYCL-compatible devices. It is structured around several key components:

SYCL Queues and Devices

SYCL uses a queue-based execution model:

sycl::device -- Represents a compute device (GPU, CPU, etc.)
sycl::queue -- A command queue bound to a device; kernels and memory operations are submitted to queues
sycl::context -- Shared state across queues for the same platform

The GGML backend maintains per-device state including SYCL queues, memory pools, and scratch buffers.

Kernel Implementation

SYCL kernels are written as C++ lambda expressions submitted via queue.parallel_for() or queue.submit(). The GGML SYCL backend implements tensor operations using two approaches:

Hand-written SYCL kernels -- Custom parallel_for kernels for operations like element-wise ops, quantized dot products, copy/conversion, and specialized matrix-vector multiply
oneDNN integration -- For performance-critical operations like matrix multiplication, the backend delegates to Intel's oneDNN (Deep Neural Network Library), which provides highly optimized GPU implementations with automatic kernel selection and tuning

oneDNN Integration

When GGML_SYCL_DNNL is enabled, the backend uses oneDNN for:

Matrix multiplication (matmul primitive)
Convolution operations
Other operations with efficient oneDNN implementations

oneDNN primitives are created with memory descriptors that specify tensor layouts and data types, and are executed on SYCL queues via the dnnl::sycl_interop API. This leverages Intel's deep hardware expertise for GPU kernel optimization.

oneMKL Integration

For BLAS operations, the backend can use oneMKL (oneAPI Math Kernel Library) to call optimized GEMM routines on Intel GPUs, similar to how the BLAS backend uses CPU BLAS libraries.

Memory Management

SYCL provides several memory models:

USM (Unified Shared Memory) -- malloc_device/malloc_shared for pointer-based memory management
Buffer/Accessor model -- SYCL buffers with accessor objects for automatic data dependency tracking

The GGML backend primarily uses USM device allocations (sycl::malloc_device) for explicit memory management, with host-device transfers via queue.memcpy().

Usage

Apply SYCL GPU computation when:

Targeting Intel discrete GPUs (Arc A-series, Data Center GPUs)
Targeting Intel integrated GPUs (Iris Xe, UHD)
The Intel oneAPI toolkit (DPC++ compiler, oneDNN, oneMKL) is installed
A portable solution across Intel hardware generations is needed

SYCL is the recommended path for Intel GPU inference because:

oneDNN provides Intel-optimized matrix multiply and convolution kernels
The oneAPI ecosystem provides a complete development stack (compiler, profiler, debugger)
SYCL supports both integrated and discrete Intel GPUs with the same code

Theoretical Basis

The SYCL execution model for GGML tensor operations:

 Initialization:
 1. Device discovery:
    Enumerate SYCL devices:
      auto devices = sycl::device::get_devices(sycl::info::device_type::gpu)
    Select Intel GPU devices based on vendor ID and device properties

 2. Queue creation:
    For each selected device:
      sycl::queue q(device, exception_handler, {in_order_property})

 3. oneDNN engine setup (if DNNL enabled):
    dnnl::engine eng = dnnl::sycl_interop::make_engine(device, context)
    dnnl::stream strm = dnnl::sycl_interop::make_stream(eng, queue)

 4. Memory allocation:
    Tensor buffers allocated via sycl::malloc_device(size, queue)

 Graph Execution:
 For each node in the computation graph:

   Case 1: oneDNN-accelerated operation (e.g., MUL_MAT):
   a. Create oneDNN memory descriptors for input/output tensors
      dnnl::memory::desc src_md({M, K}, data_type, format_tag)
      dnnl::memory::desc wei_md({K, N}, data_type, format_tag)
      dnnl::memory::desc dst_md({M, N}, data_type, format_tag)

   b. Create matmul primitive descriptor
      dnnl::matmul::primitive_desc pd(eng, src_md, wei_md, dst_md)

   c. Create and execute primitive
      dnnl::matmul matmul_prim(pd)
      matmul_prim.execute(stream, {
        {DNNL_ARG_SRC, src_mem},
        {DNNL_ARG_WEIGHTS, wei_mem},
        {DNNL_ARG_DST, dst_mem}
      })

   Case 2: Custom SYCL kernel (e.g., element-wise op):
   a. Submit parallel_for kernel:
      queue.parallel_for(sycl::range<1>(n_elements), [=](sycl::id<1> idx) {
        // Each work-item processes one element
        dst[idx] = activation_fn(src[idx])
      })

   Case 3: Quantized operations:
   a. Submit kernel with subgroup operations:
      queue.parallel_for(sycl::nd_range<1>(global, local), [=](sycl::nd_item<1> item) {
        auto sg = item.get_sub_group()
        // Cooperative dequantize and dot product within subgroup
        float partial = quantized_dot(src0_q, src1_f, item)
        // Subgroup reduction
        float result = sycl::reduce_over_group(sg, partial, sycl::plus<>())
        if (sg.get_local_linear_id() == 0)
          dst[row] = result
      })

 Synchronization:
 queue.wait()  // Wait for all submitted operations to complete

 Memory Transfer:
 Host to device: queue.memcpy(dev_ptr, host_ptr, size).wait()
 Device to host: queue.memcpy(host_ptr, dev_ptr, size).wait()

Related Pages

Implementation:Ggml_org_Ggml_Sycl_backend
Ggml_org_Ggml_Sycl_backend -- The backend implementation that applies this principle
Ggml_org_Ggml_Vulkan_GPU_Computation -- Alternative cross-platform GPU compute using Vulkan
Ggml_org_Ggml_BLAS_Matrix_Multiplication -- CPU-side BLAS approach that SYCL/oneMKL extends to GPU

Page Connections

Double-click a node to navigate. Hold to expand connections.

Principle

Implementation

Heuristic

Environment