[UR][L0] Memory resident limit to enabled peers

Previously, USM memory was made resident on all peer devices regardless
of whether peer access had been explicitly enabled by the user. This
change limits memory residency to only those peer devices for which
peer access has been enabled via urUsmP2PEnablePeerAccessExp.

To support this, a per-device `peers` vector is introduced
(ur_device_handle_t_::PeerStatus) that tracks the peer-access state
(ENABLED, DISABLED, NO_CONNECTION) for every device on the platform.
The vector is populated at platform initialisation time using
zeDeviceGetP2PProperties and updated atomically when
urUsmP2PEnablePeerAccessExp / urUsmP2PDisablePeerAccessExp are called.

Key changes:
- device.hpp / device.cpp: introduce the PeerStatus enum
  (ENABLED, DISABLED, NO_CONNECTION) and peers vector on
  ur_device_handle_t_; add stream operators for logging device id and
  PeerStatus values.
- platform.cpp: populate `peers` for every device during device-cache
  initialisation.
- v2/context.cpp: add changeResidentDevice(),
  getDevicesWhoseAllocationsCanBeAccessedFrom(), and
  getDevicesWhichCanAccessAllocationsPresentOn() to update memory
  residency in existing USM pools when peer-access state changes.
  Fix urContextRelease to only remove the context from
  Platform->Contexts and delete it when the reference count actually
  reaches zero, preventing live contexts from being missed during
  subsequent peer-access updates.
  Replace UR_ASSERT with a conditional erase so that contexts created
  via urContextCreateWithNativeHandle (which are not tracked in
  Platform->Contexts) are released without error.
  Fix changeResidentDevice to hold a shared lock on the context Mutex
  while iterating usmPoolHandles, preventing a data race with
  concurrent addUsmPool/removeUsmPool calls.
  Fix a data race in getDevicesWhichCanAccessAllocationsPresentOn where
  peerCandidateDevice->peers.size() was read in UR_FASSERT before the
  scoped_lock was acquired; the lock is now moved before the assert,
  consistent with getDevicesWhoseAllocationsCanBeAccessedFrom.
- v2/usm.cpp: build the resident-devices list from the source device
  plus only the explicitly enabled peer devices (via
  getDevicesWhichCanAccessAllocationsPresentOn), instead of passing
  all platform devices to umfLevelZeroMemoryProviderParamsSetResidentDevices.
  The previous code computed the enabled-peer set but discarded it,
  making memory resident on all devices unconditionally.
- v2/usm_p2p.cpp: new file implementing urUsmP2PEnablePeerAccessExp,
  urUsmP2PDisablePeerAccessExp, and urUsmP2PPeerAccessGetInfoExp for
  the v2 adapter. Fix a data race where peers[] was written under a
  shared (read) lock; changed to an exclusive scoped_lock.
  Copy Platform->Contexts under ContextsMutex and iterate outside the
  critical section to avoid holding the mutex during heavy
  changeResidentDevice calls, reducing deadlock risk.
- CMakeLists.txt: move usm_p2p.cpp from the v1 build into v2/usm_p2p.cpp
  in the v2 build so the v2 adapter uses its own P2P implementation.
- usm_p2p.cpp (v1): add diagnostic log messages to
  urUsmP2PEnablePeerAccessExp and urUsmP2PDisablePeerAccessExp
  clarifying that the operations are no-ops in v1 because P2P is
  always enabled there.
- sycl/source/device.cpp: add a cross-platform check in
  ext_oneapi_enable_peer_access / ext_oneapi_disable_peer_access and
  throw sycl::exception(errc::invalid) when the two devices belong to
  different platforms. Fix exception messages: "Can not" -> "Cannot".
- ur_pool_manager.hpp: propagate resident-devices information through
  pool descriptors.
- backtrace.hpp / backtrace_lin.cpp: replace #define MAX_BACKTRACE_FRAMES
  in backtrace.hpp with a static constexpr int; remove the now-redundant
  redeclaration from backtrace_lin.cpp. This avoids macro leakage and
  improves type safety.
- test/memory_residency.cpp: fix memory leak in
  allocationExistsOnPeerWithEnabledAccess by freeing the USM allocation.

Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
Co-authored-by: Lukasz Dorau <lukasz.dorau@intel.com>

Signed-off-by: Lukasz Dorau <lukasz.dorau@intel.com>

Author: ldorau

sycl/source/device.cpp

@@ -212,19 +212,33 @@ bool device::has(aspect Aspect) const { return impl->has(Aspect); }
 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) {
-    detail::adapter_impl &Adapter = impl->getAdapter();
-    Adapter.call<detail::UrApiKind::urUsmP2PEnablePeerAccessExp>(Device, Peer);
+
+  if (Device == Peer)
+    return;
+
+  if (peer.get_platform() != get_platform()) {
+    throw exception(errc::invalid,
+                    "Cannot enable peer access between different platforms");
   }
+
+  impl->getAdapter().call<detail::UrApiKind::urUsmP2PEnablePeerAccessExp>(
+      Device, Peer);
 }
 
 void device::ext_oneapi_disable_peer_access(const device &peer) {
   ur_device_handle_t Device = impl->getHandleRef();
   ur_device_handle_t Peer = peer.impl->getHandleRef();
-  if (Device != Peer) {
-    detail::adapter_impl &Adapter = impl->getAdapter();
-    Adapter.call<detail::UrApiKind::urUsmP2PDisablePeerAccessExp>(Device, Peer);
+
+  if (Device == Peer)
+    return;
+
+  if (peer.get_platform() != get_platform()) {
+    throw exception(errc::invalid,
+                    "Cannot disable peer access between different platforms");
   }
+
+  impl->getAdapter().call<detail::UrApiKind::urUsmP2PDisablePeerAccessExp>(
+      Device, Peer);
 }
 
 bool device::ext_oneapi_can_access_peer(const device &peer,

unified-runtime/source/adapters/level_zero/CMakeLists.txt

@@ -150,7 +150,6 @@ if(UR_BUILD_ADAPTER_L0_V2)
         ${CMAKE_CURRENT_SOURCE_DIR}/helpers/kernel_helpers.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/helpers/memory_helpers.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/helpers/mutable_helpers.cpp
-        ${CMAKE_CURRENT_SOURCE_DIR}/usm_p2p.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/virtual_mem.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/../../ur/ur.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/sampler.hpp
@@ -194,6 +193,7 @@ if(UR_BUILD_ADAPTER_L0_V2)
         ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_immediate_in_order.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_immediate_out_of_order.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/v2/usm.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/v2/usm_p2p.cpp
     )
     install_ur_library(ur_adapter_level_zero_v2)
 

unified-runtime/source/adapters/level_zero/context.cpp

@@ -41,6 +41,8 @@ ur_result_t urContextCreate(
 
     Context->initialize();
     *RetContext = reinterpret_cast<ur_context_handle_t>(Context);
+    // TODO: delete below 'if' when memory isolation in the context is
+    // implemented in the driver
     if (IndirectAccessTrackingEnabled) {
       std::scoped_lock<ur_shared_mutex> Lock(Platform->ContextsMutex);
       Platform->Contexts.push_back(*RetContext);

unified-runtime/source/adapters/level_zero/device.cpp

@@ -2353,3 +2353,24 @@ void ZeUSMImportExtension::doZeUSMRelease(ze_driver_handle_t DriverHandle,
                                           void *HostPtr) {
   ZE_CALL_NOCHECK(zexDriverReleaseImportedPointer, (DriverHandle, HostPtr));
 }
+
+std::ostream &operator<<(std::ostream &os,
+                         ur_device_handle_t_ const &device_handle) {
+  if (device_handle.Id.has_value()) {
+    return os << device_handle.Id.value();
+  }
+  return os << "NONE";
+}
+
+std::ostream &operator<<(std::ostream &os,
+                         ur_device_handle_t_::PeerStatus peer_status) {
+  switch (peer_status) {
+  case ur_device_handle_t_::PeerStatus::DISABLED:
+    return os << "DISABLED";
+  case ur_device_handle_t_::PeerStatus::ENABLED:
+    return os << "ENABLED";
+  case ur_device_handle_t_::PeerStatus::NO_CONNECTION:
+    return os << "NO_CONNECTION";
+  }
+  return os << "UNKNOWN";
+}

unified-runtime/source/adapters/level_zero/device.hpp

@@ -261,17 +261,29 @@ struct ur_device_handle_t_ : ur_object {
   std::unordered_map<ur_exp_image_native_handle_t, ze_image_handle_t>
       ZeOffsetToImageHandleMap;
 
-  // unique ephemeral identifer of the device in the adapter
+  // Devices which user enabled p2p access by
+  // urUsmP2P(Enable|Disable)PeerAccessExp. Devices are indexed by device id.
+  enum class PeerStatus : char { ENABLED, DISABLED, NO_CONNECTION };
+  std::vector<PeerStatus>
+      peers; // info if our device can access given peer device allocations
+
+  // unique ephemeral identifier of the device in the adapter
   std::optional<DeviceId> Id;
 
   ur::RefCount RefCount;
 };
 
+std::ostream &operator<<(std::ostream &os,
+                         ur_device_handle_t_ const &device_handle);
+std::ostream &operator<<(std::ostream &os,
+                         ur_device_handle_t_::PeerStatus peer_status);
+
 // Collects a flat vector of unique devices for USM memory pool creation.
 // Traverses the input devices and their sub-devices, ensuring each Level Zero
 // device handle appears only once in the result.
 inline std::vector<ur_device_handle_t> CollectDevicesForUsmPoolCreation(
     const std::vector<ur_device_handle_t> &Devices) {
+
   std::vector<ur_device_handle_t> DevicesAndSubDevices;
   std::unordered_set<ze_device_handle_t> Seen;
 

unified-runtime/source/adapters/level_zero/platform.cpp

@@ -746,9 +746,9 @@ ur_platform_handle_t_::getDeviceFromNativeHandle(ze_device_handle_t ZeDevice) {
   std::shared_lock<ur_shared_mutex> Lock(URDevicesCacheMutex);
   auto it = std::find_if(URDevicesCache.begin(), URDevicesCache.end(),
                          [&](std::unique_ptr<ur_device_handle_t_> &D) {
-                           return D.get()->ZeDevice == ZeDevice &&
-                                  (D.get()->RootDevice == nullptr ||
-                                   D.get()->RootDevice->RootDevice == nullptr);
+                           return D->ZeDevice == ZeDevice &&
+                                  (D->RootDevice == nullptr ||
+                                   D->RootDevice->RootDevice == nullptr);
                          });
   if (it != URDevicesCache.end()) {
     return (*it).get();
@@ -914,6 +914,43 @@ ur_result_t ur_platform_handle_t_::populateDeviceCacheIfNeeded() {
     ZeDeviceSynchronizeSupported = Supported;
   }
 
+  for (auto &dev : URDevicesCache) {
+    dev->peers = std::vector<ur_device_handle_t_::PeerStatus>(
+        URDevicesCache.size(), ur_device_handle_t_::PeerStatus::NO_CONNECTION);
+
+    for (size_t peerId = 0; peerId < URDevicesCache.size(); ++peerId) {
+      if (peerId == dev->Id.value())
+        continue;
+
+      ZeStruct<ze_device_p2p_properties_t> p2pProperties;
+      ZE2UR_CALL(zeDeviceGetP2PProperties,
+                 (dev->ZeDevice, URDevicesCache[peerId]->ZeDevice,
+                  &p2pProperties));
+      if (!(p2pProperties.flags & ZE_DEVICE_P2P_PROPERTY_FLAG_ACCESS)) {
+        UR_LOG(INFO,
+               "p2p access to memory of dev:{} from dev:{} not possible due to "
+               "lack of p2p property",
+               peerId, dev->Id.value());
+        continue;
+      }
+
+      ze_bool_t p2p;
+      ZE2UR_CALL(zeDeviceCanAccessPeer,
+                 (dev->ZeDevice, URDevicesCache[peerId]->ZeDevice, &p2p));
+      if (!p2p) {
+        UR_LOG(INFO,
+               "p2p access to memory of dev:{} from dev:{} not possible due to "
+               "no connection",
+               peerId, dev->Id.value());
+        continue;
+      }
+
+      UR_LOG(INFO, "p2p access to memory of dev:{} from dev:{} can be enabled",
+             peerId, dev->Id.value());
+      dev->peers[peerId] = ur_device_handle_t_::PeerStatus::DISABLED;
+    }
+  }
+
   return UR_RESULT_SUCCESS;
 }
 

unified-runtime/source/adapters/level_zero/platform.hpp

@@ -97,11 +97,10 @@ struct ur_platform_handle_t_ : ur::handle_base<ur::level_zero::ddi_getter>,
                                      uint32_t VersionMinor,
                                      uint32_t VersionBuild);
 
-  // Keep track of all contexts in the platform. This is needed to manage
-  // a lifetime of memory allocations in each context when there are kernels
-  // with indirect access.
-  // TODO: should be deleted when memory isolation in the context is implemented
-  // in the driver.
+  // Keep track of all contexts in the platform. In v1 L0 this is needed to
+  // manage a lifetime of memory allocations in each context when there are
+  // kernels with indirect access. In v2 it is used during
+  // ext_oneapi_enable_peer_access and ext_oneapi_disable_peer_access calls.
   std::list<ur_context_handle_t> Contexts;
   ur_shared_mutex ContextsMutex;
 

unified-runtime/source/adapters/level_zero/usm_p2p.cpp

@@ -12,17 +12,23 @@
 
 namespace ur::level_zero {
 
-ur_result_t urUsmP2PEnablePeerAccessExp(ur_device_handle_t /*commandDevice*/,
-                                        ur_device_handle_t /*peerDevice*/) {
+ur_result_t urUsmP2PEnablePeerAccessExp(ur_device_handle_t commandDevice,
+                                        ur_device_handle_t peerDevice) {
 
-  // L0 has peer devices enabled by default
+  UR_LOG(INFO,
+         "user enables peer access to memory of {} from {}, ignored, in V1 P2P "
+         "is always enabled",
+         *peerDevice, *commandDevice);
   return UR_RESULT_SUCCESS;
 }
 
-ur_result_t urUsmP2PDisablePeerAccessExp(ur_device_handle_t /*commandDevice*/,
-                                         ur_device_handle_t /*peerDevice*/) {
+ur_result_t urUsmP2PDisablePeerAccessExp(ur_device_handle_t commandDevice,
+                                         ur_device_handle_t peerDevice) {
 
-  // L0 has peer devices enabled by default
+  UR_LOG(INFO,
+         "user disables peer access to memory of {} from {}, ignored, in V1 "
+         "P2P is always enabled",
+         *peerDevice, *commandDevice);
   return UR_RESULT_SUCCESS;
 }