[TEST][UR][L0] Disable USM pooling on Xe2+ devices with L0 driver 1.18+

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
 

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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;
 }

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"));

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;

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.

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")

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

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

@@ -258,7 +258,17 @@ 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. 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) {
     return USMFreeHelper(Queue->Context, Mem);
   }