[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:
- 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.
- 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.
- v2/usm.cpp: pass the list of enabled peer devices to the USM
  pool/provider so that zeMemMakeResident is called only for those
  devices.
- 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.
- ur_pool_manager.hpp: propagate resident-devices information through
  pool descriptors.

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: lslusarczyk

.github/copilot-instructions.md

@@ -1,4 +1,164 @@
-When performing a code review, pay close attention to code modifying a function's
-control flow. Could the change result in the corruption of performance profile
-data? Could the change result in invalid debug information, in particular for
-branches and calls?
+# Intel LLVM (DPC++ / oneAPI) — Copilot Instructions
+
+## Repository overview
+
+This is Intel's LLVM fork (`intel/llvm`, `sycl` branch), hosting:
+- **oneAPI DPC++ compiler** — a SYCL implementation built on top of Clang/LLVM
+- **DPC++ runtime library** (`sycl/`)
+- **Unified Runtime (UR)** (`unified-runtime/`) — backend abstraction over Level-Zero, OpenCL, CUDA, HIP
+- **LLVM-SPIRV translator** (`llvm-spirv/`)
+- **SYCL JIT** (`sycl-jit/`), instrumentation (`xpti/`, `xptifw/`), device libraries (`libdevice/`)
+
+For non-SYCL/DPC++ changes, submit patches upstream to https://llvm.org/ directly.
+
+## Build
+
+Use the buildbot scripts (default output directory: `./build`):
+
+```bash
+python buildbot/configure.py   # Run CMake; use --help for full flag list
+python buildbot/compile.py     # Build; default target is deploy-sycl-toolchain
+```
+
+Common `configure.py` flags: `--cuda`, `--hip`, `--native_cpu`, `--shared-libs`, `-t Debug`, `-o <build_dir>`, `--werror`  
+Common `compile.py` flags: `-o <build_dir>`, `-t <target>` (e.g. `clang`, `llvm-spirv`), `-j <N>`
+
+After build, add to environment:
+```bash
+export PATH=$DPCPP_HOME/llvm/build/bin:$PATH
+export LD_LIBRARY_PATH=$DPCPP_HOME/llvm/build/lib:$LD_LIBRARY_PATH
+```
+
+## Testing
+
+### In-tree LIT tests (device-independent)
+
+Run the full check via the buildbot helper:
+```bash
+python buildbot/check.py
+```
+
+Or run individual check targets from the build directory:
+```bash
+ninja check-sycl        # DPC++ headers + runtime unit tests
+ninja check-clang       # Clang (including SYCL FE tests)
+ninja check-llvm        # Core LLVM
+ninja check-llvm-spirv  # LLVM-SPIRV translator
+```
+
+Run a single LIT test file:
+```bash
+llvm-lit sycl/test/path/to/test.cpp
+llvm-lit clang/test/path/to/test.cpp
+```
+
+### E2E tests (require hardware / backend runtimes)
+
+```bash
+# Run all E2E tests on all available devices
+llvm-lit sycl/test-e2e
+
+# Limit to specific devices
+llvm-lit --param sycl_devices="level_zero:gpu;opencl:cpu" sycl/test-e2e
+
+# Run a single E2E test
+llvm-lit sycl/test-e2e/Basic/vector_add.cpp
+```
+
+### Unified Runtime tests (from UR build directory)
+
+```bash
+make check-unified-runtime              # All UR tests
+python -m lit test/conformance/device/ -a --filter="urDeviceGetInfo"  # Single UR conformance test
+```
+
+## Code formatting
+
+Use `git-clang-format` against the `sycl` branch merge base. Match the CI `clang-format` version (see `.github/workflows/pr-code-format.yml`):
+
+```bash
+# Check format of staged changes
+./clang/tools/clang-format/git-clang-format `git merge-base origin/sycl HEAD`
+
+# Also fix unstaged changes
+./clang/tools/clang-format/git-clang-format -f `git merge-base origin/sycl HEAD`
+```
+
+## Architecture
+
+### DPC++ compilation pipeline
+
+1. **Clang front-end** (`clang/`) — parses SYCL source, outlines device code, generates integration headers, enforces device-code restrictions (no exceptions, no virtual calls)
+2. **Device middle-end / back-end** — applies LLVM passes (address space inference, etc.), translates to SPIR-V via `llvm-spirv`, or targets NVPTX/AMDGCPU directly
+3. **Host compilation** — standard Clang compilation of host code
+4. **Linking** — `clang-offload-wrapper` embeds device images into the fat binary
+5. **Runtime** (`sycl/source/`) — manages command queues, buffers, kernels; dispatches work to backend via Unified Runtime
+6. **Unified Runtime** (`unified-runtime/`) — adapters translate UR calls to Level-Zero, OpenCL, CUDA, or HIP
+
+### Key directories
+
+| Path | Contents |
+|---|---|
+| `clang/lib/Sema/`, `clang/lib/CodeGen/` | SYCL-specific FE logic |
+| `sycl/include/sycl/` | Public SYCL headers |
+| `sycl/source/` | DPC++ runtime library implementation |
+| `sycl/test/` | Device-independent LIT + unit tests |
+| `sycl/test-e2e/` | End-to-end tests (need real hardware) |
+| `sycl/unittests/` | GoogleTest unit tests (use UrMock, no device required) |
+| `unified-runtime/source/adapters/` | Level-Zero, OpenCL, CUDA, HIP adapters |
+| `unified-runtime/scripts/core/*.yml` | UR API spec (source of truth — run `make generate` after edits) |
+| `llvm-spirv/` | LLVM IR ↔ SPIR-V translator |
+| `sycl-jit/` | JIT/runtime compilation support |
+| `xpti/`, `xptifw/` | Instrumentation (XPTI tracing) |
+
+### Unified Runtime code generation
+
+UR API is defined in YAML (`unified-runtime/scripts/core/*.yml`). After any YAML changes:
+```bash
+# From unified-runtime build directory
+make generate
+```
+Do **not** hand-edit generated files (`ur_api.h`, `ur_ddi.h`, `ur_libapi.cpp`, `ur_ldrddi.cpp`, `ur_valddi.cpp`, etc.).
+
+## Key conventions
+
+### Commit messages & PRs
+
+- SYCL/DPC++ changes must include `[SYCL]` in the title. Add sub-tags for components: `[UR]`, `[CUDA]`, `[HIP]`, `[Doc]`, `[NFC]`.
+- Commits are squash-merged; the PR description becomes the commit message — keep it accurate.
+- Gatekeepers: `@intel/llvm-gatekeepers`.
+
+### ABI/API stability
+
+- Public DPC++ APIs must remain ABI/API compatible. Experimental APIs are exempted (placed in experimental namespaces).
+- Deprecation order: implement new API → add deprecation warning → wait → remove from headers → remove from library.
+- See `sycl/doc/developer/ABIPolicyGuide.md`.
+
+### Test placement
+
+| Test type | Location | When to use |
+|---|---|---|
+| Clang FE / compile-time checks | `clang/test/ASTSYCL/`, `sycl/test/` | Checking diagnostics, IR, `static_assert` |
+| Runtime unit tests (no device) | `sycl/unittests/` | Runtime behavior, use `UrMock` |
+| Device-independent LIT | `sycl/test/` | Header/IR checks, syntax-only |
+| E2E (requires backend) | `sycl/test-e2e/` | Full stack, actual device execution |
+
+- Use `-fsyntax-only` when only checking compilation succeeds.
+- Device code checks go under `sycl/test/check_device_code/`.
+- Use `sycl::` namespace, not `cl::sycl::`.
+- Always redirect compiler output to `%t` (temp) to avoid race conditions in parallel lit runs.
+
+### SYCL-specific C++ rules
+
+- Never `#include <iostream>` in library/header files. Use `sycl/detail/iostream_proxy.hpp` instead (avoids static constructors).
+- All identifiers in `llvm/sycl` headers must contain at least one lowercase letter (avoids macro conflicts).
+- Target C++17.
+
+### Release notes
+
+Record changes in `sycl/ReleaseNotes.md` under "upcoming release" when changing public interfaces, fixing bugs, or making observable behavioral changes. Reference PRs as `intel/llvm#NNNNN`.
+
+### Code review focus
+
+- When reviewing changes to control flow, check whether the change could corrupt performance profile data or produce invalid debug information (especially for branches and calls).
+- For LLDB code specifically, see `.github/instructions/lldb.instructions.md` for naming conventions (`snake_case` variables, `UpperCamelCase` methods, `m_`/`s_`/`g_` prefixes) and other LLDB-specific rules.

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,
+                    "Can not 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,
+                    "Can not 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;
 
+  // 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 identifer 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,44 @@ 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_THROWS(
+          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_THROWS(
+          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",
+         *commandDevice, *peerDevice);
   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",
+         *commandDevice, *peerDevice);
   return UR_RESULT_SUCCESS;
 }