[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-17T02:03:45Z

@@ -0,0 +1,9 @@
+# The SYCL backend follows the project's existing PyTorch >= 2.10 floor.


Let's remove this comment; LLMs tend to add a lot of comments, so I'll abbreviate requests to remove comments and keep code self-documenting with --selfdocument or --nodocument

vbharadwaj-bk · 2026-09-17T02:04:07Z

+        cache: 'pip'
+        cache-dependency-path: '**/requirements_sycl_ci.txt'
+
+    # The SYCL backend compiles the extension with the oneAPI DPC++ driver


--selfdocument

vbharadwaj-bk · 2026-09-17T02:06:57Z

+    # No Intel GPU is attached to this runner; only the build is exercised.
+    - name: Install oneAPI DPC++ and oneMKL
+      run: |
+        wget -qO- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \


Hmm - have you timed how long this takes, perhaps on your fork? If it takes too long, not thrilled about including it in the CI since our minutes are limited.

vbharadwaj-bk · 2026-09-17T02:08:23Z

-`OpenEquivariance <https://github.com/PASSIONLab/OpenEquivariance>`_ is a CUDA and 
-HIP kernel generator for the Clebsch-Gordon
+`OpenEquivariance <https://github.com/PASSIONLab/OpenEquivariance>`_ is a CUDA, 
+HIP, and SYCL kernel generator for the Clebsch-Gordon


Let's hold off README changes until the PR is in place; want to make sure that people know SYCL is in a beta state, and will receive less support (probably) from us as maintainers.

vbharadwaj-bk · 2026-09-17T02:29:55Z

+    Determines which GPU backend this PyTorch build targets.
+
+    Returns one of ``"cuda"``, ``"hip"`` or ``"sycl"``. HIP builds report a
+    ``torch.version.cuda`` of ``None``, so HIP must be tested first.


--nodocument

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 · 2026-09-20T18:30:02Z

 __launch_bounds__({{schedule.launch_config.num_threads}})
 {%- endmacro %}
+
+{# Declares the per-block shared memory buffer `s`. CUDA and HIP size the


--nodocument

vbharadwaj-bk · 2026-09-20T18:30:42Z

+    )
+    add_stable_extension(oeq_stable_sycl SYCL_BACKEND "${SYCL_LINK_LIBS}")
+
+    # -fsycl is required on both the compile and the link line.


--nodocument

vbharadwaj-bk · 2026-09-20T18:31:16Z

@@ -0,0 +1,115 @@
+{#
+Compatibility shim that lets the CUDA/HIP-flavored kernel templates compile as


--nodocument

vbharadwaj-bk · 2026-09-20T18:32:43Z

+    operator size_t() const { return x; }
+};
+
+static inline OeqIndex1D oeq_thread_idx() { return {twi::get_nd_item<1>().get_local_id(0)}; }


Hmm - I feel like a lot of these functions should perhaps be prefixed _sycl_thread_idx, or similar? Prefixing with oeq seems to imply distinct backends, whereas these functions are only called in the SYCL path.

Ditto for all remaining functions in this file.

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

		@@ -0,0 +1,9 @@
		# The SYCL backend follows the project's existing PyTorch >= 2.10 floor.

		@@ -0,0 +1,115 @@
		{#
		Compatibility shim that lets the CUDA/HIP-flavored kernel templates compile as

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!

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!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

abagusetty commented Sep 3, 2026 •

edited

Loading