Add SYCL backend - #22
Open
zjin-lcf wants to merge 2 commits into
Open
Conversation
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.
Contributor
|
Is there a use-case for SYCL? Our expectation was that those working on SYCL would consume the OpenCL kernels. |
- Launch the non-tiled cstream/bstream kernels via an explicit nd_range instead of a basic parallel_for. The Level Zero runtime otherwise auto-picks a poor work-group size for cstream, nearly halving its bandwidth (~113 -> ~210 GB/s on Arc B580); bstream is unaffected. - Re-assert __restrict on the pointers used inside the cstream/bstream kernel bodies (lost when the launcher pointers are captured by value). - Flatten the shared-memory local_accessor in bstream-msplit and cstream-ksplit from multi-dimensional to 1D with compile-time constant offsets. IGC does not constant-fold the runtime-held strides of a multi-dim local_accessor, so the flat form removes redundant address arithmetic (bstream-msplit ~208 -> ~257 GB/s on Arc B580). Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds a SYCL code-generation backend to GiMMiK, mirroring the structure of the existing OpenCL backend.
gimmik/sycl.py—SYCLMatMul(platform='sycl'), registered ingimmik/__init__.py.gimmik/kernels/sycl/— Mako templates for thecstream,bstream,bstream-msplit, andcstream-ksplitkernels. Each emits a self-contained launcher of the formsycl::event kname(sycl::queue& q, ...)that submits the kernel to a queue. Shared-memory kernels usesycl::local_accessorandnd_range.betaterm, both static-nand dynamic-nsignatures, and the fp32float2vectorised 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.float2vector kernels. All pass.