From 14960ce3bfa0e4bd3ffc95a3e6c06c8b7ccb7e30 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Tue, 6 Oct 2026 14:09:11 +0000
Subject: [PATCH 01/12] [UR][L0] Disable USM pooling on Xe2+ devices with L0
 driver 1.17+

Use a proxy pool instead of a disjoint pool in both the L0 v1 and v2
adapters for devices that are Xe2 or newer and have an L0 driver of
at least the minimum version defined by
UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR and
UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR in
common/usm_pooling_disabled.hpp (currently 1.17). Host pools are not
associated with any device, so pooling of host memory is disabled only
if all devices in the context meet this condition.
---
 .../source/adapters/level_zero/CMakeLists.txt |  1 +
 .../adapters/level_zero/common/device.cpp     | 32 +++++++++++++++++++
 .../adapters/level_zero/common/device.hpp     | 14 ++++++++
 .../common/usm_pooling_disabled.hpp           | 13 ++++++++
 .../source/adapters/level_zero/usm.cpp        | 16 +++++++---
 .../source/adapters/level_zero/v2/usm.cpp     |  4 +--
 6 files changed, 73 insertions(+), 7 deletions(-)
 create mode 100644 unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp

diff --git a/unified-runtime/source/adapters/level_zero/CMakeLists.txt b/unified-runtime/source/adapters/level_zero/CMakeLists.txt
index 733d810c5bf7f..58fa8d2705b1a 100644
--- a/unified-runtime/source/adapters/level_zero/CMakeLists.txt
+++ b/unified-runtime/source/adapters/level_zero/CMakeLists.txt
@@ -6,6 +6,7 @@ if(UR_BUILD_ADAPTER_L0 OR UR_BUILD_ADAPTER_L0_V2)
     set(UR_L0_COMMON_SOURCES
         ${CMAKE_CURRENT_SOURCE_DIR}/common/device.hpp
         ${CMAKE_CURRENT_SOURCE_DIR}/common/platform.hpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/common/usm_pooling_disabled.hpp
 
         ${CMAKE_CURRENT_SOURCE_DIR}/common/adapter.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/common/device.cpp
diff --git a/unified-runtime/source/adapters/level_zero/common/device.cpp b/unified-runtime/source/adapters/level_zero/common/device.cpp
index ab1a06c9e06ce..41f688e429203 100644
--- a/unified-runtime/source/adapters/level_zero/common/device.cpp
+++ b/unified-runtime/source/adapters/level_zero/common/device.cpp
@@ -10,7 +10,9 @@
 #include "adapter.hpp"
 #include "logger/ur_logger.hpp"
 #include "platform.hpp"
+#include "ur_pool_manager.hpp"
 #include "ur_util.hpp"
+#include "usm_pooling_disabled.hpp"
 #include <algorithm>
 #include <climits>
 #if defined(__linux__)
@@ -2025,6 +2027,36 @@ ur_device_handle_t_::useImmediateCommandLists() {
   }
 }
 
+bool ur_device_handle_t_::isUsmPoolingDisabled() {
+  return isBMGOrNewer() && Platform->isDriverVersionNewerOrSimilar(
+                               UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR,
+                               UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR, 0);
+}
+
+bool isUsmPoolingDisabled(const usm::pool_descriptor &Desc) {
+  bool Disabled = false;
+  if (Desc.hDevice) {
+    Disabled = common_cast(Desc.hDevice)->isUsmPoolingDisabled();
+  } else {
+    const auto &Devices = common_cast(Desc.hContext)->getDevices();
+    Disabled =
+        !Devices.empty() && std::all_of(Devices.begin(), Devices.end(),
+                                        [](ur_device_handle_t Device) {
+                                          return Device->isUsmPoolingDisabled();
+                                        });
+  }
+
+  if (Disabled) {
+    UR_LOG(INFO,
+           "USM pooling is disabled on Xe2 or newer devices with L0 driver "
+           "{}.{} or newer, desc:{}",
+           UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR,
+           UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR,
+           logger::makeStringFromStreamable(Desc));
+  }
+  return Disabled;
+}
+
 bool ur_device_handle_t_::useRelaxedAllocationLimits() {
   static const bool EnableRelaxedAllocationLimits = [] {
     auto UrRet = ur_getenv("UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS");
diff --git a/unified-runtime/source/adapters/level_zero/common/device.hpp b/unified-runtime/source/adapters/level_zero/common/device.hpp
index 19f38691b6817..2222f8ccd9d1d 100644
--- a/unified-runtime/source/adapters/level_zero/common/device.hpp
+++ b/unified-runtime/source/adapters/level_zero/common/device.hpp
@@ -29,6 +29,10 @@
 #include <ze_api.h>
 #include <zes_api.h>
 
+namespace usm {
+struct pool_descriptor;
+} // namespace usm
+
 namespace ur::level_zero {
 
 typedef size_t DeviceId;
@@ -247,6 +251,11 @@ struct ur_device_handle_t_ : ur::level_zero::ur_object_t {
             ZeDeviceIpVersionExt->ipVersion >= 0x05004000);
   }
 
+  // USM pooling in the adapter is disabled on Xe2 or newer devices with
+  // L0 driver UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer
+  // (see usm_pooling_disabled.hpp).
+  bool isUsmPoolingDisabled();
+
   bool isIntegrated() {
     return (ZeDeviceProperties->flags & ZE_DEVICE_PROPERTY_FLAG_INTEGRATED);
   }
@@ -350,4 +359,9 @@ inline std::vector<ur_device_handle_t> CollectDevicesForUsmPoolCreation(
   return DevicesAndSubDevices;
 }
 
+// Returns true if USM pooling is disabled for the pool described by Desc.
+// Host pools are not associated with any device, so pooling of host memory is
+// disabled only if it is disabled for all devices in the context.
+bool isUsmPoolingDisabled(const usm::pool_descriptor &Desc);
+
 } // namespace ur::level_zero
diff --git a/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp b/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp
new file mode 100644
index 0000000000000..c5c1c8c5851e7
--- /dev/null
+++ b/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp
@@ -0,0 +1,13 @@
+//===--------- usm_pooling_disabled.hpp - Level Zero Adapter -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM
+// Exceptions. See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+// Minimum L0 driver version (major.minor) starting from which USM pooling
+// in the adapter is disabled on Xe2 or newer devices.
+#define UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR 1
+#define UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR 17
diff --git a/unified-runtime/source/adapters/level_zero/usm.cpp b/unified-runtime/source/adapters/level_zero/usm.cpp
index fc255e9048826..bedd615eb6d4a 100644
--- a/unified-runtime/source/adapters/level_zero/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/usm.cpp
@@ -1106,7 +1106,7 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
       common_cast(DevicesAndSubDevices));
   for (auto &Desc : Descriptors) {
     umf::pool_unique_handle_t Pool = nullptr;
-    if (IsProxy) {
+    if (IsProxy || isUsmPoolingDisabled(Desc)) {
       Pool = usm::makeProxyPool(MakeProvider(&Desc));
     } else {
       auto &PoolConfig =
@@ -1174,11 +1174,17 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
   }
 
   for (auto &Desc : Descriptors) {
-    auto &PoolConfig =
-        DisjointPoolConfigs.Configs[DescToDisjointPoolMemType(Desc)];
+    umf::pool_unique_handle_t Pool = nullptr;
+    if (isUsmPoolingDisabled(Desc)) {
+      Pool = usm::makeProxyPool(MakeProvider(&Desc));
+    } else {
+      auto &PoolConfig =
+          DisjointPoolConfigs.Configs[DescToDisjointPoolMemType(Desc)];
+      Pool = usm::makeDisjointPool(MakeProvider(&Desc), PoolConfig);
+    }
 
-    std::unique_ptr<UsmPool> usmPool = std::make_unique<UsmPool>(
-        this, usm::makeDisjointPool(MakeProvider(&Desc), PoolConfig));
+    std::unique_ptr<UsmPool> usmPool =
+        std::make_unique<UsmPool>(this, std::move(Pool));
     auto Ret = umf::umf2urResult(
         umfPoolSetTag(usmPool->UmfPool.get(), usmPool.get(), nullptr));
     if (Ret) {
diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.cpp b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
index ff00971ad3dba..a8c71545f38a9 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
@@ -167,7 +167,7 @@ ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t hContext,
       v2_cast(this), v2_cast(hContext), common_cast(devicesAndSubDevices));
   for (auto &desc : descriptors) {
     std::unique_ptr<UsmPool> usmPool;
-    if (disjointPoolConfigs.has_value()) {
+    if (disjointPoolConfigs.has_value() && !isUsmPoolingDisabled(desc)) {
       auto &poolConfig =
           disjointPoolConfigs.value().Configs[descToDisjoinPoolMemType(desc)];
       auto pool = usm::makeDisjointPool(makeProvider(desc), poolConfig);
@@ -232,7 +232,7 @@ ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t hContext,
 
   for (auto &desc : descriptors) {
     std::unique_ptr<UsmPool> usmPool;
-    if (disjointPoolConfigs.has_value()) {
+    if (disjointPoolConfigs.has_value() && !isUsmPoolingDisabled(desc)) {
       auto &poolConfig =
           disjointPoolConfigs.value().Configs[descToDisjoinPoolMemType(desc)];
       auto pool = usm::makeDisjointPool(makeProvider(desc), poolConfig);

From 840a0cf67b47bf31255c9781b70471916b3de33c Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 07:17:26 +0000
Subject: [PATCH 02/12] [UR][L0v2] Tolerate proxy-pool NOT_SUPPORTED in
 appendUSMFreeExp

When USM pooling is disabled for Xe2+ devices with at least the minimum
L0 driver version (see isUsmPoolingDisabled()), the adapter uses a UMF
proxy pool instead of a disjoint pool. The proxy pool's
malloc_usable_size implementation always returns
UMF_RESULT_ERROR_NOT_SUPPORTED, which appendUSMFreeExp() treated as
fatal, causing urEnqueueUSMFreeExp to fail with UR_RESULT_ERROR_UNKNOWN.
This mirrors the tolerance pattern already present in usm.cpp (v1) and
v2/usm.cpp for the same proxy-pool quirk.

Also relax enqueue_alloc.cpp's SuccessReuse test: the pointer-reuse
assertion doesn't hold when pooling is disabled and a proxy pool (with
no reuse semantics) is used, so skip it for such devices via a new
isUsmPoolingDisabledForDevice() helper that mirrors the adapter-side
check and uses the UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR/MINOR
constants.
---
 .../level_zero/v2/command_list_manager.cpp    |  3 +-
 .../test/adapters/level_zero/CMakeLists.txt   |  3 ++
 .../adapters/level_zero/enqueue_alloc.cpp     | 43 ++++++++++++++++++-
 3 files changed, 47 insertions(+), 2 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index 8c97ade034533..6bfc2b8cd4e76 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -956,7 +956,8 @@ ur_result_t ur_command_list_manager::appendUSMFreeExp(
 
   size_t size = 0;
   umfRet = umfPoolMallocUsableSize(hPool, pMem, &size);
-  if (umfRet != UMF_RESULT_SUCCESS) {
+  if (umfRet != UMF_RESULT_SUCCESS &&
+      umfRet != UMF_RESULT_ERROR_NOT_SUPPORTED) {
     UR_LOG(ERR, "enqueueUSMFreeExp: failed to retrieve usable malloc size");
     return UR_RESULT_ERROR_UNKNOWN;
   }
diff --git a/unified-runtime/test/adapters/level_zero/CMakeLists.txt b/unified-runtime/test/adapters/level_zero/CMakeLists.txt
index 2a28b1deb4498..4b2da90bb8bd7 100644
--- a/unified-runtime/test/adapters/level_zero/CMakeLists.txt
+++ b/unified-runtime/test/adapters/level_zero/CMakeLists.txt
@@ -44,6 +44,9 @@ else()
     add_l0_loader_kernels_test(kernel_create urKernelCreateWithNativeHandle.cpp)
     add_l0_loader_kernels_test(event_create urEventCreateWithNativeHandle.cpp)
     add_conformance_kernels_test(enqueue_alloc enqueue_alloc.cpp)
+    target_include_directories(enqueue_alloc-test PRIVATE
+        ${PROJECT_SOURCE_DIR}/source/adapters/level_zero/common
+    )
     add_conformance_kernels_test(usm_ops_ordering_l0 urEnqueueUSMOperationsOrderingL0.cpp)
 endif()
 
diff --git a/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp b/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
index 46583845da70d..7784be393f75f 100644
--- a/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
+++ b/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
@@ -9,6 +9,7 @@
 #include <thread>
 
 #include "unified-runtime/ur_api.h"
+#include "usm_pooling_disabled.hpp"
 #include "uur/utils.h"
 #include <uur/fixtures.h>
 
@@ -33,6 +34,38 @@ std::ostream &operator<<(std::ostream &os, EnqueueAllocTestParam param) {
   return os;
 }
 
+// Mirrors ur_device_handle_t_::isUsmPoolingDisabled() in
+// adapters/level_zero/common/device.cpp: USM pooling (and therefore the
+// address-reuse optimization it enables) is disabled by the adapter for
+// Xe2-or-newer devices when running on L0 driver
+// UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer. On such
+// configurations the adapter falls back to a UMF proxy pool, which has no
+// reuse semantics, so tests must not assume a freed allocation's address
+// will be reused.
+bool isUsmPoolingDisabledForDevice(ur_device_handle_t device) {
+  uint32_t vendorId = 0;
+  if (urDeviceGetInfo(device, UR_DEVICE_INFO_VENDOR_ID, sizeof(vendorId),
+                      &vendorId, nullptr) != UR_RESULT_SUCCESS) {
+    return false;
+  }
+
+  uint32_t ipVersion = 0;
+  if (urDeviceGetInfo(device, UR_DEVICE_INFO_IP_VERSION, sizeof(ipVersion),
+                      &ipVersion, nullptr) != UR_RESULT_SUCCESS) {
+    return false;
+  }
+  constexpr uint32_t BMGMinIpVersion = 0x05004000;
+  bool isBMGOrNewer = (vendorId == 0x8086) && (ipVersion >= BMGMinIpVersion);
+  if (!isBMGOrNewer) {
+    return false;
+  }
+
+  auto [major, minor, patch] = uur::getDriverVersion(device);
+  return std::make_tuple(major, minor, patch) >=
+         std::make_tuple(UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR,
+                         UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR, 0);
+}
+
 struct urL0EnqueueAllocTest
     : uur::urKernelExecutionTestWithParam<EnqueueAllocTestParam> {
   void SetUp() override {
@@ -235,7 +268,15 @@ TEST_P(urL0EnqueueAllocTest, SuccessReuse) {
   ur_event_handle_t allocEvent2 = nullptr;
   ASSERT_SUCCESS(enqueueUSMAllocFunc(queue, nullptr, sizeof(DATA), nullptr, 1,
                                      &freeEvent, &ptr2, &allocEvent2));
-  ASSERT_EQ(ptr2, ptr); // Memory should be reused from previous allocation.
+  // Memory should normally be reused from the previous allocation. However,
+  // on Xe2-or-newer devices with L0 driver
+  // UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer, the adapter
+  // disables USM pooling and falls back to a UMF proxy pool (see
+  // ur_device_handle_t_::isUsmPoolingDisabled()), which has no reuse
+  // semantics, so the address is not guaranteed to be reused there.
+  if (!isUsmPoolingDisabledForDevice(device)) {
+    ASSERT_EQ(ptr2, ptr);
+  }
   ASSERT_NE(allocEvent2, nullptr);
 
   ASSERT_SUCCESS(urEventRelease(allocEvent));

From 5947ca91a685347518581ad7b5ec017ea90d64df Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 07:50:37 +0000
Subject: [PATCH 03/12] [UR][TEST] Fix use-after-free in
 urGraphPopulatedExpTest::TearDown

TearDown() called urUSMFree(context, deviceMem) before resetData(),
which enqueues a USM fill on deviceMem. This used the pointer after it
had already been freed.

This was masked when USM pooling kept freed memory mapped/valid for
reuse, but is exposed now that pooling is disabled for Xe2+ devices with
at least the minimum L0 driver version (a UMF proxy pool returns memory
immediately on free), causing urEnqueueUSMFill to fail with
UR_RESULT_ERROR_UNKNOWN in TearDown for every test in the
urGraphPopulatedExpTest hierarchy on affected devices (seen in CI as
widespread exp_graph-test failures on Arc B580).

Fix: reset the data while the allocation is still valid, then free it.
---
 unified-runtime/test/conformance/exp_graph/fixtures.h | 9 ++++++++-
 1 file changed, 8 insertions(+), 1 deletion(-)

diff --git a/unified-runtime/test/conformance/exp_graph/fixtures.h b/unified-runtime/test/conformance/exp_graph/fixtures.h
index f8445392e9b52..bec386505be82 100644
--- a/unified-runtime/test/conformance/exp_graph/fixtures.h
+++ b/unified-runtime/test/conformance/exp_graph/fixtures.h
@@ -91,8 +91,15 @@ struct urGraphPopulatedExpTest : urGraphExpTest {
 
   void TearDown() override {
     if (deviceMem) {
-      ASSERT_SUCCESS(urUSMFree(context, deviceMem));
+      // Reset the data while deviceMem is still a valid allocation. Freeing
+      // first and then enqueuing a fill on the freed pointer is a
+      // use-after-free: it could appear to work when USM pooling keeps freed
+      // memory mapped for reuse, but is not actually valid and fails (or
+      // crashes) once pooling is disabled, e.g. on Xe2+ devices with L0
+      // driver UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer
+      // where a UMF proxy pool (no pooling) is used instead.
       resetData();
+      ASSERT_SUCCESS(urUSMFree(context, deviceMem));
     }
 
     UUR_RETURN_ON_FATAL_FAILURE(urGraphExpTest::TearDown());

From 6a826ff6988b3982288f5c0d8c9ee7d9ded86c96 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 08:32:34 +0000
Subject: [PATCH 04/12] [UR][L0] Fix proxy-pool USM usage stats and USM
 leak-check test

Proxy pools (used when USM pooling is disabled, e.g. on Xe2/BMG devices
with at least the minimum L0 driver version) always fail
malloc_usable_size with UMF_RESULT_ERROR_NOT_SUPPORTED and report size
0. Since AllocationStats relies on that call to track allocation sizes,
UR_USM_POOL_INFO_USED_CURRENT_EXP/USED_HIGH_EXP always reported 0 for
proxy-pool-backed allocations, even though memory was actually in use.

Fix by tracking whether a UsmPool is backed by a proxy pool
(IsProxy/isProxy), and for such pools querying the underlying memory
provider's own statistics (umf.provider.by_handle.{}.stats.
allocated_memory / peak_memory) via umfCtlGet, summing them with the
regular AllocationStats-based accounting used for non-proxy pools.
This mirrors the existing approach already used for
getTotalReservedSize()/getPeakReservedSize().

getTotalUsedSize()/getPeakUsedSize() now return ur_result_t (with the
size returned via an output parameter) instead of throwing through
the UR C API boundary, so UMF errors are propagated cleanly to
urUSMPoolGetInfoExp callers instead of risking an uncaught C++
exception crossing an extern "C" entry point.

Also fix USM/usm_leak_check.cpp: direct_usm() only freed the shared
allocation (p1) and relied on pooling to implicitly free the host/
device allocations (p2/p3) at process exit. With USM pooling disabled
on some devices, this now leaks and fails the test. Explicitly free
p2 and p3 as well.
---
 sycl/test-e2e/USM/usm_leak_check.cpp          |  6 +-
 .../source/adapters/level_zero/usm.cpp        | 59 +++++++++++++++----
 .../source/adapters/level_zero/usm.hpp        |  9 ++-
 .../source/adapters/level_zero/v2/usm.cpp     | 59 +++++++++++++++----
 .../source/adapters/level_zero/v2/usm.hpp     | 11 ++--
 5 files changed, 112 insertions(+), 32 deletions(-)

diff --git a/sycl/test-e2e/USM/usm_leak_check.cpp b/sycl/test-e2e/USM/usm_leak_check.cpp
index a859ea6347e05..0c553cf828fd5 100644
--- a/sycl/test-e2e/USM/usm_leak_check.cpp
+++ b/sycl/test-e2e/USM/usm_leak_check.cpp
@@ -22,9 +22,11 @@ void direct_usm(queue &Q) {
   auto p1 = malloc_shared(1024, Q);
   auto p2 = malloc_host(1024, Q);
   auto p3 = malloc_device(1024, Q);
-  // Host and Device allocations, pooled by default will be automatically freed
-  // Shared is not pooled by default, so it needs to be explicitly freed
+  // All allocations must be freed explicitly: pooling is disabled on some
+  // devices, so host and device allocations may not be cached until teardown.
   free(p1, Q.get_context());
+  free(p2, Q.get_context());
+  free(p3, Q.get_context());
 }
 
 template <typename T> class K;
diff --git a/unified-runtime/source/adapters/level_zero/usm.cpp b/unified-runtime/source/adapters/level_zero/usm.cpp
index bedd615eb6d4a..8b30852cdba31 100644
--- a/unified-runtime/source/adapters/level_zero/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/usm.cpp
@@ -659,13 +659,15 @@ ur_result_t UR_APICALL urUSMPoolGetInfoExp(::ur_usm_pool_handle_t hPoolOpque,
     value = hPool->getTotalReservedSize();
     break;
   case UR_USM_POOL_INFO_USED_CURRENT_EXP:
-    value = hPool->getTotalUsedSize();
+    if (auto Ret = hPool->getTotalUsedSize(value))
+      return Ret;
     break;
   case UR_USM_POOL_INFO_RESERVED_HIGH_EXP:
     value = hPool->getPeakReservedSize();
     break;
   case UR_USM_POOL_INFO_USED_HIGH_EXP:
-    value = hPool->getPeakUsedSize();
+    if (auto Ret = hPool->getPeakUsedSize(value))
+      return Ret;
     break;
   default:
     // Unknown enumerator
@@ -1076,8 +1078,9 @@ MakeProvider(ProviderParams *Params = nullptr) {
   return nullptr;
 }
 
-UsmPool::UsmPool(ur_usm_pool_handle_t UrPool, umf::pool_unique_handle_t UmfPool)
-    : UrPool(UrPool), UmfPool(std::move(UmfPool)),
+UsmPool::UsmPool(ur_usm_pool_handle_t UrPool, umf::pool_unique_handle_t UmfPool,
+                 bool IsProxy)
+    : UrPool(UrPool), UmfPool(std::move(UmfPool)), IsProxy(IsProxy),
       AsyncPool(
           [](::ur_event_handle_t Event) {
             return urEventReleaseInternal(v1_cast(Event));
@@ -1106,7 +1109,8 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
       common_cast(DevicesAndSubDevices));
   for (auto &Desc : Descriptors) {
     umf::pool_unique_handle_t Pool = nullptr;
-    if (IsProxy || isUsmPoolingDisabled(Desc)) {
+    bool UseProxy = IsProxy || isUsmPoolingDisabled(Desc);
+    if (UseProxy) {
       Pool = usm::makeProxyPool(MakeProvider(&Desc));
     } else {
       auto &PoolConfig =
@@ -1115,7 +1119,7 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
     }
 
     std::unique_ptr<UsmPool> usmPool =
-        std::make_unique<UsmPool>(this, std::move(Pool));
+        std::make_unique<UsmPool>(this, std::move(Pool), UseProxy);
     auto Ret = umf::umf2urResult(
         umfPoolSetTag(usmPool->UmfPool.get(), usmPool.get(), nullptr));
     if (Ret) {
@@ -1175,7 +1179,8 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
 
   for (auto &Desc : Descriptors) {
     umf::pool_unique_handle_t Pool = nullptr;
-    if (isUsmPoolingDisabled(Desc)) {
+    bool UseProxy = isUsmPoolingDisabled(Desc);
+    if (UseProxy) {
       Pool = usm::makeProxyPool(MakeProvider(&Desc));
     } else {
       auto &PoolConfig =
@@ -1184,7 +1189,7 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context,
     }
 
     std::unique_ptr<UsmPool> usmPool =
-        std::make_unique<UsmPool>(this, std::move(Pool));
+        std::make_unique<UsmPool>(this, std::move(Pool), UseProxy);
     auto Ret = umf::umf2urResult(
         umfPoolSetTag(usmPool->UmfPool.get(), usmPool.get(), nullptr));
     if (Ret) {
@@ -1408,12 +1413,42 @@ size_t v1::ur_usm_pool_handle_t_::getPeakReservedSize() {
   return Ret == UMF_RESULT_SUCCESS ? MaxPeakSize : 0;
 }
 
-size_t v1::ur_usm_pool_handle_t_::getTotalUsedSize() {
-  return AllocStats.getCurrent();
+ur_result_t v1::ur_usm_pool_handle_t_::getProxyUsedSize(bool Peak,
+                                                        size_t &Size) {
+  umf_result_t Ret = UMF_RESULT_SUCCESS;
+  PoolManager.forEachPool([&](UsmPool *Pool) {
+    if (!Pool->IsProxy)
+      return true;
+    umf_memory_provider_handle_t Provider = nullptr;
+    Ret = umfPoolGetMemoryProvider(Pool->UmfPool.get(), &Provider);
+    if (Ret != UMF_RESULT_SUCCESS)
+      return false;
+    size_t ProviderSize = 0;
+    Ret = umfCtlGet(Peak ? "umf.provider.by_handle.{}.stats.peak_memory"
+                         : "umf.provider.by_handle.{}.stats.allocated_memory",
+                    &ProviderSize, sizeof(ProviderSize), Provider);
+    if (Ret != UMF_RESULT_SUCCESS)
+      return false;
+    // Proxy pools do not cache allocations, so provider statistics track
+    // current and peak usage even though malloc_usable_size is unsupported.
+    Size += ProviderSize;
+    return true;
+  });
+  if (Ret != UMF_RESULT_SUCCESS) {
+    UR_LOG(ERR, "Failed to query proxy pool memory usage: {}", Ret);
+    return umf::umf2urResult(Ret);
+  }
+  return UR_RESULT_SUCCESS;
+}
+
+ur_result_t v1::ur_usm_pool_handle_t_::getTotalUsedSize(size_t &UsedSize) {
+  UsedSize = AllocStats.getCurrent();
+  return getProxyUsedSize(false, UsedSize);
 }
 
-size_t v1::ur_usm_pool_handle_t_::getPeakUsedSize() {
-  return AllocStats.getPeak();
+ur_result_t v1::ur_usm_pool_handle_t_::getPeakUsedSize(size_t &PeakSize) {
+  PeakSize = AllocStats.getPeak();
+  return getProxyUsedSize(true, PeakSize);
 }
 
 bool v1::ur_usm_pool_handle_t_::hasPool(const umf_memory_pool_handle_t Pool) {
diff --git a/unified-runtime/source/adapters/level_zero/usm.hpp b/unified-runtime/source/adapters/level_zero/usm.hpp
index f665756a2591a..2e2be2c9108a2 100644
--- a/unified-runtime/source/adapters/level_zero/usm.hpp
+++ b/unified-runtime/source/adapters/level_zero/usm.hpp
@@ -22,10 +22,12 @@ namespace ur::level_zero::v1 {
 extern usm::DisjointPoolAllConfigs DisjointPoolConfigInstance;
 
 struct UsmPool {
-  UsmPool(ur_usm_pool_handle_t UrPool, umf::pool_unique_handle_t UmfPool);
+  UsmPool(ur_usm_pool_handle_t UrPool, umf::pool_unique_handle_t UmfPool,
+          bool IsProxy);
   // Parent pool.
   ur_usm_pool_handle_t UrPool;
   umf::pool_unique_handle_t UmfPool;
+  bool IsProxy;
   // 'AsyncPool' needs to be declared after 'UmfPool' so its destructor is
   // invoked first.
   EnqueuedPool AsyncPool;
@@ -80,14 +82,15 @@ struct ur_usm_pool_handle_t_ : ur_object_t {
   void cleanupPoolsForQueue(ur_queue_handle_t Queue);
   size_t getTotalReservedSize();
   size_t getPeakReservedSize();
-  size_t getTotalUsedSize();
-  size_t getPeakUsedSize();
+  ur_result_t getTotalUsedSize(size_t &UsedSize);
+  ur_result_t getPeakUsedSize(size_t &PeakSize);
   UsmPool *getPool(const usm::pool_descriptor &Desc);
 
   ur_context_handle_t Context;
   ur::RefCount RefCount;
 
 private:
+  ur_result_t getProxyUsedSize(bool Peak, size_t &Size);
   usm::pool_manager<usm::pool_descriptor, UsmPool> PoolManager;
   AllocationStats AllocStats;
 };
diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.cpp b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
index a8c71545f38a9..32c4584133431 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
@@ -167,14 +167,16 @@ ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t hContext,
       v2_cast(this), v2_cast(hContext), common_cast(devicesAndSubDevices));
   for (auto &desc : descriptors) {
     std::unique_ptr<UsmPool> usmPool;
-    if (disjointPoolConfigs.has_value() && !isUsmPoolingDisabled(desc)) {
+    bool isProxy =
+        !disjointPoolConfigs.has_value() || isUsmPoolingDisabled(desc);
+    if (!isProxy) {
       auto &poolConfig =
           disjointPoolConfigs.value().Configs[descToDisjoinPoolMemType(desc)];
       auto pool = usm::makeDisjointPool(makeProvider(desc), poolConfig);
-      usmPool = std::make_unique<UsmPool>(this, std::move(pool));
+      usmPool = std::make_unique<UsmPool>(this, std::move(pool), isProxy);
     } else {
       auto pool = usm::makeProxyPool(makeProvider(desc));
-      usmPool = std::make_unique<UsmPool>(this, std::move(pool));
+      usmPool = std::make_unique<UsmPool>(this, std::move(pool), isProxy);
     }
     UMF_CALL_THROWS(
         umfPoolSetTag(usmPool->umfPool.get(), usmPool.get(), nullptr));
@@ -232,14 +234,16 @@ ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t hContext,
 
   for (auto &desc : descriptors) {
     std::unique_ptr<UsmPool> usmPool;
-    if (disjointPoolConfigs.has_value() && !isUsmPoolingDisabled(desc)) {
+    bool isProxy =
+        !disjointPoolConfigs.has_value() || isUsmPoolingDisabled(desc);
+    if (!isProxy) {
       auto &poolConfig =
           disjointPoolConfigs.value().Configs[descToDisjoinPoolMemType(desc)];
       auto pool = usm::makeDisjointPool(makeProvider(desc), poolConfig);
-      usmPool = std::make_unique<UsmPool>(this, std::move(pool));
+      usmPool = std::make_unique<UsmPool>(this, std::move(pool), isProxy);
     } else {
       auto pool = usm::makeProxyPool(makeProvider(desc));
-      usmPool = std::make_unique<UsmPool>(this, std::move(pool));
+      usmPool = std::make_unique<UsmPool>(this, std::move(pool), isProxy);
     }
     UMF_CALL_THROWS(
         umfPoolSetTag(usmPool->umfPool.get(), usmPool.get(), nullptr));
@@ -437,11 +441,42 @@ size_t ur_usm_pool_handle_t_::getPeakReservedSize() {
   return umfRet == UMF_RESULT_SUCCESS ? maxPeakSize : 0;
 }
 
-size_t ur_usm_pool_handle_t_::getTotalUsedSize() {
-  return allocStats.getCurrent();
+ur_result_t ur_usm_pool_handle_t_::getProxyUsedSize(bool peak, size_t &size) {
+  umf_result_t ret = UMF_RESULT_SUCCESS;
+  poolManager.forEachPool([&](UsmPool *pool) {
+    if (!pool->isProxy)
+      return true;
+    umf_memory_provider_handle_t provider = nullptr;
+    ret = umfPoolGetMemoryProvider(pool->umfPool.get(), &provider);
+    if (ret != UMF_RESULT_SUCCESS)
+      return false;
+    size_t providerSize = 0;
+    ret = umfCtlGet(peak ? "umf.provider.by_handle.{}.stats.peak_memory"
+                         : "umf.provider.by_handle.{}.stats.allocated_memory",
+                    &providerSize, sizeof(providerSize), provider);
+    if (ret != UMF_RESULT_SUCCESS)
+      return false;
+    // Proxy pools do not cache allocations, so provider statistics track
+    // current and peak usage even though malloc_usable_size is unsupported.
+    size += providerSize;
+    return true;
+  });
+  if (ret != UMF_RESULT_SUCCESS) {
+    UR_LOG(ERR, "Failed to query proxy pool memory usage: {}", ret);
+    return umf::umf2urResult(ret);
+  }
+  return UR_RESULT_SUCCESS;
 }
 
-size_t ur_usm_pool_handle_t_::getPeakUsedSize() { return allocStats.getPeak(); }
+ur_result_t ur_usm_pool_handle_t_::getTotalUsedSize(size_t &usedSize) {
+  usedSize = allocStats.getCurrent();
+  return getProxyUsedSize(false, usedSize);
+}
+
+ur_result_t ur_usm_pool_handle_t_::getPeakUsedSize(size_t &peakSize) {
+  peakSize = allocStats.getPeak();
+  return getProxyUsedSize(true, peakSize);
+}
 
 void ur_usm_pool_handle_t_::changeResidentDevice(ur_device_handle_t hDevice,
                                                  ur_device_handle_t peerDevice,
@@ -607,13 +642,15 @@ ur_result_t urUSMPoolGetInfoExp(::ur_usm_pool_handle_t hPoolOpque,
     value = hPool->getTotalReservedSize();
     break;
   case UR_USM_POOL_INFO_USED_CURRENT_EXP:
-    value = hPool->getTotalUsedSize();
+    if (auto ret = hPool->getTotalUsedSize(value))
+      return ret;
     break;
   case UR_USM_POOL_INFO_RESERVED_HIGH_EXP:
     value = hPool->getPeakReservedSize();
     break;
   case UR_USM_POOL_INFO_USED_HIGH_EXP:
-    value = hPool->getPeakUsedSize();
+    if (auto ret = hPool->getPeakUsedSize(value))
+      return ret;
     break;
   default:
     // Unknown enumerator
diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.hpp b/unified-runtime/source/adapters/level_zero/v2/usm.hpp
index 2fb53dbbfafb1..4fc17a4d59479 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.hpp
@@ -79,8 +79,8 @@ struct ur_usm_pool_handle_t_ : v2::ur_object_t {
   void cleanupPoolsForQueue(void *hQueue);
   size_t getTotalReservedSize();
   size_t getPeakReservedSize();
-  size_t getTotalUsedSize();
-  size_t getPeakUsedSize();
+  ur_result_t getTotalUsedSize(size_t &usedSize);
+  ur_result_t getPeakUsedSize(size_t &peakSize);
   void changeResidentDevice(ur_device_handle_t hDevice,
                             ur_device_handle_t peerDevice, bool isAdding);
 
@@ -89,14 +89,16 @@ struct ur_usm_pool_handle_t_ : v2::ur_object_t {
   ur::RefCount RefCount;
 
 private:
+  ur_result_t getProxyUsedSize(bool peak, size_t &size);
   ur_context_handle_t hContext;
   usm::pool_manager<usm::pool_descriptor, UsmPool> poolManager;
   AllocationStats allocStats;
 };
 
 struct UsmPool {
-  UsmPool(ur_usm_pool_handle_t urPool, umf::pool_unique_handle_t umfPool)
-      : urPool(urPool), umfPool(std::move(umfPool)),
+  UsmPool(ur_usm_pool_handle_t urPool, umf::pool_unique_handle_t umfPool,
+          bool isProxy)
+      : urPool(urPool), umfPool(std::move(umfPool)), isProxy(isProxy),
         asyncPool(
             [](::ur_event_handle_t hEvent) {
               return v2_cast(hEvent)->release();
@@ -106,6 +108,7 @@ struct UsmPool {
             }) {}
   ur_usm_pool_handle_t urPool;
   umf::pool_unique_handle_t umfPool;
+  bool isProxy;
   // 'asyncPool' needs to be declared after 'umfPool' so its destructor is
   // invoked first.
   EnqueuedPool asyncPool;

From 36bfab738f23a327d79cab65a48daf0f8ca18235 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 10:43:56 +0000
Subject: [PATCH 05/12] [SYCL][E2E] Add level_zero_usm_pooling_disabled lit
 feature

The Level Zero adapter disables USM pooling on Xe2 or newer Intel GPUs
(IP version >= 20.1.0) with at least the minimum L0 driver version
(UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR/MINOR). Detect this
condition in lit.cfg.py from sycl-ls output and mark the tests which
verify pooling behavior (usm_device_read_only.cpp, usm_pooling.cpp) as
UNSUPPORTED on such devices. The minimum driver version is mirrored in
lit.cfg.py as L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR/MINOR.
---
 .../level_zero/usm_device_read_only.cpp       |  5 +++
 sycl/test-e2e/USM/usm_pooling.cpp             |  5 +++
 sycl/test-e2e/lit.cfg.py                      | 35 +++++++++++++++++++
 3 files changed, 45 insertions(+)

diff --git a/sycl/test-e2e/Adapters/level_zero/usm_device_read_only.cpp b/sycl/test-e2e/Adapters/level_zero/usm_device_read_only.cpp
index c1205071829b7..536d13f84b489 100644
--- a/sycl/test-e2e/Adapters/level_zero/usm_device_read_only.cpp
+++ b/sycl/test-e2e/Adapters/level_zero/usm_device_read_only.cpp
@@ -8,6 +8,11 @@
 // UNSUPPORTED: windows && (gpu-intel-dg2 || arch-intel_gpu_bmg_g21)
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22099
 
+// UNSUPPORTED: level_zero_usm_pooling_disabled
+// UNSUPPORTED-INTENDED: USM pooling is disabled in the Level Zero adapter on
+// Xe2 or newer devices with L0 driver at least
+// UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR.
+
 // RUN: %{build} -o %t.out
 // RUN: env SYCL_UR_TRACE=2 UR_L0_DEBUG=1 %{run} %t.out 2>&1 | FileCheck %s
 
diff --git a/sycl/test-e2e/USM/usm_pooling.cpp b/sycl/test-e2e/USM/usm_pooling.cpp
index 8c9582e3808ab..87fccbf4ed0e4 100644
--- a/sycl/test-e2e/USM/usm_pooling.cpp
+++ b/sycl/test-e2e/USM/usm_pooling.cpp
@@ -8,6 +8,11 @@
 // UNSUPPORTED: windows
 // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/21556
 
+// UNSUPPORTED: level_zero_usm_pooling_disabled
+// UNSUPPORTED-INTENDED: USM pooling is disabled in the Level Zero adapter on
+// Xe2 or newer devices with L0 driver at least
+// UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR.
+
 // Allocate 2 items of 2MB. Free 2. Allocate 3 more of 2MB.
 
 // With no pooling: 1,2,3,4,5 allocs lead to ZE call.
diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py
index c0aef4146238b..858b3012e9289 100644
--- a/sycl/test-e2e/lit.cfg.py
+++ b/sycl/test-e2e/lit.cfg.py
@@ -1105,6 +1105,12 @@ def get_sycl_ls_verbose(sycl_device, env):
 
 # Version of the driver for a given device. Empty for non-Intel devices.
 config.intel_driver_ver = {}
+# Minimum L0 driver version starting from which the Level Zero adapter
+# disables USM pooling on Xe2 or newer Intel GPUs. Must be kept in sync with
+# UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR/MINOR defined in
+# unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp.
+L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR = 1
+L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR = 17
 for full_name, sycl_device in zip(
     config.sycl_devices, remove_level_zero_suffix(config.sycl_devices)
 ):
@@ -1122,14 +1128,35 @@ def get_sycl_ls_verbose(sycl_device, env):
     # See format.py's parse_min_intel_driver_req for explanation.
     is_intel_driver = False
     intel_driver_ver = {}
+    # Per-device flags used to detect whether the Level Zero adapter disables
+    # USM pooling (see isUsmPoolingDisabled() in the UR L0 adapter).
+    usm_pooling_disabled = []
     sycl_ls_sp = get_sycl_ls_verbose(sycl_device, env)
     offload_assigned_backend = ""
     for line in sycl_ls_sp.stdout.splitlines():
         if re.match(r" *Vendor *: Intel\(R\) Corporation", line):
             is_intel_driver = True
+        if re.match(r" *Version *:", line) and sycl_device.startswith("level_zero"):
+            # Device IP version of Intel GPUs, e.g. "20.1.0" for BMG.
+            ip_ver = re.match(r" *Version *: *([0-9]+)\.([0-9]+)\.([0-9]+)", line)
+            if ip_ver:
+                usm_pooling_disabled.append(
+                    is_intel_driver and tuple(map(int, ip_ver.groups())) >= (20, 1, 0)
+                )
         if re.match(r" *Driver *:", line):
             _, driver_str = line.split(":", 1)
             driver_str = driver_str.strip()
+            if usm_pooling_disabled and sycl_device.startswith("level_zero"):
+                l0_ver = re.match(r"([0-9]+)\.([0-9]+)\.", driver_str)
+                usm_pooling_disabled[-1] = bool(
+                    usm_pooling_disabled[-1]
+                    and l0_ver
+                    and tuple(map(int, l0_ver.groups()))
+                    >= (
+                        L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR,
+                        L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR,
+                    )
+                )
             if sycl_device.endswith("cpu"):
                 intel_driver_ver["cpu"] = driver_str
             else:
@@ -1255,6 +1282,14 @@ def get_sycl_ls_verbose(sycl_device, env):
     features.update(device_family)
     features.update(device_name_features)
 
+    # The Level Zero adapter disables USM pooling on Xe2 or newer Intel GPUs
+    # (IP version >= 20.1.0) with L0 driver
+    # L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer, so tests
+    # checking pooling behavior cannot pass there.
+    if usm_pooling_disabled and all(usm_pooling_disabled):
+        features.add("level_zero_usm_pooling_disabled")
+        lit_config.note("USM pooling is disabled for {}".format(sycl_device))
+
     be, dev = sycl_device.split(":")
     if dev.isdigit():
         backend_devices = available_devices.get(be, "gpu")

From 693b2fe95787150175f69d5561af4b69dec22e6c Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 10:57:23 +0000
Subject: [PATCH 06/12] [UR][L0v2] Reject too large host USM allocations

The UMF Level Zero provider always passes the relaxed allocation limits
descriptor to zeMemAllocHost, so some newer L0 drivers accept absurd
sizes like SIZE_MAX. With the disjoint pool this was masked by a failing
umfPoolMallocUsableSize() call, but with the proxy pool (USM pooling
disabled on Xe2+ with at least the minimum L0 driver version) the
allocation succeeds and USM/badmalloc.cpp fails.

Reject host allocations larger than maxMemAllocSize unless relaxed
allocation limits are enabled, matching the V1 adapter behavior.
---
 .../source/adapters/level_zero/v2/usm.cpp      | 18 ++++++++++++++++++
 1 file changed, 18 insertions(+)

diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.cpp b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
index 32c4584133431..c949e981b9b25 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
@@ -285,6 +285,24 @@ ur_result_t ur_usm_pool_handle_t_::allocate(
     return UR_RESULT_ERROR_INVALID_VALUE;
   }
 
+  // The UMF Level Zero provider always passes the relaxed allocation limits
+  // descriptor to zeMemAllocHost, which makes some drivers accept absurd
+  // sizes (e.g. SIZE_MAX). Reject host allocations larger than
+  // maxMemAllocSize unless relaxed allocation limits are enabled, the same
+  // way the V1 adapter does (it calls zeMemAllocHost without that descriptor).
+  if (type == UR_USM_TYPE_HOST) {
+    bool relaxedLimits = false;
+    uint64_t maxAllocSize = 0;
+    for (auto *device : hContext->getDevices()) {
+      relaxedLimits |= device->useRelaxedAllocationLimits();
+      maxAllocSize = std::max<uint64_t>(
+          maxAllocSize, device->ZeDeviceProperties->maxMemAllocSize);
+    }
+    if (!relaxedLimits && maxAllocSize > 0 && size > maxAllocSize) {
+      return UR_RESULT_ERROR_INVALID_USM_SIZE;
+    }
+  }
+
   auto deviceFlags = getDeviceFlags(pUSMDesc);
 
   auto pool = getPool(usm::pool_descriptor{

From 8da3d191b298f3b030553471d679d135872e9ab1 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 12:15:42 +0000
Subject: [PATCH 07/12] [SYCL][E2E] Free USM allocations in Graph tests checked
 for L0 leaks

Many Graph tests run with UR_L0_LEAKS_DEBUG=1 but never free their USM
allocations. This was hidden by the UMF disjoint pool, which releases
all of its slabs on destruction. With USM pooling disabled (Xe2+ with at
least the minimum L0 driver version) the proxy pool forwards every
allocation directly to the driver, so unfreed allocations are reported
as leaks.

Add the missing free() calls.
---
 sycl/test-e2e/Graph/Inputs/raw_kernel_arg.cpp                | 1 +
 sycl/test-e2e/Graph/Inputs/whole_update_raw_arg.cpp          | 3 +++
 sycl/test-e2e/Graph/Inputs/whole_update_subgraph.cpp         | 4 ++++
 sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp   | 5 +++++
 sycl/test-e2e/Graph/RecordReplay/transitive_queue.cpp        | 3 +++
 .../test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp | 3 +++
 sycl/test-e2e/Graph/Update/regression_padded_parameter.cpp   | 1 +
 sycl/test-e2e/Graph/Update/update_before_finalize.cpp        | 2 ++
 sycl/test-e2e/Graph/Update/update_nd_range.cpp               | 2 ++
 sycl/test-e2e/Graph/Update/update_ndrange_to_range.cpp       | 2 ++
 sycl/test-e2e/Graph/Update/update_range.cpp                  | 2 ++
 sycl/test-e2e/Graph/Update/update_range_to_ndrange.cpp       | 2 ++
 .../Update/update_with_indices_multiple_exec_graphs.cpp      | 3 +++
 sycl/test-e2e/Graph/Update/update_with_indices_ordering.cpp  | 3 +++
 sycl/test-e2e/Graph/Update/update_with_indices_ptr.cpp       | 3 +++
 sycl/test-e2e/Graph/Update/update_with_indices_ptr_3D.cpp    | 3 +++
 .../Graph/Update/update_with_indices_ptr_double_update.cpp   | 4 ++++
 .../Graph/Update/update_with_indices_ptr_multiple_nodes.cpp  | 3 +++
 ...ate_with_indices_ptr_multiple_nodes_different_indices.cpp | 3 +++
 .../Graph/Update/update_with_indices_ptr_multiple_params.cpp | 4 ++++
 .../Graph/Update/update_with_indices_ptr_subgraph.cpp        | 3 +++
 sycl/test-e2e/Graph/Update/update_with_indices_scalar.cpp    | 2 ++
 sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg.cpp    | 3 +++
 sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp     | 2 ++
 24 files changed, 66 insertions(+)

diff --git a/sycl/test-e2e/Graph/Inputs/raw_kernel_arg.cpp b/sycl/test-e2e/Graph/Inputs/raw_kernel_arg.cpp
index e2cc930a0da55..37eb07459b3c3 100644
--- a/sycl/test-e2e/Graph/Inputs/raw_kernel_arg.cpp
+++ b/sycl/test-e2e/Graph/Inputs/raw_kernel_arg.cpp
@@ -45,5 +45,6 @@ int main() {
     assert(HostData[i] == (i + Scalar));
   }
 
+  free(Ptr, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Inputs/whole_update_raw_arg.cpp b/sycl/test-e2e/Graph/Inputs/whole_update_raw_arg.cpp
index 5ddc4f8852364..56bf9745d3653 100644
--- a/sycl/test-e2e/Graph/Inputs/whole_update_raw_arg.cpp
+++ b/sycl/test-e2e/Graph/Inputs/whole_update_raw_arg.cpp
@@ -75,5 +75,8 @@ __kernel void RawArgKernel(int scalar, __global int *out) {
     assert(HostDataA[i] == (i + ScalarA));
     assert(HostDataB[i] == (i + ScalarB));
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Inputs/whole_update_subgraph.cpp b/sycl/test-e2e/Graph/Inputs/whole_update_subgraph.cpp
index 00c3b455b15c5..3818b99adbc5f 100644
--- a/sycl/test-e2e/Graph/Inputs/whole_update_subgraph.cpp
+++ b/sycl/test-e2e/Graph/Inputs/whole_update_subgraph.cpp
@@ -54,6 +54,8 @@ int main() {
   } catch (sycl::exception &e) {
     // TODO The subgraph update feature is not implemented yet. For now this
     // is the expected behaviour.
+    free(DataDevice, Queue);
+    free(DataDeviceUpdate, Queue);
     return 0;
   }
   assert(!GotException);
@@ -69,5 +71,7 @@ int main() {
     assert(check_value(i, 4, DataHostUpdate[i], "DataHostUpdate"));
   }
 
+  free(DataDevice, Queue);
+  free(DataDeviceUpdate, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp b/sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp
index b3f413f1d1eec..9bd0f283dd045 100644
--- a/sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp
+++ b/sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp
@@ -127,5 +127,10 @@ int main() {
     assert(HostData[I] == Data[I]);
   }
 
+  free(PtrFrom, Queue);
+  free(PtrTo, Queue);
+  free(PtrA, Queue);
+  free(PtrB, Queue);
+  free(PtrC, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/RecordReplay/transitive_queue.cpp b/sycl/test-e2e/Graph/RecordReplay/transitive_queue.cpp
index 8378c8f1271ed..6a692a078a2f5 100644
--- a/sycl/test-e2e/Graph/RecordReplay/transitive_queue.cpp
+++ b/sycl/test-e2e/Graph/RecordReplay/transitive_queue.cpp
@@ -143,5 +143,8 @@ int main() {
     assert(check_value(i, ReferenceC[i] / 2 + 2, DataC[i], "DataC"));
   }
 
+  free(PtrA, Q1);
+  free(PtrB, Q1);
+  free(PtrC, Q1);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp b/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp
index 36981b12dedd5..89c39b09786ae 100644
--- a/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp
+++ b/sycl/test-e2e/Graph/RecordReplay/transitive_queue_barrier.cpp
@@ -81,5 +81,8 @@ int main() {
     assert(check_value(i, (PatternA + PatternB), DataC[i], "DataC"));
   }
 
+  free(PtrA, Q1);
+  free(PtrB, Q1);
+  free(PtrC, Q1);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/regression_padded_parameter.cpp b/sycl/test-e2e/Graph/Update/regression_padded_parameter.cpp
index 8296a9b408d8d..ce9dd94e799c0 100644
--- a/sycl/test-e2e/Graph/Update/regression_padded_parameter.cpp
+++ b/sycl/test-e2e/Graph/Update/regression_padded_parameter.cpp
@@ -43,5 +43,6 @@ int main() {
     assert(HostDataA[i] == PaddedKernelParam.val1 * PaddedKernelParam.val2[i]);
   }
 
+  free(PtrA, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_before_finalize.cpp b/sycl/test-e2e/Graph/Update/update_before_finalize.cpp
index 3a7298a3a4c6d..40f215b48ce97 100644
--- a/sycl/test-e2e/Graph/Update/update_before_finalize.cpp
+++ b/sycl/test-e2e/Graph/Update/update_before_finalize.cpp
@@ -51,5 +51,7 @@ int main() {
     assert(HostDataB[i] == i);
   }
 
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_nd_range.cpp b/sycl/test-e2e/Graph/Update/update_nd_range.cpp
index 7a177a6092d33..a4ea838673150 100644
--- a/sycl/test-e2e/Graph/Update/update_nd_range.cpp
+++ b/sycl/test-e2e/Graph/Update/update_nd_range.cpp
@@ -49,5 +49,7 @@ int main() {
   for (size_t i = 0; i < N; i++) {
     assert(HostDataA[i] == (i >= 512 ? i : i * 2));
   }
+
+  free(PtrA, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_ndrange_to_range.cpp b/sycl/test-e2e/Graph/Update/update_ndrange_to_range.cpp
index 85bb495a9abc5..13d09b2ebea60 100644
--- a/sycl/test-e2e/Graph/Update/update_ndrange_to_range.cpp
+++ b/sycl/test-e2e/Graph/Update/update_ndrange_to_range.cpp
@@ -48,5 +48,7 @@ int main() {
   for (size_t i = 0; i < Size; i++) {
     assert(HostDataA[i] == (i >= UpdateSize ? i : i * 2));
   }
+
+  free(PtrA, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_range.cpp b/sycl/test-e2e/Graph/Update/update_range.cpp
index bdf5f61f5a938..ad2192ba171a0 100644
--- a/sycl/test-e2e/Graph/Update/update_range.cpp
+++ b/sycl/test-e2e/Graph/Update/update_range.cpp
@@ -49,5 +49,7 @@ int main() {
   for (size_t i = 0; i < N; i++) {
     assert(HostDataA[i] == (i >= 512 ? i : i * 2));
   }
+
+  free(PtrA, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_range_to_ndrange.cpp b/sycl/test-e2e/Graph/Update/update_range_to_ndrange.cpp
index 560c942365e49..3fc995bb085c1 100644
--- a/sycl/test-e2e/Graph/Update/update_range_to_ndrange.cpp
+++ b/sycl/test-e2e/Graph/Update/update_range_to_ndrange.cpp
@@ -49,5 +49,7 @@ int main() {
   for (size_t i = 0; i < Size; i++) {
     assert(HostDataA[i] == (i >= UpdateSize ? i : i * 2));
   }
+
+  free(PtrA, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_multiple_exec_graphs.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_multiple_exec_graphs.cpp
index 0a161b520f1f9..757749326c1f3 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_multiple_exec_graphs.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_multiple_exec_graphs.cpp
@@ -67,5 +67,8 @@ int main() {
     assert(HostDataA[i] == i * 3);
     assert(HostDataB[i] == i);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ordering.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ordering.cpp
index a5808f15b8ecb..a5e710facf7a7 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ordering.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ordering.cpp
@@ -68,5 +68,8 @@ int main() {
     assert(HostDataA[i] == i * NumKernelLoops * NumSubmitLoops);
     assert(HostDataB[i] == i * NumKernelLoops * NumSubmitLoops);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr.cpp
index 1d5ff53f49f68..c5d08d9e6084d 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr.cpp
@@ -60,5 +60,8 @@ int main() {
     assert(HostDataA[i] == i);
     assert(HostDataB[i] == i);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_3D.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_3D.cpp
index 0cc4209b0d149..b141beaa63eab 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_3D.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_3D.cpp
@@ -77,5 +77,8 @@ int main() {
     assert(HostDataA[i] == Ref);
     assert(HostDataB[i] == Ref);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_double_update.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_double_update.cpp
index 279a96bf5e035..d5895e11360f8 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_double_update.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_double_update.cpp
@@ -70,5 +70,9 @@ int main() {
     // Check that PtrUnused was never actually used in a kernel
     assert(HostDataUnused[i] == 0);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
+  free(PtrUnused, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes.cpp
index d079b72d1e66a..fa43f1b2cfbc7 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes.cpp
@@ -74,5 +74,8 @@ int main() {
     assert(HostDataA[i] == i * 2);
     assert(HostDataB[i] == i * 2);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes_different_indices.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes_different_indices.cpp
index d2c50997aecf2..0676e49b2e1ab 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes_different_indices.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_nodes_different_indices.cpp
@@ -77,5 +77,8 @@ int main() {
     assert(HostDataA[i] == i * 2);
     assert(HostDataB[i] == i * 2);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_params.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_params.cpp
index 0d9dd827d1691..268a232dd72b8 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_params.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_multiple_params.cpp
@@ -74,5 +74,9 @@ int main() {
   for (size_t i = 0; i < N; i++) {
     assert(OutData[i] == HostDataB[i] + (HostDataA[i] * HostDataC[i]));
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
+  free(PtrC, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_subgraph.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_subgraph.cpp
index 58ac32b59dad0..ab753c586ce61 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_ptr_subgraph.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_ptr_subgraph.cpp
@@ -70,5 +70,8 @@ int main() {
     assert(HostDataA[i] == i * 2);
     assert(HostDataB[i] == 0);
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_scalar.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_scalar.cpp
index 859f10dab2c14..298615cee16ba 100644
--- a/sycl/test-e2e/Graph/Update/update_with_indices_scalar.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_indices_scalar.cpp
@@ -55,5 +55,7 @@ int main() {
   for (size_t i = 0; i < N; i++) {
     assert(HostData[i] == 99);
   }
+
+  free(DeviceData, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg.cpp b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg.cpp
index c3c9a4b140407..ea971c6263906 100644
--- a/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg.cpp
+++ b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg.cpp
@@ -78,5 +78,8 @@ int main() {
     assert(HostDataA[i] == (i + ScalarA));
     assert(HostDataB[i] == (i + ScalarB));
   }
+
+  free(PtrA, Queue);
+  free(PtrB, Queue);
   return 0;
 }
diff --git a/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp b/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp
index d83677509f077..d40cd595d1741 100644
--- a/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp
+++ b/sycl/test-e2e/Graph/ValidUsage/linear_graph_copy_D2H.cpp
@@ -103,6 +103,8 @@ int main() {
   free(PtrA, Queue);
   free(PtrB, Queue);
   free(PtrC, Queue);
+  free(PtrBHost, Queue);
+  free(PtrCHost, Queue);
 
   for (size_t i = 0; i < Size; i++) {
     assert(check_value(i, ReferenceA[i], DataA[i], "DataA"));

From 37006c0d0301449cf6fc005ad6af0234cb7ca33b Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Thu, 8 Oct 2026 15:23:12 +0000
Subject: [PATCH 08/12] [UR][L0] Get async-free size via zeMemGetAddressRange
 for proxy pools

The UMF proxy pool (used when USM pooling is disabled) does not support
umfPoolMallocUsableSize(). In that case urEnqueueUSMFreeExp (V1) freed
the memory synchronously and returned early without signaling the event
or executing the command list, causing DEVICE_LOST errors and hangs.
V2 appendUSMFreeExp inserted the allocation into the async pool with
size 0, so it could never be reused.

Query the allocation size with zeMemGetAddressRange instead when
umfPoolMallocUsableSize() returns UMF_RESULT_ERROR_NOT_SUPPORTED.
---
 .../source/adapters/level_zero/async_alloc.cpp         | 10 +++++++++-
 .../adapters/level_zero/v2/command_list_manager.cpp    |  9 +++++++--
 2 files changed, 16 insertions(+), 3 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/async_alloc.cpp b/unified-runtime/source/adapters/level_zero/async_alloc.cpp
index 754752a53d7dc..76b713e03c243 100644
--- a/unified-runtime/source/adapters/level_zero/async_alloc.cpp
+++ b/unified-runtime/source/adapters/level_zero/async_alloc.cpp
@@ -258,7 +258,15 @@ ur_result_t urEnqueueUSMFreeExp(
 
   size_t Size = 0;
   UmfRet = umfPoolMallocUsableSize(UmfPool, Mem, &Size);
-  if (UmfRet != UMF_RESULT_SUCCESS) {
+  if (UmfRet == UMF_RESULT_ERROR_NOT_SUPPORTED) {
+    // Proxy pools (used when USM pooling is disabled) do not track allocation
+    // sizes. Query the size from the driver so that the allocation can still
+    // be deferred to the async pool: freeing it synchronously here could
+    // release memory still used by previously enqueued work, and a size of 0
+    // would make the allocation never reusable.
+    ZE2UR_CALL(zeMemGetAddressRange,
+               (Queue->Context->ZeContext, Mem, nullptr, &Size));
+  } else if (UmfRet != UMF_RESULT_SUCCESS) {
     return USMFreeHelper(Queue->Context, Mem);
   }
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index 6bfc2b8cd4e76..fdf8e3d7dc857 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -956,8 +956,13 @@ ur_result_t ur_command_list_manager::appendUSMFreeExp(
 
   size_t size = 0;
   umfRet = umfPoolMallocUsableSize(hPool, pMem, &size);
-  if (umfRet != UMF_RESULT_SUCCESS &&
-      umfRet != UMF_RESULT_ERROR_NOT_SUPPORTED) {
+  if (umfRet == UMF_RESULT_ERROR_NOT_SUPPORTED) {
+    // Proxy pools (used when USM pooling is disabled) do not track allocation
+    // sizes. Query the size from the driver, otherwise the allocation would
+    // be inserted into the async pool with size 0 and never be reused.
+    ZE2UR_CALL(zeMemGetAddressRange,
+               (hContext->getZeHandle(), pMem, nullptr, &size));
+  } else if (umfRet != UMF_RESULT_SUCCESS) {
     UR_LOG(ERR, "enqueueUSMFreeExp: failed to retrieve usable malloc size");
     return UR_RESULT_ERROR_UNKNOWN;
   }

From cdfd8e92ee7085c415dd840aa6255707f80a23d5 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Fri, 16 Oct 2026 07:07:54 +0000
Subject: [PATCH 09/12] [UR][L0] Use blocking free in V1 when USM pooling is
 disabled

When USM pooling is disabled (Xe2 with at least the minimum L0 driver
version), V1 uses UMF proxy pools, so every USM free goes directly to
zeMemFree and the memory is unmapped immediately. With the disjoint
pool, freed memory stayed mapped, so GPU accesses still in flight were
harmless; now they can trigger GPU page faults, engine resets and
DEVICE_LOST cascades.

The V2 adapter already sets the BLOCKING_FREE policy on its UMF Level
Zero provider. Do the same in V1's L0MemoryProvider when pooling is
disabled by calling zeMemFreeExt with
ZE_DRIVER_MEMORY_FREE_POLICY_EXT_FLAG_BLOCKING_FREE.
---
 .../source/adapters/level_zero/usm.cpp        | 29 +++++++++++++++++--
 .../source/adapters/level_zero/usm.hpp        |  1 +
 2 files changed, 27 insertions(+), 3 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/usm.cpp b/unified-runtime/source/adapters/level_zero/usm.cpp
index 8b30852cdba31..de80c8e990daf 100644
--- a/unified-runtime/source/adapters/level_zero/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/usm.cpp
@@ -726,10 +726,20 @@ ur_result_t UR_APICALL urUSMHostAllocUnregisterExp(
   return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
 }
 
-static ur_result_t USMFreeImpl(ur_context_handle_t Context, void *Ptr) {
+static ur_result_t USMFreeImpl(ur_context_handle_t Context, void *Ptr,
+                               bool BlockingFree = false) {
   ur_result_t Res = UR_RESULT_SUCCESS;
   if (checkL0LoaderTeardown()) {
-    auto ZeResult = ZE_CALL_NOCHECK(zeMemFree, (Context->ZeContext, Ptr));
+    ze_result_t ZeResult;
+    if (BlockingFree) {
+      ze_memory_free_ext_desc_t FreeDesc = {};
+      FreeDesc.stype = ZE_STRUCTURE_TYPE_MEMORY_FREE_EXT_DESC;
+      FreeDesc.freePolicy = ZE_DRIVER_MEMORY_FREE_POLICY_EXT_FLAG_BLOCKING_FREE;
+      ZeResult =
+          ZE_CALL_NOCHECK(zeMemFreeExt, (Context->ZeContext, &FreeDesc, Ptr));
+    } else {
+      ZeResult = ZE_CALL_NOCHECK(zeMemFree, (Context->ZeContext, Ptr));
+    }
     // Handle When the driver is already released
     if (ZeResult == ZE_RESULT_ERROR_UNINITIALIZED) {
       Res = UR_RESULT_SUCCESS;
@@ -756,6 +766,19 @@ L0MemoryProvider::initialize(ur_context_handle_t Ctx,
   Context = Ctx;
   Device = Dev;
 
+  // Without USM pooling every free goes directly to the driver, so the memory
+  // may still be in use by the device. Wait for it like the V2 adapter does.
+  if (Dev) {
+    BlockingFree = Dev->isUsmPoolingDisabled();
+  } else {
+    const auto &Devices = Ctx->getDevices();
+    BlockingFree =
+        !Devices.empty() &&
+        std::all_of(Devices.begin(), Devices.end(), [](ur_device_handle_t D) {
+          return D->isUsmPoolingDisabled();
+        });
+  }
+
   return UMF_RESULT_SUCCESS;
 }
 
@@ -776,7 +799,7 @@ enum umf_result_t L0MemoryProvider::alloc(size_t Size, size_t Align,
 enum umf_result_t L0MemoryProvider::free(void *Ptr, size_t Size) {
   (void)Size;
 
-  auto Res = USMFreeImpl(Context, Ptr);
+  auto Res = USMFreeImpl(Context, Ptr, BlockingFree);
   if (Res != UR_RESULT_SUCCESS) {
     getLastStatusRef() = Res;
     return UMF_RESULT_ERROR_MEMORY_PROVIDER_SPECIFIC;
diff --git a/unified-runtime/source/adapters/level_zero/usm.hpp b/unified-runtime/source/adapters/level_zero/usm.hpp
index 2e2be2c9108a2..48667867319bb 100644
--- a/unified-runtime/source/adapters/level_zero/usm.hpp
+++ b/unified-runtime/source/adapters/level_zero/usm.hpp
@@ -177,6 +177,7 @@ class L0MemoryProvider : public USMMemoryProviderBase {
   size_t MinPageSize = 0;
   bool MinPageSizeCached = false;
   AllocationStats AllocStats;
+  bool BlockingFree = false;
 
 public:
   umf_result_t initialize(ur_context_handle_t Ctx,

From 1ea9ed593db7c59a93bc4aab65ff75c7cb0853cc Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Fri, 16 Oct 2026 08:35:25 +0000
Subject: [PATCH 10/12] [UR][L0] Clean up USM pooling disabled detection

- Add isUsmPoolingDisabled(Context, Device) and use it both for pool
  descriptors and in L0MemoryProvider::initialize() to decide on
  blocking free, instead of duplicating the all-devices check.
- Log the pooling-disabled message at DEBUG level, since it is emitted
  on every pool creation.
- Document that the V2 host allocation size check applies to all
  devices and uses the largest maxMemAllocSize in the context.
- Document why zeMemGetAddressRange gives the correct size in the async
  free paths (USM free always gets the allocation's base pointer).
- List in usm_pooling_disabled.hpp all places duplicating the
  detection logic that must be kept in sync.

Signed-off-by: Lukasz Dorau <lukasz.dorau@intel.com>
---
 .../adapters/level_zero/async_alloc.cpp       |  4 ++-
 .../adapters/level_zero/common/device.cpp     | 29 +++++++++++--------
 .../adapters/level_zero/common/device.hpp     | 11 +++++--
 .../common/usm_pooling_disabled.hpp           |  7 +++++
 .../source/adapters/level_zero/usm.cpp        | 11 +------
 .../level_zero/v2/command_list_manager.cpp    |  3 ++
 .../source/adapters/level_zero/v2/usm.cpp     |  3 ++
 7 files changed, 42 insertions(+), 26 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/async_alloc.cpp b/unified-runtime/source/adapters/level_zero/async_alloc.cpp
index 76b713e03c243..24c77063557a7 100644
--- a/unified-runtime/source/adapters/level_zero/async_alloc.cpp
+++ b/unified-runtime/source/adapters/level_zero/async_alloc.cpp
@@ -263,7 +263,9 @@ ur_result_t urEnqueueUSMFreeExp(
     // sizes. Query the size from the driver so that the allocation can still
     // be deferred to the async pool: freeing it synchronously here could
     // release memory still used by previously enqueued work, and a size of 0
-    // would make the allocation never reusable.
+    // would make the allocation never reusable. zeMemGetAddressRange returns
+    // the size of the whole allocation, which is correct here because a USM
+    // free must be given the base pointer of an allocation.
     ZE2UR_CALL(zeMemGetAddressRange,
                (Queue->Context->ZeContext, Mem, nullptr, &Size));
   } else if (UmfRet != UMF_RESULT_SUCCESS) {
diff --git a/unified-runtime/source/adapters/level_zero/common/device.cpp b/unified-runtime/source/adapters/level_zero/common/device.cpp
index 41f688e429203..b49a3f37e533f 100644
--- a/unified-runtime/source/adapters/level_zero/common/device.cpp
+++ b/unified-runtime/source/adapters/level_zero/common/device.cpp
@@ -2033,21 +2033,26 @@ bool ur_device_handle_t_::isUsmPoolingDisabled() {
                                UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR, 0);
 }
 
+bool isUsmPoolingDisabled(const ur_context_common_t *Context,
+                          ur_device_handle_t Device) {
+  if (Device)
+    return Device->isUsmPoolingDisabled();
+
+  const auto &Devices = Context->getDevices();
+  return !Devices.empty() &&
+         std::all_of(Devices.begin(), Devices.end(),
+                     [](ur_device_handle_t Device) {
+                       return Device->isUsmPoolingDisabled();
+                     });
+}
+
 bool isUsmPoolingDisabled(const usm::pool_descriptor &Desc) {
-  bool Disabled = false;
-  if (Desc.hDevice) {
-    Disabled = common_cast(Desc.hDevice)->isUsmPoolingDisabled();
-  } else {
-    const auto &Devices = common_cast(Desc.hContext)->getDevices();
-    Disabled =
-        !Devices.empty() && std::all_of(Devices.begin(), Devices.end(),
-                                        [](ur_device_handle_t Device) {
-                                          return Device->isUsmPoolingDisabled();
-                                        });
-  }
+  bool Disabled =
+      isUsmPoolingDisabled(common_cast(Desc.hContext),
+                           Desc.hDevice ? common_cast(Desc.hDevice) : nullptr);
 
   if (Disabled) {
-    UR_LOG(INFO,
+    UR_LOG(DEBUG,
            "USM pooling is disabled on Xe2 or newer devices with L0 driver "
            "{}.{} or newer, desc:{}",
            UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR,
diff --git a/unified-runtime/source/adapters/level_zero/common/device.hpp b/unified-runtime/source/adapters/level_zero/common/device.hpp
index 2222f8ccd9d1d..f7bd764b498e2 100644
--- a/unified-runtime/source/adapters/level_zero/common/device.hpp
+++ b/unified-runtime/source/adapters/level_zero/common/device.hpp
@@ -359,9 +359,14 @@ inline std::vector<ur_device_handle_t> CollectDevicesForUsmPoolCreation(
   return DevicesAndSubDevices;
 }
 
-// Returns true if USM pooling is disabled for the pool described by Desc.
-// Host pools are not associated with any device, so pooling of host memory is
-// disabled only if it is disabled for all devices in the context.
+// Returns true if USM pooling is disabled for memory of the given Device in
+// Context. A null Device denotes host memory, which is not associated with any
+// device, so its pooling is disabled only if it is disabled for all devices in
+// the context.
+bool isUsmPoolingDisabled(const ur_context_common_t *Context,
+                          ur_device_handle_t Device);
+
+// Same as above for the pool described by Desc.
 bool isUsmPoolingDisabled(const usm::pool_descriptor &Desc);
 
 } // namespace ur::level_zero
diff --git a/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp b/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp
index c5c1c8c5851e7..d622346d7316d 100644
--- a/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp
+++ b/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp
@@ -9,5 +9,12 @@
 
 // Minimum L0 driver version (major.minor) starting from which USM pooling
 // in the adapter is disabled on Xe2 or newer devices.
+//
+// The detection logic (Xe2 or newer = Intel GPU with IP version >= 0x05004000,
+// i.e. 20.1.0, plus this driver version) is duplicated in the following places,
+// which must be kept in sync with ur_device_handle_t_::isUsmPoolingDisabled():
+// - sycl/test-e2e/lit.cfg.py (level_zero_usm_pooling_disabled feature),
+// - unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
+//   (isUsmPoolingDisabledForDevice()).
 #define UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR 1
 #define UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MINOR 17
diff --git a/unified-runtime/source/adapters/level_zero/usm.cpp b/unified-runtime/source/adapters/level_zero/usm.cpp
index de80c8e990daf..36abd5f8cb4f2 100644
--- a/unified-runtime/source/adapters/level_zero/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/usm.cpp
@@ -768,16 +768,7 @@ L0MemoryProvider::initialize(ur_context_handle_t Ctx,
 
   // Without USM pooling every free goes directly to the driver, so the memory
   // may still be in use by the device. Wait for it like the V2 adapter does.
-  if (Dev) {
-    BlockingFree = Dev->isUsmPoolingDisabled();
-  } else {
-    const auto &Devices = Ctx->getDevices();
-    BlockingFree =
-        !Devices.empty() &&
-        std::all_of(Devices.begin(), Devices.end(), [](ur_device_handle_t D) {
-          return D->isUsmPoolingDisabled();
-        });
-  }
+  BlockingFree = ur::level_zero::isUsmPoolingDisabled(Ctx, Dev);
 
   return UMF_RESULT_SUCCESS;
 }
diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
index fdf8e3d7dc857..c8df02056a409 100644
--- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp
@@ -960,6 +960,9 @@ ur_result_t ur_command_list_manager::appendUSMFreeExp(
     // Proxy pools (used when USM pooling is disabled) do not track allocation
     // sizes. Query the size from the driver, otherwise the allocation would
     // be inserted into the async pool with size 0 and never be reused.
+    // zeMemGetAddressRange returns the size of the whole allocation, which is
+    // correct here because a USM free must be given the base pointer of an
+    // allocation.
     ZE2UR_CALL(zeMemGetAddressRange,
                (hContext->getZeHandle(), pMem, nullptr, &size));
   } else if (umfRet != UMF_RESULT_SUCCESS) {
diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.cpp b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
index c949e981b9b25..9a558efa15c91 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
@@ -290,6 +290,9 @@ ur_result_t ur_usm_pool_handle_t_::allocate(
   // sizes (e.g. SIZE_MAX). Reject host allocations larger than
   // maxMemAllocSize unless relaxed allocation limits are enabled, the same
   // way the V1 adapter does (it calls zeMemAllocHost without that descriptor).
+  // This applies to all devices, not only to those with USM pooling disabled.
+  // Host memory is accessible from every device in the context, so the
+  // largest maxMemAllocSize among them is used as the limit.
   if (type == UR_USM_TYPE_HOST) {
     bool relaxedLimits = false;
     uint64_t maxAllocSize = 0;

From 8e274efa8c0fa4c545d3b7f1b5ea3270b5da36cb Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Tue, 20 Oct 2026 07:15:49 +0000
Subject: [PATCH 11/12] [UR][L0] Address USM pooling review comments

Cache the effective host USM allocation limit in the V2 context instead
of recomputing it on every allocation. Preserve the existing relaxed-limit
behavior and simplify explanatory comments in the allocation and test
code.
---
 .../source/adapters/level_zero/v2/context.cpp |  9 +++++++++
 .../source/adapters/level_zero/v2/context.hpp |  6 ++++++
 .../source/adapters/level_zero/v2/usm.cpp     | 20 ++++---------------
 .../adapters/level_zero/enqueue_alloc.cpp     |  9 +--------
 .../test/conformance/exp_graph/fixtures.h     |  8 +-------
 5 files changed, 21 insertions(+), 31 deletions(-)

diff --git a/unified-runtime/source/adapters/level_zero/v2/context.cpp b/unified-runtime/source/adapters/level_zero/v2/context.cpp
index 637d2fd55c95f..7dec7db40ef8d 100644
--- a/unified-runtime/source/adapters/level_zero/v2/context.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/context.cpp
@@ -110,6 +110,15 @@ ur_context_handle_t_::ur_context_handle_t_(ze_context_handle_t hContext,
                                  this, v2::QUEUE_IMMEDIATE,
                                  v2::EVENT_FLAGS_PROFILING_ENABLED)),
       defaultUSMPool(this, nullptr), asyncPool(this, nullptr) {
+  bool relaxedLimits = false;
+  for (auto *device : getDevices()) {
+    relaxedLimits |= device->useRelaxedAllocationLimits();
+    maxHostAllocSize = std::max<uint64_t>(
+        maxHostAllocSize, device->ZeDeviceProperties->maxMemAllocSize);
+  }
+  if (relaxedLimits)
+    maxHostAllocSize = 0;
+
   UR_LOG(INFO, "UR context created with {} devices", numDevices);
 }
 
diff --git a/unified-runtime/source/adapters/level_zero/v2/context.hpp b/unified-runtime/source/adapters/level_zero/v2/context.hpp
index fc6270a6fd986..8a7a618127e47 100644
--- a/unified-runtime/source/adapters/level_zero/v2/context.hpp
+++ b/unified-runtime/source/adapters/level_zero/v2/context.hpp
@@ -34,6 +34,8 @@ struct ur_context_handle_t_ : ur_context_common_t {
   ur_usm_pool_handle_t getDefaultUSMPool();
   ur_usm_pool_handle_t getAsyncPool();
 
+  uint64_t getMaxHostAllocSize() const { return maxHostAllocSize; }
+
   void addUsmPool(ur_usm_pool_handle_t hPool);
   void removeUsmPool(ur_usm_pool_handle_t hPool);
   void changeResidentDevice(ur_device_handle_t hDevice,
@@ -87,6 +89,10 @@ struct ur_context_handle_t_ : ur_context_common_t {
   ur_shared_mutex GraphMapMutex;
 
 private:
+  // Largest maxMemAllocSize among the context's devices.
+  // Zero means relaxed allocation limits are enabled or no limit is available.
+  uint64_t maxHostAllocSize = 0;
+
   const raii::ze_context_handle_t hContext;
   command_list_cache_t commandListCache;
   event_pool_cache eventPoolCacheImmediate;
diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.cpp b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
index 9a558efa15c91..89024be9c3b23 100644
--- a/unified-runtime/source/adapters/level_zero/v2/usm.cpp
+++ b/unified-runtime/source/adapters/level_zero/v2/usm.cpp
@@ -285,23 +285,11 @@ ur_result_t ur_usm_pool_handle_t_::allocate(
     return UR_RESULT_ERROR_INVALID_VALUE;
   }
 
-  // The UMF Level Zero provider always passes the relaxed allocation limits
-  // descriptor to zeMemAllocHost, which makes some drivers accept absurd
-  // sizes (e.g. SIZE_MAX). Reject host allocations larger than
-  // maxMemAllocSize unless relaxed allocation limits are enabled, the same
-  // way the V1 adapter does (it calls zeMemAllocHost without that descriptor).
-  // This applies to all devices, not only to those with USM pooling disabled.
-  // Host memory is accessible from every device in the context, so the
-  // largest maxMemAllocSize among them is used as the limit.
+  // UMF requests relaxed limits from the driver even when UR does not.
+  // Enforce the cached context limit regardless of pooling mode.
   if (type == UR_USM_TYPE_HOST) {
-    bool relaxedLimits = false;
-    uint64_t maxAllocSize = 0;
-    for (auto *device : hContext->getDevices()) {
-      relaxedLimits |= device->useRelaxedAllocationLimits();
-      maxAllocSize = std::max<uint64_t>(
-          maxAllocSize, device->ZeDeviceProperties->maxMemAllocSize);
-    }
-    if (!relaxedLimits && maxAllocSize > 0 && size > maxAllocSize) {
+    const auto maxAllocSize = hContext->getMaxHostAllocSize();
+    if (maxAllocSize > 0 && size > maxAllocSize) {
       return UR_RESULT_ERROR_INVALID_USM_SIZE;
     }
   }
diff --git a/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp b/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
index 7784be393f75f..ac17009aa5126 100644
--- a/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
+++ b/unified-runtime/test/adapters/level_zero/enqueue_alloc.cpp
@@ -34,14 +34,7 @@ std::ostream &operator<<(std::ostream &os, EnqueueAllocTestParam param) {
   return os;
 }
 
-// Mirrors ur_device_handle_t_::isUsmPoolingDisabled() in
-// adapters/level_zero/common/device.cpp: USM pooling (and therefore the
-// address-reuse optimization it enables) is disabled by the adapter for
-// Xe2-or-newer devices when running on L0 driver
-// UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer. On such
-// configurations the adapter falls back to a UMF proxy pool, which has no
-// reuse semantics, so tests must not assume a freed allocation's address
-// will be reused.
+// Mirrors how the Level Zero adapter determines whether USM pooling is disabled.
 bool isUsmPoolingDisabledForDevice(ur_device_handle_t device) {
   uint32_t vendorId = 0;
   if (urDeviceGetInfo(device, UR_DEVICE_INFO_VENDOR_ID, sizeof(vendorId),
diff --git a/unified-runtime/test/conformance/exp_graph/fixtures.h b/unified-runtime/test/conformance/exp_graph/fixtures.h
index bec386505be82..4493b86a291b2 100644
--- a/unified-runtime/test/conformance/exp_graph/fixtures.h
+++ b/unified-runtime/test/conformance/exp_graph/fixtures.h
@@ -91,13 +91,7 @@ struct urGraphPopulatedExpTest : urGraphExpTest {
 
   void TearDown() override {
     if (deviceMem) {
-      // Reset the data while deviceMem is still a valid allocation. Freeing
-      // first and then enqueuing a fill on the freed pointer is a
-      // use-after-free: it could appear to work when USM pooling keeps freed
-      // memory mapped for reuse, but is not actually valid and fails (or
-      // crashes) once pooling is disabled, e.g. on Xe2+ devices with L0
-      // driver UR_L0_USM_POOLING_DISABLED_MIN_DRIVER_MAJOR.MINOR or newer
-      // where a UMF proxy pool (no pooling) is used instead.
+      // Reset the data before freeing the allocation.
       resetData();
       ASSERT_SUCCESS(urUSMFree(context, deviceMem));
     }

From a0356fdcd1e4b884ac510bc8611436bbdccf2888 Mon Sep 17 00:00:00 2001
From: Lukasz Dorau <lukasz.dorau@intel.com>
Date: Tue, 20 Oct 2026 11:18:46 +0000
Subject: [PATCH 12/12] [SYCL][E2E] Free USM in
 dyn_work_group_memory_multiple_graphs leak check

The test allocated device USM but never freed it, so the L0 leak check
reported LEAK = 1. Match the other dyn_work_group_memory Graph tests.
---
 .../Graph/Update/dyn_work_group_memory_multiple_graphs.cpp       | 1 +
 1 file changed, 1 insertion(+)

diff --git a/sycl/test-e2e/Graph/Update/dyn_work_group_memory_multiple_graphs.cpp b/sycl/test-e2e/Graph/Update/dyn_work_group_memory_multiple_graphs.cpp
index f0195bd3f0521..e618d5ac905c8 100644
--- a/sycl/test-e2e/Graph/Update/dyn_work_group_memory_multiple_graphs.cpp
+++ b/sycl/test-e2e/Graph/Update/dyn_work_group_memory_multiple_graphs.cpp
@@ -71,5 +71,6 @@ int main() {
     assert(check_value(i, HostData[i], HostOutputCompare[i], "HostData"));
   }
 
+  free(Ptr, Queue);
   return 0;
 }