Implementation:Ggml org Ggml Sycl convert
| Knowledge Sources | |
|---|---|
| Domains | ML_Infrastructure, GPU_Compute, Quantization |
| Last Updated | 2025-05-15 12:00 GMT |
Overview
SYCL kernels for dequantizing tensors from quantized formats to f32 or f16, including support for non-contiguous tensor layouts.
Description
convert.cpp implements the full-tensor dequantization path for the SYCL backend. When quantized tensors need to be converted to floating-point representation (for operations that lack native quantized kernels), this module provides the GPU kernels to do so. The implementation uses a layered template approach:
- dequantize_block: A device kernel that processes individual elements, computing block indices and quant sub-indices, then calling the appropriate dequantize_kernel function pointer to produce a dfloat2 pair of values.
- dequantize_block_sycl: A host-side launcher that computes work-group dimensions and dispatches the kernel via sycl::parallel_for.
- Per-type specializations: Dedicated functions (dequantize_row_q2_K_sycl, dequantize_row_q3_K_sycl, etc.) for K-quant types that require different work-group sizes.
- Function pointer lookup: ggml_get_to_fp32_sycl and ggml_get_to_fp16_sycl return the appropriate conversion function pointer for a given ggml_type, with optional BF16 support when compiled with Intel LLVM.
The file also supports non-contiguous dequantization kernels that handle strided memory access patterns.
Usage
Called by the main SYCL backend when tensor operations require dequantized inputs. The function pointer lookup functions are used during matrix multiplication setup to select the correct dequantization path based on the source tensor's quantization type.
Code Reference
Source Location
- Repository: GGML
- File: src/ggml-sycl/convert.cpp
- Lines: 676
Signatures
// Core dequantization kernel template
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, const sycl::nd_item<3> &item_ct1);
// SYCL launcher template
template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
static void dequantize_block_sycl(const void *__restrict__ vx,
dst_t *__restrict__ y, const int64_t k,
dpct::queue_ptr stream);
// Function pointer lookup
to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst);
to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor * dst);
I/O Contract
Inputs
| Name | Type | Required | Description |
|---|---|---|---|
| vx | const void * | Yes | Pointer to quantized source data |
| k | int64_t | Yes | Number of elements to dequantize |
| stream | dpct::queue_ptr | Yes | SYCL queue for kernel submission |
| type | ggml_type | Yes | Quantization type for function pointer lookup |
Outputs
| Name | Type | Description |
|---|---|---|
| y | dst_t * (float or sycl::half) | Dequantized output buffer in f32 or f16 |
Usage Examples
// Get the appropriate dequantization function for Q4_0 type
to_fp32_sycl_t to_fp32 = ggml_get_to_fp32_sycl(GGML_TYPE_Q4_0, tensor);
// Dequantize the tensor data to float
to_fp32(quantized_data, float_output, num_elements, sycl_stream);