From 834a25fbb8b476826077df0f7d1ae75767828902 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Aug 2025 05:45:57 -0700 Subject: [PATCH 1/5] [SYCL][Docs] Fix sycl_ext_oneapi_peer_access implementation and extension 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 --- .../sycl_ext_oneapi_peer_access.asciidoc | 16 ++++++++++------ sycl/source/device.cpp | 12 ++++++++++-- sycl/source/feature_test.hpp.in | 1 + sycl/test-e2e/USM/P2P/p2p_access.cpp | 1 - sycl/test-e2e/USM/P2P/p2p_atomics.cpp | 1 - 5 files changed, 21 insertions(+), 10 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_peer_access.asciidoc b/sycl/doc/extensions/supported/sycl_ext_oneapi_peer_access.asciidoc index ebacc013bbd9a..be1f26564e9bc 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_peer_access.asciidoc +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_peer_access.asciidoc @@ -49,8 +49,8 @@ This extension is implemented and fully supported by DPC++. This extension is currently implemented in DPC++ for all GPU devices and backends; however, only the CUDA, HIP and Level Zero backends allows peer to peer memory access. Other backends report false from the -`ext_oneapi_can_access_peer` query. Peer-Peer memory access is only possible -between two devices from the same backend. +`ext_oneapi_can_access_peer` query, unless both devices are the same. Peer +memory access is only possible between two devices from the same SYCL platform. == Overview @@ -153,13 +153,17 @@ functions may access USM device allocations on the peer device subject to the normal rules about context as described in the core SYCL specification. If this device does not support peer access (as defined by `peer_access::access_supported`), throws an `exception` with the -`errc::feature_not_supported` error code. If access is already enabled, -throws an exception with the `errc::invalid` error code. +`errc::feature_not_supported` error code. + +Calling this function with `peer` for which access has already been enabled will +result in undefined behavior. |void ext_oneapi_disable_peer_access(const device &peer) -|Disables access to the peer device's memory from this device. If peer access -is not enabled, throws an `exception` with the `errc::invalid` error code. +|Disables access to the peer device's memory from this device. + +Calling this function with `peer` for which access is not enabled will result in +undefined behavior. |=== diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index dc206b55dac44..4fb6cd6a66b3a 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -220,6 +220,9 @@ void device::ext_oneapi_enable_peer_access(const device &peer) { ur_device_handle_t Device = impl->getHandleRef(); ur_device_handle_t Peer = peer.impl->getHandleRef(); 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."); detail::adapter_impl &Adapter = impl->getAdapter(); Adapter.call(Device, Peer); } @@ -255,9 +258,14 @@ bool device::ext_oneapi_can_access_peer(const device &peer, }(); detail::adapter_impl &Adapter = impl->getAdapter(); int value = 0; - Adapter.call( - Device, Peer, UrAttr, sizeof(int), &value, nullptr); + auto Err = + Adapter.call_nocheck( + Device, Peer, UrAttr, sizeof(int), &value, nullptr); + // If the backend doesn't support P2P access, neither does its devices. + if (Err == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) + return false; + checkUrResult(Err); return value == 1; } diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 1e26cf0b8a23e..c9080177b1cda 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -124,6 +124,7 @@ inline namespace _V1 { #define SYCL_KHR_DEFAULT_CONTEXT 1 #define SYCL_EXT_INTEL_EVENT_MODE 1 #define SYCL_EXT_ONEAPI_TANGLE 1 +#define SYCL_EXT_ONEAPI_PEER_ACCESS 1 // Unfinished KHR extensions. These extensions are only available if the // __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined. diff --git a/sycl/test-e2e/USM/P2P/p2p_access.cpp b/sycl/test-e2e/USM/P2P/p2p_access.cpp index b86e687428895..deae6b9be74af 100644 --- a/sycl/test-e2e/USM/P2P/p2p_access.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_access.cpp @@ -1,4 +1,3 @@ -// REQUIRES: cuda || hip || level_zero // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp index c36b331ca46c4..e6e64720452ad 100644 --- a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp @@ -1,4 +1,3 @@ -// REQUIRES: cuda || hip || level_zero // RUN: %{build} %if target-nvidia %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_61 %} -o %t.out // RUN: %{run} %t.out From f2d1f8f92b6c8b7659efe2f47bbb9b6091795809 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Aug 2025 06:08:15 -0700 Subject: [PATCH 2/5] fix build error Signed-off-by: Larsen, Steffen --- sycl/source/device.cpp | 2 +- sycl/test-e2e/USM/P2P/p2p_copy.cpp | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 4fb6cd6a66b3a..22af2d3ba9050 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -265,7 +265,7 @@ bool device::ext_oneapi_can_access_peer(const device &peer, // If the backend doesn't support P2P access, neither does its devices. if (Err == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) return false; - checkUrResult(Err); + Adapter.checkUrResult(Err); return value == 1; } diff --git a/sycl/test-e2e/USM/P2P/p2p_copy.cpp b/sycl/test-e2e/USM/P2P/p2p_copy.cpp index 1f4d2733a055c..3f1ae89f7211c 100644 --- a/sycl/test-e2e/USM/P2P/p2p_copy.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_copy.cpp @@ -1,4 +1,3 @@ -// REQUIRES: cuda || hip || level_zero // RUN: %{build} -o %t.out // RUN: %{run} %t.out From 3f2e1835b79de527e836bbc3dafffb16d5a7deba Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Aug 2025 07:03:40 -0700 Subject: [PATCH 3/5] Stop requiring GPU devices Signed-off-by: Larsen, Steffen --- sycl/test-e2e/USM/P2P/p2p_access.cpp | 3 +-- sycl/test-e2e/USM/P2P/p2p_atomics.cpp | 3 +-- sycl/test-e2e/USM/P2P/p2p_copy.cpp | 3 +-- 3 files changed, 3 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/USM/P2P/p2p_access.cpp b/sycl/test-e2e/USM/P2P/p2p_access.cpp index deae6b9be74af..69d4a8b09ff93 100644 --- a/sycl/test-e2e/USM/P2P/p2p_access.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_access.cpp @@ -9,8 +9,7 @@ using namespace sycl; int main() { - - auto Devs = platform(gpu_selector_v).get_devices(info::device_type::gpu); + auto Devs = platform().get_devices(); if (Devs.size() < 2) { std::cout << "Cannot test P2P capabilities, at least two devices are " diff --git a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp index e6e64720452ad..54fa069a313a1 100644 --- a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp @@ -17,8 +17,7 @@ using namespace sycl; constexpr size_t N = 512; int main() { - - auto Devs = platform(gpu_selector_v).get_devices(info::device_type::gpu); + auto Devs = platform().get_devices(); if (Devs.size() < 2) { std::cout << "Cannot test P2P capabilities, at least two devices are " diff --git a/sycl/test-e2e/USM/P2P/p2p_copy.cpp b/sycl/test-e2e/USM/P2P/p2p_copy.cpp index 3f1ae89f7211c..7a60a31db4658 100644 --- a/sycl/test-e2e/USM/P2P/p2p_copy.cpp +++ b/sycl/test-e2e/USM/P2P/p2p_copy.cpp @@ -14,8 +14,7 @@ using namespace sycl; constexpr int N = 100; int main() { - - auto Devs = platform(gpu_selector_v).get_devices(info::device_type::gpu); + auto Devs = platform().get_devices(); if (Devs.size() < 2) { std::cout << "Cannot test P2P capabilities, at least two devices are " From f13dc2eacb058040c6a37f98a405bdf77888046d Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Aug 2025 07:45:22 -0700 Subject: [PATCH 4/5] Fix errc Signed-off-by: Larsen, Steffen --- sycl/source/device.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 22af2d3ba9050..c9735657df086 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -221,7 +221,7 @@ void device::ext_oneapi_enable_peer_access(const device &peer) { ur_device_handle_t Peer = peer.impl->getHandleRef(); if (Device != Peer) { if (!ext_oneapi_can_access_peer(peer)) - throw sycl::exception(make_error_code(errc::invalid), + throw sycl::exception(make_error_code(errc::feature_not_supported), "Peer access is not allowed between the devices."); detail::adapter_impl &Adapter = impl->getAdapter(); Adapter.call(Device, Peer); From 21c380fd0e46d114b8289d6f6a1f9e01e7ce2ef7 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Aug 2025 20:36:31 -0700 Subject: [PATCH 5/5] Remove die-call from OpenCL P2P Signed-off-by: Larsen, Steffen --- unified-runtime/source/adapters/opencl/usm_p2p.cpp | 6 ------ 1 file changed, 6 deletions(-) diff --git a/unified-runtime/source/adapters/opencl/usm_p2p.cpp b/unified-runtime/source/adapters/opencl/usm_p2p.cpp index 8a264bc11c8bd..9ee9920ad87d3 100644 --- a/unified-runtime/source/adapters/opencl/usm_p2p.cpp +++ b/unified-runtime/source/adapters/opencl/usm_p2p.cpp @@ -13,16 +13,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PEnablePeerAccessExp([[maybe_unused]] ur_device_handle_t commandDevice, [[maybe_unused]] ur_device_handle_t peerDevice) { - - die("Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PDisablePeerAccessExp([[maybe_unused]] ur_device_handle_t commandDevice, [[maybe_unused]] ur_device_handle_t peerDevice) { - - die("Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -32,7 +28,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PPeerAccessGetInfoExp( [[maybe_unused]] ur_exp_peer_info_t propName, [[maybe_unused]] size_t propSize, [[maybe_unused]] void *pPropValue, [[maybe_unused]] size_t *pPropSizeRet) { - - die("Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; }