[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 <[email protected]>

Author: steffenlarsen

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.
 
 |===
 

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<detail::UrApiKind::urUsmP2PEnablePeerAccessExp>(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<detail::UrApiKind::urUsmP2PPeerAccessGetInfoExp>(
-      Device, Peer, UrAttr, sizeof(int), &value, nullptr);
+  auto Err =
+      Adapter.call_nocheck<detail::UrApiKind::urUsmP2PPeerAccessGetInfoExp>(
+          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<errc>(Err);
   return value == 1;
 }
 

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.

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
 

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