[SYCL] Adding SYCL backend to support for Intel XPUs by abagusetty · Pull Request #215 · PASSIONLab/OpenEquivariance

abagusetty · 2026-09-03T17:54:01Z

Major changes:

Added feature support to target Intel GPUs (on ANL Aurora) via SYCL backend and generalize device selection so CUDA, HIP and SYCL are handled seamlessly

vbharadwaj-bk · 2026-09-04T01:35:14Z

Great feature, thanks! This will take some time to review and test - so bear with us, but with some back and forth we'll see where we land.

Kernels are generated as SYCL free functions and compiled at runtime through the oneAPI kernel compiler. The Jinja templates are shared with CUDA and HIP via templates/sycl_compat.cuh; the backend is now a string rather than an is_hip boolean, and JITKernel::execute also takes argument sizes.

Mirrors the CUDA job. GitHub offers no Intel GPU runner, so only the build is exercised.

…ement

The stable-ABI headers the extension includes are not all present before 2.10, which the precompiled-extension check in extlib has required all along. The SYCL guard now enforces the same floor.

Upstream PASSIONLab#217 replaced HIP's __threadfence_block() with the amdgcn wave barrier, so no template emits the token any more.

vbharadwaj-bk · 2026-09-17T02:31:18Z

Got a review in progress for this, about 20/50 files in; thanks for the patience, it's going

abagusetty · 2026-09-17T13:51:24Z

Got a review in progress for this, about 20/50 files in; thanks for the patience, it's going

Thanks @vbharadwaj-bk Appreciate for going through this. Not a small PR!

vbharadwaj-bk

@abagusetty I read every one of the + 1375 -270 line diff 😅 actually, I think this will work nicely with our current setup and I have left comments in the files. Much of this is mechanical; high-level, so you can have an LLM take care of some of this

LLMs tend to leave a lot of inline / block comments. Let's remove 99% of them, I think there were only one or two comments throughout the entire diff that I felt were necessary to understand something.
Let's roll back the changes to group_mm / symmetric contraction, as there is another PR in the queue to address the BLAS dependency and this should no longer be required. Or worst case, we won't support symmetric contraction on SYCL just yet.
There are just a few requests for code logic changes in certain files.
Let's hold off on the README / changelog things for now, we'll work with you on those once this batch of changes is in.

I didn't see any big logical gaps, but I also haven't tested this for regression on the existing backends; I will do that once this round of edits is done. In case there are bugs filed specifically for the SYCL backend, we will probably refer them to you to address, since we don't have Intel hardware. Are you ok with this, at least for a little while?

^^ Once this round of edits is done, and if we can refer bugs to you for some period of time, then I think with one more review, we can merge this.

vbharadwaj-bk · 2026-09-20T18:29:11Z

+    Builds the Jinja environment used to render the kernel templates.
+
+    :param backend: one of ``"cuda"``, ``"hip"`` or ``"sycl"``.
+    :param warp_size: size of a warp / wavefront / sub-group. Only consulted by


Hmm - warp size is needed by HIP as well.. but maybe that's a separate bug we need to address.

vbharadwaj-bk

Thanks for the thorough revision @abagusetty - things look good, added just a couple more --nodocuments (pesky, sorry), but there are no more logic changes remaining so I started the CUDA / HIP regression testing on our end.

Remaining items:

CUDA / HIP regression testing
Resolution of these last --nodocuments (in parallel)
(strictly after the last two) CI-ready label
Merge

vbharadwaj-bk · 2026-09-22T02:30:39Z


 class GroupMM:
    def __init__(self, dtype, num_elements, batch_size):
+        # group_gemm is not built for SYCL, so the operator does not exist there.


--nodocument, redundant

vbharadwaj-bk · 2026-09-22T02:30:54Z



-if extlib.BUILT_EXTENSION:
+# group_gemm is not built for SYCL, so the operator does not exist there.


--nodocument

vbharadwaj-bk · 2026-09-22T02:31:06Z

    else:
        from torch.profiler import ProfilerActivity, profile, record_function

+        # Profiler activity is per-accelerator.


--nodocument

vbharadwaj-bk · 2026-09-22T02:32:10Z

+namespace syclex = sycl::ext::oneapi::experimental;
+namespace twi = sycl::ext::oneapi::this_work_item;
+
+// The sub-group size is fixed to the warp size the schedule was generated


--nodocument

vbharadwaj-bk · 2026-09-22T02:32:41Z

+
+#define __launch_bounds__(...)
+
+// All generated kernels are launched as 1D nd_ranges, so only .x is meaningful.


--nodocument

vbharadwaj-bk · 2026-09-22T02:33:06Z

+    return ref.fetch_add(val);
+}
+
+// Templated on both operands to keep the mixed-width call sites working.


--nodocument

vbharadwaj-bk · 2026-09-22T02:36:29Z

+}
+
+/*
+* Registered after the first compile, not at static init: libsycl-jit is


--nodocument, got it, I think this actually becomes pretty clear just reading the code.

vbharadwaj-bk · 2026-09-22T05:41:31Z

Update: regression tests pass; with the short round of edits above, should be good to go.

abagusetty · 2026-09-22T11:32:30Z

Appreciate for the keen eye @vbharadwaj-bk, will go through rest of them. I was sure, I didnt address all of your comments :)

Happy to take on any bugs related to SYCL and `ll try to host a CI on intel hardware on our end or will just fall back to manual testing

abagusetty · 2026-09-22T14:18:25Z

@vbharadwaj-bk Thanks again for your time. I did manage to resolve most, please let me know if anything get missed

abagusetty added 8 commits September 16, 2026 13:33

ci: verify the SYCL extension build on a GPU-less runner

a552549

Mirrors the CUDA job. GitHub offers no Intel GPU runner, so only the build is exercised.

ci: pin the SYCL job to a verified PyTorch

ec8b0af

docs: align the SYCL PyTorch floor with the project's existing requir…

7db6cc1

…ement

docs: correct the AOTI PyTorch floor to 2.10

cb16f22

The stable-ABI headers the extension includes are not all present before 2.10, which the precompiled-extension check in extlib has required all along. The SYCL guard now enforces the same floor.

docs: shorten the CHANGELOG entry for the AOTI version floor

4f25be5

docs: simplify the CHANGELOG entry for the float64 literal suffix

353e326

fix: drop the dead __threadfence_block shim from sycl_compat.cuh

b7381ce

Upstream PASSIONLab#217 replaced HIP's __threadfence_block() with the amdgcn wave barrier, so no template emits the token any more.

abagusetty force-pushed the xpu-support branch from 69ae39e to b7381ce Compare September 16, 2026 18:36

vbharadwaj-bk requested changes Sep 20, 2026

View reviewed changes

abagusetty added 2 commits September 21, 2026 10:38

address PR comments

5fe3c77

fix SYCL CI: link libsycl explicitly and bump the torch pin to 2.13

ff37c65

vbharadwaj-bk reviewed Sep 22, 2026

View reviewed changes

abagusetty added 2 commits September 22, 2026 06:43

address a few more PR comments

a9f3912

cleanup a few more comments

f8e3afd



		if extlib.BUILT_EXTENSION:
		# group_gemm is not built for SYCL, so the operator does not exist there.


		#define __launch_bounds__(...)

		// All generated kernels are launched as 1D nd_ranges, so only .x is meaningful.

Conversation

abagusetty commented Sep 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

vbharadwaj-bk commented Sep 4, 2026

Uh oh!

vbharadwaj-bk commented Sep 17, 2026

Uh oh!

abagusetty commented Sep 17, 2026

Uh oh!

vbharadwaj-bk left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

vbharadwaj-bk left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

vbharadwaj-bk commented Sep 22, 2026

Uh oh!

abagusetty commented Sep 22, 2026

Uh oh!

abagusetty commented Sep 22, 2026

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

abagusetty commented Sep 3, 2026 •

edited

Loading

vbharadwaj-bk left a comment •

edited

Loading