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..c9735657df086 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::feature_not_supported), + "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; + Adapter.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..69d4a8b09ff93 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 @@ -10,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 c36b331ca46c4..54fa069a313a1 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 @@ -18,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 1f4d2733a055c..7a60a31db4658 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 @@ -15,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 " 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; }