Add SYCL backend - #22
Open
zjin-lcf wants to merge 5 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>
Annotate the SYCL kernel lambdas with [[sycl::reqd_work_group_size(...)]] so the launch geometry is known at compile time. On AMD (DPC++ amdgcn backend) this lets the compiler size the register file for the actual work-group instead of the conservative default, eliminating the register spilling that otherwise appeared in the cstream/bstream kernels. The hint matches the geometry already used at launch, so Intel and NVIDIA codegen are unaffected. Co-authored-by: Cursor <cursoragent@cursor.com>
- Emit width-2 (sycl::double2 / sycl::float2) variants of every kernel when the leading dimension is even, at both single and double precision (previously float2 only, and only for cstream/bstream-msplit). - Add preload-C variants of bstream-msplit and cstream-ksplit that load C up-front; emitted only for beta != 0, as extra autotune candidates. - Hoist every used row of B into a register in cstream before the dot products. IGC does not eliminate redundant vector-typed (width > 1) loads, so the previous form reloaded B once per non-zero and made the width-2 cstream memory-bound (~66 -> ~190 GB/s on Arc B580); width-1 also improves (~210 -> ~226 GB/s). Co-authored-by: Cursor <cursoragent@cursor.com>
- bench_gen.py: honour GMK_BETA / GMK_ALIGNE so the width-2 and preload-C kernels are emitted; wrap width>1 kernels (which take vector pointers) in a uniform double* entry point; tag preloaded kernels; and wipe the kernel dir between runs so a stale kernel from a different config is never picked up by the build glob. - validate_gen.py: add fp64 double2 and preload cases (beta != 0) so the correctness suite covers the new variants. - ocl_bench.cpp: allow overriding the matched OpenCL vendor via GMK_OCL_VENDOR (defaults to Intel). 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, with a compile-time[[sycl::reqd_work_group_size]]annotation.betaterm, and both static-nand dynamic-nsignatures.Kernel variants / tuning
sycl::double2/sycl::float2) variants of every kernel when the leading dimension is even, at both single and double precision.bstream-msplitandcstream-ksplit(load C up-front), emitted forbeta != 0as extra autotune candidates.nd_rangelaunches, 1D flattenedlocal_accessors with compile-time offsets, and hoisting the B-row loads incstreamso the width-2 kernel does not become memory-bound (IGC does not CSE vector-typed loads). Per-kernel Arc B580 numbers are in the commit messages.Benchmarks (
bench/)Test plan
SYCLMatMul(...).kernels(dtype)generates all variants for fp32/fp64.float2/double2) and preload-C variants. All pass (70/70 SYCL, 46/46 OpenCL).