From fa1080884e9356246c001c7675362e95360c0e83 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=C5=81ukasz=20=C5=9Alusarczyk?= Date: Fri, 11 Jul 2025 12:52:56 +0200 Subject: [PATCH 1/2] [SYCL][Graph] test exposing zeVirtualMemQueryPageSize not using allocation size --- .../Explicit/async_alloc_different_sizes.cpp | 10 ++++++ .../Inputs/async_alloc_different_sizes.cpp | 34 +++++++++++++++++++ .../async_alloc_different_sizes.cpp | 10 ++++++ 3 files changed, 54 insertions(+) create mode 100644 sycl/test-e2e/Graph/AsyncAlloc/Explicit/async_alloc_different_sizes.cpp create mode 100644 sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp create mode 100644 sycl/test-e2e/Graph/AsyncAlloc/RecordReplay/async_alloc_different_sizes.cpp diff --git a/sycl/test-e2e/Graph/AsyncAlloc/Explicit/async_alloc_different_sizes.cpp b/sycl/test-e2e/Graph/AsyncAlloc/Explicit/async_alloc_different_sizes.cpp new file mode 100644 index 0000000000000..aa3183d1e5d01 --- /dev/null +++ b/sycl/test-e2e/Graph/AsyncAlloc/Explicit/async_alloc_different_sizes.cpp @@ -0,0 +1,10 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +#define GRAPH_E2E_EXPLICIT + +#include "../Inputs/async_alloc_different_sizes.cpp" diff --git a/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp new file mode 100644 index 0000000000000..40ea023b4dd3c --- /dev/null +++ b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp @@ -0,0 +1,34 @@ +// Tests async allocations with different sizes. + +#include "../../graph_common.hpp" +#include + +void asyncAllocWorksWithSize(size_t Size) { + queue Queue{}; + + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + void *AsyncPtr = nullptr; + // Add alloc node + auto AllocNode = add_node(Graph, Queue, [&](handler &CGH) { + AsyncPtr = exp_ext::async_malloc(CGH, usm::alloc::device, Size); + }); + + add_node( + Graph, Queue, + [&](handler &CGH) { + depends_on_helper(CGH, AllocNode); + exp_ext::async_free(CGH, AsyncPtr); + }, + AllocNode); + + auto GraphExec = Graph.finalize(); +} + +int main() { + asyncAllocWorksWithSize(1); + asyncAllocWorksWithSize(131); + asyncAllocWorksWithSize(10071); + asyncAllocWorksWithSize(1007177); + // asyncAllocWorksWithSize(191439360); // BUG +} \ No newline at end of file diff --git a/sycl/test-e2e/Graph/AsyncAlloc/RecordReplay/async_alloc_different_sizes.cpp b/sycl/test-e2e/Graph/AsyncAlloc/RecordReplay/async_alloc_different_sizes.cpp new file mode 100644 index 0000000000000..9c04c83782937 --- /dev/null +++ b/sycl/test-e2e/Graph/AsyncAlloc/RecordReplay/async_alloc_different_sizes.cpp @@ -0,0 +1,10 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +#define GRAPH_E2E_RECORD_REPLAY + +#include "../Inputs/async_alloc_different_sizes.cpp" \ No newline at end of file From 889a26acf72b3819c17454bfe141f5ae65e3384b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=C5=81ukasz=20=C5=9Alusarczyk?= Date: Fri, 11 Jul 2025 16:29:45 +0200 Subject: [PATCH 2/2] [SYCL][Grapg] async_malloc use allocation size for zeVirtualMemQueryPageSize --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 3 + .../ext/oneapi/virtual_mem/virtual_mem.hpp | 5 ++ sycl/source/detail/graph/memory_pool.cpp | 3 +- sycl/source/virtual_mem.cpp | 15 ++++- .../Inputs/async_alloc_different_sizes.cpp | 2 +- unified-runtime/include/ur_api.h | 3 + unified-runtime/include/ur_ddi.h | 4 +- unified-runtime/include/ur_print.hpp | 5 ++ .../scripts/core/virtual_memory.yml | 3 + .../source/adapters/cuda/virtual_mem.cpp | 1 + .../source/adapters/hip/virtual_mem.cpp | 4 +- .../level_zero/ur_interface_loader.hpp | 4 +- .../adapters/level_zero/virtual_mem.cpp | 7 ++- .../source/adapters/mock/ur_mockddi.cpp | 5 +- .../adapters/native_cpu/virtual_mem.cpp | 4 +- .../source/adapters/opencl/virtual_mem.cpp | 4 +- .../sanitizer_common/sanitizer_utils.cpp | 7 ++- .../loader/layers/tracing/ur_trcddi.cpp | 10 +++- .../loader/layers/validation/ur_valddi.cpp | 7 ++- unified-runtime/source/loader/ur_ldrddi.cpp | 6 +- unified-runtime/source/loader/ur_libapi.cpp | 6 +- unified-runtime/source/ur_api.cpp | 2 + .../enqueue/urEnqueueKernelLaunch.cpp | 4 +- .../testing/include/uur/fixtures.h | 15 +++-- .../urVirtualMemGranularityGetInfo.cpp | 59 ++++++++++--------- 25 files changed, 124 insertions(+), 64 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 8e5c59ef0bdea..33368f2b9c197 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -28,6 +28,9 @@ set(UR_BUILD_TESTS "${SYCL_UR_BUILD_TESTS}" CACHE BOOL "" FORCE) # UR tests require the examples to be built set(UR_BUILD_EXAMPLES "${SYCL_UR_BUILD_TESTS}" CACHE BOOL "" FORCE) +option(SYCL_UR_FORMAT_CPP_STYLE "Format code style of UR C++ sources" OFF) +set(UR_FORMAT_CPP_STYLE "${SYCL_UR_FORMAT_CPP_STYLE}" CACHE BOOL "" FORCE) + # Here we override the defaults to unified-runtime set(UR_BUILD_XPTI_LIBS OFF CACHE BOOL "") set(UR_ENABLE_SYMBOLIZER ON CACHE BOOL "Enable symbolizer for sanitizer layer.") diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/virtual_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/virtual_mem.hpp index 74a42354eaa01..5d7cfb49756e5 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/virtual_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/virtual_mem.hpp @@ -30,6 +30,11 @@ __SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, const context &SyclContext, granularity_mode Mode = granularity_mode::recommended); +__SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, + const context &SyclContext, + granularity_mode Mode, + size_t allocationSize); + __SYCL_EXPORT size_t get_mem_granularity(const context &SyclContext, granularity_mode Mode = granularity_mode::recommended); diff --git a/sycl/source/detail/graph/memory_pool.cpp b/sycl/source/detail/graph/memory_pool.cpp index 869ef5985093d..d6043c468f781 100644 --- a/sycl/source/detail/graph/memory_pool.cpp +++ b/sycl/source/detail/graph/memory_pool.cpp @@ -42,7 +42,8 @@ void *graph_mem_pool::malloc(size_t Size, usm::alloc AllocType, context_impl &CtxImpl = *getSyclObjImpl(MContext); adapter_impl &Adapter = CtxImpl.getAdapter(); - size_t Granularity = get_mem_granularity(MDevice, MContext); + const size_t Granularity = get_mem_granularity( + MDevice, MContext, granularity_mode::recommended, Size); uintptr_t StartPtr = 0; size_t AlignedSize = alignByteSize(Size, Granularity); // See if we can find an allocation to reuse diff --git a/sycl/source/virtual_mem.cpp b/sycl/source/virtual_mem.cpp index cb2f3630af5e5..da22b6c05d7fa 100644 --- a/sycl/source/virtual_mem.cpp +++ b/sycl/source/virtual_mem.cpp @@ -24,7 +24,8 @@ namespace ext::oneapi::experimental { __SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, const context &SyclContext, - granularity_mode Mode) { + granularity_mode Mode, + size_t allocationSize) { if (!SyclDevice.has(aspect::ext_oneapi_virtual_mem)) throw sycl::exception( sycl::make_error_code(sycl::errc::feature_not_supported), @@ -45,13 +46,15 @@ __SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, #ifndef NDEBUG size_t InfoOutputSize = 0; Adapter->call( - urCtx, urDevice, GranularityQuery, 0, nullptr, &InfoOutputSize); + urCtx, urDevice, allocationSize, GranularityQuery, 0, nullptr, + &InfoOutputSize); assert(InfoOutputSize == sizeof(size_t) && "Unexpected output size of granularity info query."); #endif // NDEBUG size_t Granularity = 0; Adapter->call( - urCtx, urDevice, GranularityQuery, sizeof(size_t), &Granularity, nullptr); + urCtx, urDevice, allocationSize, GranularityQuery, sizeof(size_t), + &Granularity, nullptr); if (Granularity == 0) throw sycl::exception( sycl::make_error_code(sycl::errc::invalid), @@ -59,6 +62,12 @@ __SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, return Granularity; } +__SYCL_EXPORT size_t get_mem_granularity(const device &SyclDevice, + const context &SyclContext, + granularity_mode Mode) { + return get_mem_granularity(SyclDevice, SyclContext, Mode, 1); +} + __SYCL_EXPORT size_t get_mem_granularity(const context &SyclContext, granularity_mode Mode) { const std::vector Devices = SyclContext.get_devices(); diff --git a/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp index 40ea023b4dd3c..7b076336a9cde 100644 --- a/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp +++ b/sycl/test-e2e/Graph/AsyncAlloc/Inputs/async_alloc_different_sizes.cpp @@ -30,5 +30,5 @@ int main() { asyncAllocWorksWithSize(131); asyncAllocWorksWithSize(10071); asyncAllocWorksWithSize(1007177); - // asyncAllocWorksWithSize(191439360); // BUG + asyncAllocWorksWithSize(191439360); } \ No newline at end of file diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index 577bb4d5b2c89..bad99cc7d8bea 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -4993,6 +4993,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -15324,6 +15326,7 @@ typedef struct ur_loader_init_params_t { typedef struct ur_virtual_mem_granularity_get_info_params_t { ur_context_handle_t *phContext; ur_device_handle_t *phDevice; + size_t *pallocationSize; ur_virtual_mem_granularity_info_t *ppropName; size_t *ppropSize; void **ppPropValue; diff --git a/unified-runtime/include/ur_ddi.h b/unified-runtime/include/ur_ddi.h index cb944b6c395d6..5f58d4c560fe3 100644 --- a/unified-runtime/include/ur_ddi.h +++ b/unified-runtime/include/ur_ddi.h @@ -1834,8 +1834,8 @@ typedef ur_result_t(UR_APICALL *ur_pfnGetUsmP2PExpProcAddrTable_t)( /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urVirtualMemGranularityGetInfo typedef ur_result_t(UR_APICALL *ur_pfnVirtualMemGranularityGetInfo_t)( - ur_context_handle_t, ur_device_handle_t, ur_virtual_mem_granularity_info_t, - size_t, void *, size_t *); + ur_context_handle_t, ur_device_handle_t, size_t, + ur_virtual_mem_granularity_info_t, size_t, void *, size_t *); /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urVirtualMemReserve diff --git a/unified-runtime/include/ur_print.hpp b/unified-runtime/include/ur_print.hpp index 7fc43237a2fbd..c7dc701db3624 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -20319,6 +20319,11 @@ inline std::ostream &operator<<( ur::details::printPtr(os, *(params->phDevice)); + os << ", "; + os << ".allocationSize = "; + + os << *(params->pallocationSize); + os << ", "; os << ".propName = "; diff --git a/unified-runtime/scripts/core/virtual_memory.yml b/unified-runtime/scripts/core/virtual_memory.yml index 61fca47d1b457..6fec73fdf9314 100644 --- a/unified-runtime/scripts/core/virtual_memory.yml +++ b/unified-runtime/scripts/core/virtual_memory.yml @@ -41,6 +41,9 @@ params: [in][optional] is the device to get the granularity from, if the device is null then the granularity is suitable for all devices in context. + - type: size_t + name: allocationSize + desc: "[in] size in bytes of allocation size which granurality we search for." - type: $x_virtual_mem_granularity_info_t name: propName desc: "[in] type of the info to query." diff --git a/unified-runtime/source/adapters/cuda/virtual_mem.cpp b/unified-runtime/source/adapters/cuda/virtual_mem.cpp index 29908ad1d4fd7..38f70e031dbca 100644 --- a/unified-runtime/source/adapters/cuda/virtual_mem.cpp +++ b/unified-runtime/source/adapters/cuda/virtual_mem.cpp @@ -18,6 +18,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( ur_context_handle_t, ur_device_handle_t hDevice, + [[maybe_unused]] size_t allocationSize, ur_virtual_mem_granularity_info_t propName, size_t propSize, void *pPropValue, size_t *pPropSizeRet) { UrReturnHelper ReturnValue(propSize, pPropValue, pPropSizeRet); diff --git a/unified-runtime/source/adapters/hip/virtual_mem.cpp b/unified-runtime/source/adapters/hip/virtual_mem.cpp index 12cf9f838ed9c..1effbbfa06357 100644 --- a/unified-runtime/source/adapters/hip/virtual_mem.cpp +++ b/unified-runtime/source/adapters/hip/virtual_mem.cpp @@ -14,8 +14,8 @@ #include "physical_mem.hpp" UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( - ur_context_handle_t, ur_device_handle_t, ur_virtual_mem_granularity_info_t, - size_t, void *, size_t *) { + ur_context_handle_t, ur_device_handle_t, size_t, + ur_virtual_mem_granularity_info_t, size_t, void *, size_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp index 5e9fad25cbf55..bbbe1fce9690a 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp @@ -165,8 +165,8 @@ ur_result_t urUSMPoolGetInfo(ur_usm_pool_handle_t hPool, void *pPropValue, size_t *pPropSizeRet); ur_result_t urVirtualMemGranularityGetInfo( ur_context_handle_t hContext, ur_device_handle_t hDevice, - ur_virtual_mem_granularity_info_t propName, size_t propSize, - void *pPropValue, size_t *pPropSizeRet); + size_t allocationSize, ur_virtual_mem_granularity_info_t propName, + size_t propSize, void *pPropValue, size_t *pPropSizeRet); ur_result_t urVirtualMemReserve(ur_context_handle_t hContext, const void *pStart, size_t size, void **ppStart); diff --git a/unified-runtime/source/adapters/level_zero/virtual_mem.cpp b/unified-runtime/source/adapters/level_zero/virtual_mem.cpp index f61c8fd43fe2f..0488d2102318c 100644 --- a/unified-runtime/source/adapters/level_zero/virtual_mem.cpp +++ b/unified-runtime/source/adapters/level_zero/virtual_mem.cpp @@ -23,8 +23,8 @@ namespace ur::level_zero { ur_result_t urVirtualMemGranularityGetInfo( ur_context_handle_t hContext, ur_device_handle_t hDevice, - ur_virtual_mem_granularity_info_t propName, size_t propSize, - void *pPropValue, size_t *pPropSizeRet) { + size_t allocationSize, ur_virtual_mem_granularity_info_t propName, + size_t propSize, void *pPropValue, size_t *pPropSizeRet) { UrReturnHelper ReturnValue(propSize, pPropValue, pPropSizeRet); switch (propName) { case UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM: @@ -34,7 +34,8 @@ ur_result_t urVirtualMemGranularityGetInfo( // aligned size. size_t PageSize; ZE2UR_CALL(zeVirtualMemQueryPageSize, - (hContext->getZeHandle(), hDevice->ZeDevice, 1, &PageSize)); + (hContext->getZeHandle(), hDevice->ZeDevice, allocationSize, + &PageSize)); return ReturnValue(PageSize); } default: diff --git a/unified-runtime/source/adapters/mock/ur_mockddi.cpp b/unified-runtime/source/adapters/mock/ur_mockddi.cpp index 0a4a45d0898b2..f4d979248dc17 100644 --- a/unified-runtime/source/adapters/mock/ur_mockddi.cpp +++ b/unified-runtime/source/adapters/mock/ur_mockddi.cpp @@ -2729,6 +2729,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -2744,7 +2746,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( ur_result_t result = UR_RESULT_SUCCESS; ur_virtual_mem_granularity_get_info_params_t params = { - &hContext, &hDevice, &propName, &propSize, &pPropValue, &pPropSizeRet}; + &hContext, &hDevice, &allocationSize, &propName, + &propSize, &pPropValue, &pPropSizeRet}; auto beforeCallback = reinterpret_cast( mock::getCallbacks().get_before_callback( diff --git a/unified-runtime/source/adapters/native_cpu/virtual_mem.cpp b/unified-runtime/source/adapters/native_cpu/virtual_mem.cpp index 131b480ac14b2..6697902564aa9 100644 --- a/unified-runtime/source/adapters/native_cpu/virtual_mem.cpp +++ b/unified-runtime/source/adapters/native_cpu/virtual_mem.cpp @@ -13,8 +13,8 @@ #include "physical_mem.hpp" UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( - ur_context_handle_t, ur_device_handle_t, ur_virtual_mem_granularity_info_t, - size_t, void *, size_t *) { + ur_context_handle_t, ur_device_handle_t, size_t, + ur_virtual_mem_granularity_info_t, size_t, void *, size_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/unified-runtime/source/adapters/opencl/virtual_mem.cpp b/unified-runtime/source/adapters/opencl/virtual_mem.cpp index 7c411d9b7b468..c7db068eca057 100644 --- a/unified-runtime/source/adapters/opencl/virtual_mem.cpp +++ b/unified-runtime/source/adapters/opencl/virtual_mem.cpp @@ -13,8 +13,8 @@ #include "physical_mem.hpp" UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( - ur_context_handle_t, ur_device_handle_t, ur_virtual_mem_granularity_info_t, - size_t, void *, size_t *) { + ur_context_handle_t, ur_device_handle_t, size_t, + ur_virtual_mem_granularity_info_t, size_t, void *, size_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp index 3539a2d2a5267..f8f7c58bf5c3c 100644 --- a/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp +++ b/unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_utils.cpp @@ -272,10 +272,13 @@ size_t GetKernelPrivateMemorySize(ur_kernel_handle_t Kernel, size_t GetVirtualMemGranularity(ur_context_handle_t Context, ur_device_handle_t Device) { size_t Size; + const size_t allocationSize = + 1; // probably we want to use actual allocation size [[maybe_unused]] auto Result = getContext()->urDdiTable.VirtualMem.pfnGranularityGetInfo( - Context, Device, UR_VIRTUAL_MEM_GRANULARITY_INFO_RECOMMENDED, - sizeof(Size), &Size, nullptr); + Context, Device, allocationSize, + UR_VIRTUAL_MEM_GRANULARITY_INFO_RECOMMENDED, sizeof(Size), &Size, + nullptr); assert(Result == UR_RESULT_SUCCESS); return Size; } diff --git a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp index e0d57228e4a61..7dfb39bd5fce4 100644 --- a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp +++ b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp @@ -2236,6 +2236,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -2255,7 +2257,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; ur_virtual_mem_granularity_get_info_params_t params = { - &hContext, &hDevice, &propName, &propSize, &pPropValue, &pPropSizeRet}; + &hContext, &hDevice, &allocationSize, &propName, + &propSize, &pPropValue, &pPropSizeRet}; uint64_t instance = getContext()->notify_begin(UR_FUNCTION_VIRTUAL_MEM_GRANULARITY_GET_INFO, "urVirtualMemGranularityGetInfo", ¶ms); @@ -2263,8 +2266,9 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( auto &logger = getContext()->logger; UR_LOG_L(logger, INFO, " ---> urVirtualMemGranularityGetInfo\n"); - ur_result_t result = pfnGranularityGetInfo( - hContext, hDevice, propName, propSize, pPropValue, pPropSizeRet); + ur_result_t result = + pfnGranularityGetInfo(hContext, hDevice, allocationSize, propName, + propSize, pPropValue, pPropSizeRet); getContext()->notify_end(UR_FUNCTION_VIRTUAL_MEM_GRANULARITY_GET_INFO, "urVirtualMemGranularityGetInfo", ¶ms, &result, diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 979eb3ef22746..2d0d8e357a1e0 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -2182,6 +2182,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -2228,8 +2230,9 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( URLOG_CTX_INVALID_REFERENCE(hDevice); } - ur_result_t result = pfnGranularityGetInfo( - hContext, hDevice, propName, propSize, pPropValue, pPropSizeRet); + ur_result_t result = + pfnGranularityGetInfo(hContext, hDevice, allocationSize, propName, + propSize, pPropValue, pPropSizeRet); return result; } diff --git a/unified-runtime/source/loader/ur_ldrddi.cpp b/unified-runtime/source/loader/ur_ldrddi.cpp index 2ddf9f21c95cf..9eb7fef7196da 100644 --- a/unified-runtime/source/loader/ur_ldrddi.cpp +++ b/unified-runtime/source/loader/ur_ldrddi.cpp @@ -1238,6 +1238,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -1258,8 +1260,8 @@ __urdlllocal ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( return UR_RESULT_ERROR_UNINITIALIZED; // forward to device-platform - return pfnGranularityGetInfo(hContext, hDevice, propName, propSize, - pPropValue, pPropSizeRet); + return pfnGranularityGetInfo(hContext, hDevice, allocationSize, propName, + propSize, pPropValue, pPropSizeRet); } /////////////////////////////////////////////////////////////////////////////// diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index 1261145424445..9bebb5831139d 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -2725,6 +2725,8 @@ ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. @@ -2742,8 +2744,8 @@ ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( if (nullptr == pfnGranularityGetInfo) return UR_RESULT_ERROR_UNINITIALIZED; - return pfnGranularityGetInfo(hContext, hDevice, propName, propSize, - pPropValue, pPropSizeRet); + return pfnGranularityGetInfo(hContext, hDevice, allocationSize, propName, + propSize, pPropValue, pPropSizeRet); } catch (...) { return exceptionToResult(std::current_exception()); } diff --git a/unified-runtime/source/ur_api.cpp b/unified-runtime/source/ur_api.cpp index cc69811f5782b..523a9486f5025 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -2410,6 +2410,8 @@ ur_result_t UR_APICALL urVirtualMemGranularityGetInfo( /// device is null then the granularity is suitable for all devices in /// context. ur_device_handle_t hDevice, + /// [in] size in bytes of allocation size which granurality we search for. + size_t allocationSize, /// [in] type of the info to query. ur_virtual_mem_granularity_info_t propName, /// [in] size in bytes of the memory pointed to by pPropValue. diff --git a/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp b/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp index 327728bb5ace3..fa3eb3f4b5483 100644 --- a/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp +++ b/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp @@ -491,11 +491,11 @@ struct urEnqueueKernelLaunchWithVirtualMemory : uur::urKernelExecutionTest { GTEST_SKIP() << "Virtual memory is not supported."; } + alloc_size = 1024; ASSERT_SUCCESS(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, + context, device, alloc_size, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, sizeof(granularity), &granularity, nullptr)); - alloc_size = 1024; virtual_page_size = uur::RoundUpToNearestFactor(alloc_size, granularity); ASSERT_SUCCESS(urPhysicalMemCreate(context, device, virtual_page_size, diff --git a/unified-runtime/test/conformance/testing/include/uur/fixtures.h b/unified-runtime/test/conformance/testing/include/uur/fixtures.h index b67eddd8f8182..fff0be4a0107e 100644 --- a/unified-runtime/test/conformance/testing/include/uur/fixtures.h +++ b/unified-runtime/test/conformance/testing/include/uur/fixtures.h @@ -976,9 +976,12 @@ struct urVirtualMemGranularityTest : urContextTest { GTEST_SKIP() << "Virtual memory is not supported."; } + const size_t allocationSize = + 1; // assuming allocations in test are small enough and minimal granularity is used ASSERT_SUCCESS(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - sizeof(granularity), &granularity, nullptr)); + context, device, allocationSize, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, sizeof(granularity), + &granularity, nullptr)); } size_t granularity; }; @@ -995,10 +998,12 @@ struct urVirtualMemGranularityTestWithParam : urContextTestWithParam { if (!virtual_memory_support) { GTEST_SKIP() << "Virtual memory is not supported."; } - + const size_t allocationSize = + 1; // assuming allocations in test are small and use smallest granularity ASSERT_SUCCESS(urVirtualMemGranularityGetInfo( - this->context, this->device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - sizeof(granularity), &granularity, nullptr)); + this->context, this->device, allocationSize, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, sizeof(granularity), + &granularity, nullptr)); ASSERT_NE(granularity, 0); } diff --git a/unified-runtime/test/conformance/virtual_memory/urVirtualMemGranularityGetInfo.cpp b/unified-runtime/test/conformance/virtual_memory/urVirtualMemGranularityGetInfo.cpp index 0507b8903a361..d94e7045a0bbe 100644 --- a/unified-runtime/test/conformance/virtual_memory/urVirtualMemGranularityGetInfo.cpp +++ b/unified-runtime/test/conformance/virtual_memory/urVirtualMemGranularityGetInfo.cpp @@ -26,14 +26,14 @@ TEST_P(urVirtualMemGranularityGetInfoTest, SuccessMinimum) { UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM; ASSERT_SUCCESS_OR_OPTIONAL_QUERY( - urVirtualMemGranularityGetInfo(context, device, property_name, 0, nullptr, - &property_size), + urVirtualMemGranularityGetInfo(context, device, 1, property_name, 0, + nullptr, &property_size), property_name); ASSERT_EQ(sizeof(size_t), property_size); size_t property_value = 0; ASSERT_QUERY_RETURNS_VALUE( - urVirtualMemGranularityGetInfo(context, device, property_name, + urVirtualMemGranularityGetInfo(context, device, 1, property_name, property_size, &property_value, nullptr), property_value); @@ -46,14 +46,14 @@ TEST_P(urVirtualMemGranularityGetInfoTest, SuccessRecommended) { UR_VIRTUAL_MEM_GRANULARITY_INFO_RECOMMENDED; ASSERT_SUCCESS_OR_OPTIONAL_QUERY( - urVirtualMemGranularityGetInfo(context, device, property_name, 0, nullptr, - &property_size), + urVirtualMemGranularityGetInfo(context, device, 1, property_name, 0, + nullptr, &property_size), property_name); ASSERT_EQ(sizeof(size_t), property_size); size_t property_value = 0; ASSERT_QUERY_RETURNS_VALUE( - urVirtualMemGranularityGetInfo(context, device, property_name, + urVirtualMemGranularityGetInfo(context, device, 1, property_name, property_size, &property_value, nullptr), property_value); @@ -62,47 +62,52 @@ TEST_P(urVirtualMemGranularityGetInfoTest, SuccessRecommended) { TEST_P(urVirtualMemGranularityGetInfoTest, InvalidNullHandleContext) { size_t property_size = 0; - ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - nullptr, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - 0, nullptr, &property_size), - UR_RESULT_ERROR_INVALID_NULL_HANDLE); + ASSERT_EQ_RESULT( + urVirtualMemGranularityGetInfo(nullptr, device, 1, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, 0, + nullptr, &property_size), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); } TEST_P(urVirtualMemGranularityGetInfoTest, InvalidEnumeration) { size_t property_size = 0; ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - context, device, + context, device, 1, UR_VIRTUAL_MEM_GRANULARITY_INFO_FORCE_UINT32, 0, nullptr, &property_size), UR_RESULT_ERROR_INVALID_ENUMERATION); } TEST_P(urVirtualMemGranularityGetInfoTest, InvalidNullPointerPropSizeRet) { - ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - 0, nullptr, nullptr), - UR_RESULT_ERROR_INVALID_NULL_POINTER); + ASSERT_EQ_RESULT( + urVirtualMemGranularityGetInfo(context, device, 1, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, 0, + nullptr, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); } TEST_P(urVirtualMemGranularityGetInfoTest, InvalidNullPointerPropValue) { - ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - sizeof(size_t), nullptr, nullptr), - UR_RESULT_ERROR_INVALID_NULL_POINTER); + ASSERT_EQ_RESULT( + urVirtualMemGranularityGetInfo(context, device, 1, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, + sizeof(size_t), nullptr, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); } TEST_P(urVirtualMemGranularityGetInfoTest, InvalidPropSizeZero) { size_t minimum = 0; - ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - 0, &minimum, nullptr), - UR_RESULT_ERROR_INVALID_SIZE); + ASSERT_EQ_RESULT( + urVirtualMemGranularityGetInfo(context, device, 1, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, 0, + &minimum, nullptr), + UR_RESULT_ERROR_INVALID_SIZE); } TEST_P(urVirtualMemGranularityGetInfoTest, InvalidSizePropSizeSmall) { size_t minimum = 0; - ASSERT_EQ_RESULT(urVirtualMemGranularityGetInfo( - context, device, UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, - sizeof(size_t) - 1, &minimum, nullptr), - UR_RESULT_ERROR_INVALID_SIZE); + ASSERT_EQ_RESULT( + urVirtualMemGranularityGetInfo(context, device, 1, + UR_VIRTUAL_MEM_GRANULARITY_INFO_MINIMUM, + sizeof(size_t) - 1, &minimum, nullptr), + UR_RESULT_ERROR_INVALID_SIZE); }