Skip to content

Add SYCL backend - #22

Open
zjin-lcf wants to merge 1 commit into
PyFR:masterfrom
zjin-lcf:add-sycl-backend
Open

Add SYCL backend#22
zjin-lcf wants to merge 1 commit into
PyFR:masterfrom
zjin-lcf:add-sycl-backend

Conversation

@zjin-lcf

Copy link
Copy Markdown

Summary

Adds a SYCL code-generation backend to GiMMiK, mirroring the structure of the existing OpenCL backend.

  • gimmik/sycl.pySYCLMatMul (platform='sycl'), registered in gimmik/__init__.py.
  • gimmik/kernels/sycl/ — Mako templates for the cstream, bstream, bstream-msplit, and cstream-ksplit kernels. Each emits a self-contained launcher of the form sycl::event kname(sycl::queue& q, ...) that submits the kernel to a queue. Shared-memory kernels use sycl::local_accessor and nd_range.
  • Supports fp32/fp64, the beta term, both static-n and dynamic-n signatures, and the fp32 float2 vectorised variants — matching the OpenCL backend's feature set.
  • bench/ — a small harness (OpenCL vs SYCL) plus a correctness-validation suite.

The generated code compiles with any SYCL 2020 compiler (tested with Intel oneAPI DPC++).

Test plan

  • SYCLMatMul(...).kernels(dtype) generates all variants for fp32/fp64.
  • Correctness suite: every generated kernel verified against a NumPy reference across {fp32, fp64} x {beta=0, beta!=0} x {static-n, dynamic-n}, including the float2 vector kernels. All pass.
  • Runs on an Intel Data Center GPU (Level Zero and OpenCL runtimes).

Add a SYCL code-generation backend mirroring the existing OpenCL backend.

- gimmik/sycl.py: SYCLMatMul (platform='sycl'), registered in __init__
- gimmik/kernels/sycl/: cstream, bstream, bstream-msplit, cstream-ksplit
  Mako templates emitting self-contained launcher functions of the form
  `sycl::event kname(sycl::queue& q, ...)`, using local_accessor and
  nd_range for the shared-memory kernels. Supports fp32/fp64, beta,
  static- and dynamic-n signatures, and the fp32 float2 vector variants.
- bench/: OpenCL-vs-SYCL benchmark and a correctness-validation suite that
  checks every generated kernel against a NumPy reference.
- README: mention SYCL support.
@FreddieWitherden

Copy link
Copy Markdown
Contributor

Is there a use-case for SYCL? Our expectation was that those working on SYCL would consume the OpenCL kernels.

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.

2 participants