Skip to content

candle Intel GPU / SYCL support - #3974

Open
Hahihula wants to merge 3 commits into
huggingface:mainfrom
Hahihula:sycle-support
Open

Hahihula wants to merge 3 commits into
huggingface:mainfrom
Hahihula:sycle-support

Conversation

@Hahihula

Copy link
Copy Markdown

Adds a sycl feature that runs candle on Intel GPUs via oneAPI DPC++ / Level Zero.

Design

  • All SYCL code is behind #[cfg(feature = "sycl")], following the ROCm PR's cfg-gated pattern. Default cargo build is byte-identical and needs no oneAPI toolchain.
  • Device kernels live in candle-sycl-kernels (excluded from the workspace). build.rs compiles the csrc/*.cpp SYCL sources into one libcandle_sycl.so via icpx; Rust links it through a flat C ABI. A crate that only depends on this can
    be built by anyone; the kernel crate is the only thing that needs icpx.
  • New GPU-free CI job (ci_sycl.yaml) compile-checks the feature + clippy in the intel/oneapi-basekit container.

Coverage

  • BackendStorage / BackendDevice full impl: elementwise (19 unary / 6 binary), reductions, cmp, where, gather/scatter/index, sort, pooling, upsample, conv1d/2d + conv_transpose (im2col + oneMKL GEMM), matmul (oneMKL gemm /
    gemm_batch).
  • GGUF quantization: GPU dequant for all 12 block types, fused quantized mat-vec (MMVQ) for the decode path, quantize / embedding.
  • candle-nn fused ops: native softmax_last_dim, rms_norm, RoPE (3 variants) kernels; sycl_fwd for the rest.
  • Allocation pool on the queue; kernels are enqueue-only (sync at host copies / synchronize()).

Status

  • On a Meteor Lake Arc iGPU: tensor 35/37, quantized 26/26, matmul / conv / grad / pool / nn ops all green. Only F8E4M3 dtypes unimplemented.
  • Qwen3-0.6B Q4_K_M GGUF generates coherent text end-to-end at ~5 tok/s decode (flat with sequence length). Further perf work (MMVQ v2 / DP4a, per-op overhead) is tracked separately.
  • CPU path unchanged; no regressions.

proof of actual usability: lucasjinreal/Crane#112

- Add an integer mat-vec kernel (`mmvq_q8`): the activation is quantized to
  int8 once per call and each weight block is dotted against it in integer
  arithmetic, matching the CPU `vec_dot` numerics. Both mat-vec kernels now
  parallelise over (row, chunk of blocks) rather than one work-item per row.
- Keep a half-precision pipeline in f16 throughout: the mat-vec reads and
  writes f16, and prefill dequantizes straight to f16 (`dequantize_f16`)
  instead of materialising f32 and casting.
- Serve `m <= 32` from the mat-vec rather than dequantize-and-GEMM.
- Matmul an unquantized GGUF weight (F32/F16/BF16) in its own dtype instead of
  widening it to F32 at load; `QMatMul::is_quantized` tells the two apart.
- Add `get_rows`, so a quantized embedding dequantizes only the gathered rows.
- Accept a non-zero `start_offset` on the indexing kernels' operands (what
  `narrow` produces) instead of rejecting it.
- Bound the allocation pool, defer frees behind a synchronize (`sycl::free`
  does not wait for enqueued work), and drain-and-retry a failed allocation.
- Default the UR loader to the Level-Zero V1 adapter; V2 cannot launch kernels
  from an in-order USM queue on Xe2.
- Bake the oneAPI runtime dirs into the library as an rpath, so a binary runs
  without `setvars.sh`.
- Add the `sycl_parity` example, a SYCL-vs-CPU numerical check.
@Hahihula
Hahihula force-pushed the sycle-support branch 2 times, most recently from ddab996 to 6ae6c36 Compare October 2, 2026 05:17
@Hahihula
Hahihula force-pushed the sycle-support branch 2 times, most recently from 68c5a9e to 6cd5b7e Compare October 20, 2026 13:58

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant