[SYCL][Docs] Fix sycl_ext_oneapi_peer_access implementation and extension by steffenlarsen · Pull Request #19787 · intel/llvm

steffenlarsen · 2025-08-13T12:53:39Z

This commit fixes the following bugs in the specification and extension:

Changes the exceptions from repeat calls to enabling and disabling to undefined behavior. The implementation did not properly issue the specified exception before.
Implement the exception from attempting to enable access between devices that do not support peer access between them.
Make the access support query return false for backends that do not support it.
Specify and implement the relaxation that the access can be enabled when both devices are the same, even if the backend doesn't support P2P.
Specify that devices need to be from the same platform, rather than with the same backend.
Add the missing feature test macro.

…sion This commit fixes the following bugs in the specification and extension: * Changes the exceptions from repeat calls to enabling and disabling to undefined behavior. The implementation did not properly issue the specified exception before. * Implement the exception from attempting to enable access between devices that do not support peer access between them. * Make the access support query return false for backends that do not support it. * Specify and implement the relaxation that the access can be enabled when both devices are the same, even if the backend doesn't support P2P. * Specify that devices need to be from the same platform, rather than with the same backend. * Add the missing feature test macro. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2025-08-13T12:54:46Z

  if (Device != Peer) {
+    if (!ext_oneapi_can_access_peer(peer))
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Peer access is not allowed between the devices.");


@gmlueck - If we worry about the additional overhead of doing this check after requiring the user to do it first, we can make it UB instead.

Shouldn't this be feature_not_supported according to the extension spec?

Are we really worried about performance here? It seems like enabling P2P is something the application will just do once in some initialization code.

Shouldn't this be feature_not_supported according to the extension spec?

Ah, right you are!

Are we really worried about performance here? It seems like enabling P2P is something the application will just do once in some initialization code.

I don't particularly worry about the performance, as you are right that it should mainly be a one-and-done call. It is just redundant when the user is doing it correctly, as they are expected to have already checked the call.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

gmlueck · 2025-08-13T13:58:18Z

+`errc::feature_not_supported` error code.
+
+Calling this function with `peer` for which access has already been enabled will
+result in undefined behavior.


What will our implementation do in this case?

I looked back at the discussion in #6104, and it seems like the CUDA backend will return an error in this case. (Maybe that ends up throwing an exception?) What is the Level Zero behavior?

Same question for the case below when you disable access that was never enabled.

I believe CUDA will cause a backend error and L0 will just allow it through without issue. The issue with trying to make the L0 case return an error is that the tracking of which P2P pathways are enabled may become somewhat costly, while the user could do it themselves if they are concerned with overlapping enabling.

I'd like @pbalcer to weigh in here, because these things will actually go through UR rather than straight to L0.

My understanding is that the L0 adapter in UR will already have to keep track of which P2P pathways are enabled, in order to avoid sharing all USM allocations with all devices (see #19257).

If we have to track it anyway, I have no objections to moving back to having an exception for this.

Just following up ... have we decided yet whether we will report the exception vs. making this UB?

I lean towards keeping it an exception. I converted these changes to draft while we wait for #19257 to get back on track. Should happen shortly, from what I've been told.

Throwing an exception on repeated attempts to enable (or disable) is unfriendly to the case where an application and its library dependencies might want to collaborate in enabling P2P access for their respective needs. Enabling P2P access a second time doesn't create a problem for any party, for example.

I note that the analogous API in CUDA (cudaDeviceEnablePeerAccess) returns a status code that notes that it was a repeated enabling, and lets the caller decide whether that is actually a problem (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__PEER.html#group__CUDART__PEER_1g2b0adabf90db37e5cfddc92cbb2589f3)>. HIP presumably does likewise. I think following this relaxed approach has merit.

Throwing is particularly unfriendly if there is no way to query the status in advance. That forces a user writing a library to wrap the call in try...catch just in case the caller already enabled peer access. It is also brittle with respect to changes in defaults (e.g. if someone would ever decide that a multi-device context should imply that all devices have P2P access enabled). This SYCL extension has a query function for whether we can enable access but not for whether access is enabled. If we think throwing an exception is the right API for repeated enable or disable, we should add a query function for whether access is enabled. This is analogous to letting the user of std::vector call size() to find out if a planned access is within bounds so as to avoid that std::vector::at(i) actually throws, while retaining the safety of a throw if either caller or callee has a bug.

My preferred solution is to not throw in these cases, i.e leave it as undefined behaviour for now.

gmlueck · 2025-08-13T13:59:33Z

  if (Device != Peer) {
+    if (!ext_oneapi_can_access_peer(peer))
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Peer access is not allowed between the devices.");


Shouldn't this be feature_not_supported according to the extension spec?

Are we really worried about performance here? It seems like enabling P2P is something the application will just do once in some initialization code.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2025-09-01T05:52:47Z

Converting to draft while #19257 makes it in.

mabraham · 2026-03-13T14:18:17Z

+`errc::feature_not_supported` error code.
+
+Calling this function with `peer` for which access has already been enabled will
+result in undefined behavior.


Throwing an exception on repeated attempts to enable (or disable) is unfriendly to the case where an application and its library dependencies might want to collaborate in enabling P2P access for their respective needs. Enabling P2P access a second time doesn't create a problem for any party, for example.

I note that the analogous API in CUDA (cudaDeviceEnablePeerAccess) returns a status code that notes that it was a repeated enabling, and lets the caller decide whether that is actually a problem (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__PEER.html#group__CUDART__PEER_1g2b0adabf90db37e5cfddc92cbb2589f3)>. HIP presumably does likewise. I think following this relaxed approach has merit.

Throwing is particularly unfriendly if there is no way to query the status in advance. That forces a user writing a library to wrap the call in try...catch just in case the caller already enabled peer access. It is also brittle with respect to changes in defaults (e.g. if someone would ever decide that a multi-device context should imply that all devices have P2P access enabled). This SYCL extension has a query function for whether we can enable access but not for whether access is enabled. If we think throwing an exception is the right API for repeated enable or disable, we should add a query function for whether access is enabled. This is analogous to letting the user of std::vector call size() to find out if a planned access is within bounds so as to avoid that std::vector::at(i) actually throws, while retaining the safety of a throw if either caller or callee has a bug.

My preferred solution is to not throw in these cases, i.e leave it as undefined behaviour for now.

mabraham · 2026-03-13T14:21:25Z

This should return false, because if the P2P feature is not supported, then trivially P2P access is impossible.

steffenlarsen requested review from a team as code owners August 13, 2025 12:53

steffenlarsen requested a review from sergey-semenov August 13, 2025 12:53

steffenlarsen had a problem deploying to WindowsCILock August 13, 2025 12:53 — with GitHub Actions Error

steffenlarsen commented Aug 13, 2025

View reviewed changes

fix build error

f2d1f8f

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 13:08 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 13:43 — with GitHub Actions Inactive

gmlueck reviewed Aug 13, 2025

View reviewed changes

Stop requiring GPU devices

3f2e183

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 14:03 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 14:33 — with GitHub Actions Inactive

Fix errc

f13dc2e

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 14:45 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock August 13, 2025 15:11 — with GitHub Actions Inactive

steffenlarsen mentioned this pull request Aug 14, 2025

memory resident limit to enabled peers #19257

Closed

Remove die-call from OpenCL P2P

21c380f

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen requested a review from a team as a code owner August 14, 2025 03:36

steffenlarsen temporarily deployed to WindowsCILock August 14, 2025 03:36 — with GitHub Actions Inactive

steffenlarsen temporarily deployed to WindowsCILock August 14, 2025 04:17 — with GitHub Actions Inactive

kbenzie approved these changes Aug 14, 2025

View reviewed changes

steffenlarsen marked this pull request as draft September 1, 2025 05:52

steffenlarsen closed this Jan 12, 2026

mabraham reviewed Mar 13, 2026

View reviewed changes

mabraham mentioned this pull request Mar 13, 2026

[SYCL][UR] pool allocators bypass peer-access API #21395

Open

Uh oh!

Conversation

steffenlarsen commented Aug 13, 2025

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!

steffenlarsen commented Sep 1, 2025

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!

5 participants