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/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/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; } 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")); 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 class K; 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 8cdf938269081..663ea10b877ef 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -1104,6 +1104,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) ): @@ -1121,14 +1127,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: @@ -1254,6 +1281,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") diff --git a/unified-runtime/source/adapters/level_zero/CMakeLists.txt b/unified-runtime/source/adapters/level_zero/CMakeLists.txt index 4c114ba958665..9695156ee892d 100644 --- a/unified-runtime/source/adapters/level_zero/CMakeLists.txt +++ b/unified-runtime/source/adapters/level_zero/CMakeLists.txt @@ -6,8 +6,8 @@ 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/ze_intel_ext.hpp - ${CMAKE_CURRENT_SOURCE_DIR}/common/adapter.cpp ${CMAKE_CURRENT_SOURCE_DIR}/common/device.cpp ${CMAKE_CURRENT_SOURCE_DIR}/common/platform.cpp diff --git a/unified-runtime/source/adapters/level_zero/async_alloc.cpp b/unified-runtime/source/adapters/level_zero/async_alloc.cpp index 754752a53d7dc..24c77063557a7 100644 --- a/unified-runtime/source/adapters/level_zero/async_alloc.cpp +++ b/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); } diff --git a/unified-runtime/source/adapters/level_zero/common/device.cpp b/unified-runtime/source/adapters/level_zero/common/device.cpp index 660946b21be2e..65c832a027d63 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 #include #if defined(__linux__) @@ -2024,6 +2026,41 @@ 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 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 = + isUsmPoolingDisabled(common_cast(Desc.hContext), + Desc.hDevice ? common_cast(Desc.hDevice) : nullptr); + + if (Disabled) { + 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, + 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 bca35786d8d51..eeceaec0410d5 100644 --- a/unified-runtime/source/adapters/level_zero/common/device.hpp +++ b/unified-runtime/source/adapters/level_zero/common/device.hpp @@ -31,6 +31,10 @@ #include #include +namespace usm { +struct pool_descriptor; +} // namespace usm + namespace ur::level_zero { typedef size_t DeviceId; @@ -305,6 +309,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); } @@ -407,4 +416,14 @@ inline std::vector CollectDevicesForUsmPoolCreation( return DevicesAndSubDevices; } +// 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 new file mode 100644 index 0000000000000..d622346d7316d --- /dev/null +++ b/unified-runtime/source/adapters/level_zero/common/usm_pooling_disabled.hpp @@ -0,0 +1,20 @@ +//===--------- 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. +// +// 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 fc255e9048826..36abd5f8cb4f2 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 @@ -724,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; @@ -754,6 +766,10 @@ 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. + BlockingFree = ur::level_zero::isUsmPoolingDisabled(Ctx, Dev); + return UMF_RESULT_SUCCESS; } @@ -774,7 +790,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; @@ -1076,8 +1092,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 +1123,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) { + bool UseProxy = IsProxy || isUsmPoolingDisabled(Desc); + if (UseProxy) { Pool = usm::makeProxyPool(MakeProvider(&Desc)); } else { auto &PoolConfig = @@ -1115,7 +1133,7 @@ v1::ur_usm_pool_handle_t_::ur_usm_pool_handle_t_(ur_context_handle_t Context, } std::unique_ptr usmPool = - std::make_unique(this, std::move(Pool)); + std::make_unique(this, std::move(Pool), UseProxy); auto Ret = umf::umf2urResult( umfPoolSetTag(usmPool->UmfPool.get(), usmPool.get(), nullptr)); if (Ret) { @@ -1174,11 +1192,18 @@ 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; + bool UseProxy = isUsmPoolingDisabled(Desc); + if (UseProxy) { + Pool = usm::makeProxyPool(MakeProvider(&Desc)); + } else { + auto &PoolConfig = + DisjointPoolConfigs.Configs[DescToDisjointPoolMemType(Desc)]; + Pool = usm::makeDisjointPool(MakeProvider(&Desc), PoolConfig); + } - std::unique_ptr usmPool = std::make_unique( - this, usm::makeDisjointPool(MakeProvider(&Desc), PoolConfig)); + std::unique_ptr usmPool = + std::make_unique(this, std::move(Pool), UseProxy); auto Ret = umf::umf2urResult( umfPoolSetTag(usmPool->UmfPool.get(), usmPool.get(), nullptr)); if (Ret) { @@ -1402,12 +1427,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..48667867319bb 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 PoolManager; AllocationStats AllocStats; }; @@ -174,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, 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..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 @@ -956,7 +956,16 @@ 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_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. + // 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) { UR_LOG(ERR, "enqueueUSMFreeExp: failed to retrieve usable malloc size"); return UR_RESULT_ERROR_UNKNOWN; } 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( + 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 ff00971ad3dba..89024be9c3b23 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; - if (disjointPoolConfigs.has_value()) { + 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(this, std::move(pool)); + usmPool = std::make_unique(this, std::move(pool), isProxy); } else { auto pool = usm::makeProxyPool(makeProvider(desc)); - usmPool = std::make_unique(this, std::move(pool)); + usmPool = std::make_unique(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; - if (disjointPoolConfigs.has_value()) { + 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(this, std::move(pool)); + usmPool = std::make_unique(this, std::move(pool), isProxy); } else { auto pool = usm::makeProxyPool(makeProvider(desc)); - usmPool = std::make_unique(this, std::move(pool)); + usmPool = std::make_unique(this, std::move(pool), isProxy); } UMF_CALL_THROWS( umfPoolSetTag(usmPool->umfPool.get(), usmPool.get(), nullptr)); @@ -281,6 +285,15 @@ ur_result_t ur_usm_pool_handle_t_::allocate( return UR_RESULT_ERROR_INVALID_VALUE; } + // 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) { + const auto maxAllocSize = hContext->getMaxHostAllocSize(); + if (maxAllocSize > 0 && size > maxAllocSize) { + return UR_RESULT_ERROR_INVALID_USM_SIZE; + } + } + auto deviceFlags = getDeviceFlags(pUSMDesc); auto pool = getPool(usm::pool_descriptor{ @@ -437,11 +450,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; +} + +ur_result_t ur_usm_pool_handle_t_::getTotalUsedSize(size_t &usedSize) { + usedSize = allocStats.getCurrent(); + return getProxyUsedSize(false, usedSize); } -size_t ur_usm_pool_handle_t_::getPeakUsedSize() { return allocStats.getPeak(); } +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 +651,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 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; diff --git a/unified-runtime/test/adapters/level_zero/CMakeLists.txt b/unified-runtime/test/adapters/level_zero/CMakeLists.txt index a35ecd6696828..3b461a75645dd 100644 --- a/unified-runtime/test/adapters/level_zero/CMakeLists.txt +++ b/unified-runtime/test/adapters/level_zero/CMakeLists.txt @@ -42,6 +42,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..ac17009aa5126 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 #include "unified-runtime/ur_api.h" +#include "usm_pooling_disabled.hpp" #include "uur/utils.h" #include @@ -33,6 +34,31 @@ std::ostream &operator<<(std::ostream &os, EnqueueAllocTestParam param) { return os; } +// 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), + &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 { void SetUp() override { @@ -235,7 +261,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)); diff --git a/unified-runtime/test/conformance/exp_graph/fixtures.h b/unified-runtime/test/conformance/exp_graph/fixtures.h index f8445392e9b52..4493b86a291b2 100644 --- a/unified-runtime/test/conformance/exp_graph/fixtures.h +++ b/unified-runtime/test/conformance/exp_graph/fixtures.h @@ -91,8 +91,9 @@ struct urGraphPopulatedExpTest : urGraphExpTest { void TearDown() override { if (deviceMem) { - ASSERT_SUCCESS(urUSMFree(context, deviceMem)); + // Reset the data before freeing the allocation. resetData(); + ASSERT_SUCCESS(urUSMFree(context, deviceMem)); } UUR_RETURN_ON_FATAL_FAILURE(urGraphExpTest::TearDown());