diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml index f76326c1c4..6662f7833d 100644 --- a/.github/workflows/cmake.yml +++ b/.github/workflows/cmake.yml @@ -20,7 +20,7 @@ jobs: compiler: [{c: gcc, cxx: g++}] libbacktrace: ['-DVAL_USE_LIBBACKTRACE_BACKTRACE=OFF'] pool_tracking: ['-DUMF_ENABLE_POOL_TRACKING=ON', '-DUMF_ENABLE_POOL_TRACKING=OFF'] - latency_tracking: ['-DUMF_ENABLE_LATENCY_TRACKING=OFF'] + latency_tracking: ['-DUR_ENABLE_LATENCY_HISTOGRAM=OFF'] include: - os: 'ubuntu-22.04' build_type: Release @@ -40,7 +40,7 @@ jobs: - os: 'ubuntu-22.04' build_type: Release compiler: {c: clang, cxx: clang++} - latency_tracking: '-DUMF_ENABLE_LATENCY_TRACKING=ON' + latency_tracking: '-DUR_ENABLE_LATENCY_HISTOGRAM=ON' runs-on: ${{ (matrix.os == 'ubuntu-22.04' && github.repository_owner == 'oneapi-src') && 'intel-ubuntu-22.04' || matrix.os }} steps: diff --git a/include/ur_print.hpp b/include/ur_print.hpp index f71cc12b32..681e8e814d 100644 --- a/include/ur_print.hpp +++ b/include/ur_print.hpp @@ -17403,6 +17403,11 @@ inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const struct return os; } +inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const ur_bool_t value) { + os << (value ? "true" : "false"); + return os; +} + namespace ur::details { /////////////////////////////////////////////////////////////////////////////// // @brief Print pointer value diff --git a/scripts/core/INTRO.rst b/scripts/core/INTRO.rst index 448e3569e2..898d4ce5f3 100644 --- a/scripts/core/INTRO.rst +++ b/scripts/core/INTRO.rst @@ -396,6 +396,14 @@ Specific environment variables can be set to control the behavior of unified run See the Layers_ section for details of the layers currently included in the runtime. +.. envvar:: UR_LOADER_PRELOAD_FILTER + + If set, the loader will read `ONEAPI_DEVICE_SELECTOR` before loading the UR Adapters to determine which backends should be loaded. + + .. note:: + + This environment variable is default enabled on Linux, but default disabled on Windows. + Service identifiers --------------------- diff --git a/scripts/templates/print.hpp.mako b/scripts/templates/print.hpp.mako index 9bf427b889..4180231ea4 100644 --- a/scripts/templates/print.hpp.mako +++ b/scripts/templates/print.hpp.mako @@ -411,6 +411,11 @@ inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const struct %endfor %endfor +inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const ur_bool_t value) { + os << (value ? "true" : "false"); + return os; +} + namespace ${x}::details { /////////////////////////////////////////////////////////////////////////////// // @brief Print pointer value diff --git a/source/adapters/cuda/device.cpp b/source/adapters/cuda/device.cpp index bbaaa27cdb..9c8a0c807c 100644 --- a/source/adapters/cuda/device.cpp +++ b/source/adapters/cuda/device.cpp @@ -57,12 +57,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(4318u); } case UR_DEVICE_INFO_MAX_COMPUTE_UNITS: { - int ComputeUnits = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &ComputeUnits, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, - hDevice->get())); - detail::ur::assertion(ComputeUnits >= 0); - return ReturnValue(static_cast(ComputeUnits)); + return ReturnValue(hDevice->getNumComputeUnits()); } case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: { return ReturnValue(MaxWorkItemDimensions); diff --git a/source/adapters/cuda/device.hpp b/source/adapters/cuda/device.hpp index 0a40329026..3654f2bb36 100644 --- a/source/adapters/cuda/device.hpp +++ b/source/adapters/cuda/device.hpp @@ -32,6 +32,7 @@ struct ur_device_handle_t_ { int MaxCapacityLocalMem{0}; int MaxChosenLocalMem{0}; bool MaxLocalMemSizeChosen{false}; + uint32_t NumComputeUnits{0}; public: ur_device_handle_t_(native_type cuDevice, CUcontext cuContext, CUevent evBase, @@ -54,6 +55,10 @@ struct ur_device_handle_t_ { sizeof(MaxWorkGroupSize), &MaxWorkGroupSize, nullptr)); + UR_CHECK_ERROR(cuDeviceGetAttribute( + reinterpret_cast(&NumComputeUnits), + CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, cuDevice)); + // Set local mem max size if env var is present static const char *LocalMemSizePtrUR = std::getenv("UR_CUDA_MAX_LOCAL_MEM_SIZE"); @@ -107,6 +112,8 @@ struct ur_device_handle_t_ { int getMaxChosenLocalMem() const noexcept { return MaxChosenLocalMem; }; bool maxLocalMemSizeChosen() { return MaxLocalMemSizeChosen; }; + + uint32_t getNumComputeUnits() const noexcept { return NumComputeUnits; }; }; int getAttribute(ur_device_handle_t Device, CUdevice_attribute Attribute); diff --git a/source/adapters/cuda/image.cpp b/source/adapters/cuda/image.cpp index e2960573aa..427fde70e6 100644 --- a/source/adapters/cuda/image.cpp +++ b/source/adapters/cuda/image.cpp @@ -759,13 +759,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp( cpy_desc.dstZ = pCopyRegion->dstOffset.z; cpy_desc.srcMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST; cpy_desc.srcHost = pSrc; - cpy_desc.srcPitch = pCopyRegion->copyExtent.width * PixelSizeBytes; - cpy_desc.srcHeight = pCopyRegion->copyExtent.height; + cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes; + cpy_desc.srcHeight = std::max(uint64_t{1}, pSrcImageDesc->height); cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_ARRAY; cpy_desc.dstArray = (CUarray)pDst; cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width; cpy_desc.Height = std::max(uint64_t{1}, pCopyRegion->copyExtent.height); - cpy_desc.Depth = pDstImageDesc->arraySize; + cpy_desc.Depth = pCopyRegion->copyExtent.depth; UR_CHECK_ERROR(cuMemcpy3DAsync(&cpy_desc, Stream)); } } else if (imageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) { @@ -855,10 +855,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp( cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST; cpy_desc.dstHost = pDst; cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes; - cpy_desc.dstHeight = pDstImageDesc->height; + cpy_desc.dstHeight = std::max(uint64_t{1}, pDstImageDesc->height); cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width; cpy_desc.Height = std::max(uint64_t{1}, pCopyRegion->copyExtent.height); - cpy_desc.Depth = pSrcImageDesc->arraySize; + cpy_desc.Depth = pCopyRegion->copyExtent.depth; UR_CHECK_ERROR(cuMemcpy3DAsync(&cpy_desc, Stream)); } } else { @@ -932,7 +932,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp( cpy_desc.dstArray = (CUarray)pDst; cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width; cpy_desc.Height = std::max(uint64_t{1}, pCopyRegion->copyExtent.height); - cpy_desc.Depth = pSrcImageDesc->arraySize; + cpy_desc.Depth = pCopyRegion->copyExtent.depth; UR_CHECK_ERROR(cuMemcpy3DAsync(&cpy_desc, Stream)); } // Synchronization is required here to handle the case of copying data diff --git a/source/adapters/cuda/kernel.cpp b/source/adapters/cuda/kernel.cpp index d43bd046dc..2061893744 100644 --- a/source/adapters/cuda/kernel.cpp +++ b/source/adapters/cuda/kernel.cpp @@ -167,10 +167,46 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetNativeHandle( UR_APIEXPORT ur_result_t UR_APICALL urKernelSuggestMaxCooperativeGroupCountExp( ur_kernel_handle_t hKernel, size_t localWorkSize, size_t dynamicSharedMemorySize, uint32_t *pGroupCountRet) { - (void)hKernel; - (void)localWorkSize; - (void)dynamicSharedMemorySize; - *pGroupCountRet = 1; + UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_KERNEL); + + // We need to set the active current device for this kernel explicitly here, + // because the occupancy querying API does not take device parameter. + ur_device_handle_t Device = hKernel->getProgram()->getDevice(); + ScopedContext Active(Device); + try { + // We need to calculate max num of work-groups using per-device semantics. + + int MaxNumActiveGroupsPerCU{0}; + UR_CHECK_ERROR(cuOccupancyMaxActiveBlocksPerMultiprocessor( + &MaxNumActiveGroupsPerCU, hKernel->get(), localWorkSize, + dynamicSharedMemorySize)); + detail::ur::assertion(MaxNumActiveGroupsPerCU >= 0); + // Handle the case where we can't have all SMs active with at least 1 group + // per SM. In that case, the device is still able to run 1 work-group, hence + // we will manually check if it is possible with the available HW resources. + if (MaxNumActiveGroupsPerCU == 0) { + size_t MaxWorkGroupSize{}; + urKernelGetGroupInfo( + hKernel, Device, UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE, + sizeof(MaxWorkGroupSize), &MaxWorkGroupSize, nullptr); + size_t MaxLocalSizeBytes{}; + urDeviceGetInfo(Device, UR_DEVICE_INFO_LOCAL_MEM_SIZE, + sizeof(MaxLocalSizeBytes), &MaxLocalSizeBytes, nullptr); + if (localWorkSize > MaxWorkGroupSize || + dynamicSharedMemorySize > MaxLocalSizeBytes || + hasExceededMaxRegistersPerBlock(Device, hKernel, localWorkSize)) + *pGroupCountRet = 0; + else + *pGroupCountRet = 1; + } else { + // Multiply by the number of SMs (CUs = compute units) on the device in + // order to retreive the total number of groups/blocks that can be + // launched. + *pGroupCountRet = Device->getNumComputeUnits() * MaxNumActiveGroupsPerCU; + } + } catch (ur_result_t Err) { + return Err; + } return UR_RESULT_SUCCESS; } diff --git a/source/adapters/level_zero/CMakeLists.txt b/source/adapters/level_zero/CMakeLists.txt index 653bf4d118..cc05d36084 100644 --- a/source/adapters/level_zero/CMakeLists.txt +++ b/source/adapters/level_zero/CMakeLists.txt @@ -118,6 +118,7 @@ if(UR_BUILD_ADAPTER_L0) ${CMAKE_CURRENT_SOURCE_DIR}/queue.hpp ${CMAKE_CURRENT_SOURCE_DIR}/sampler.hpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/kernel_helpers.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/helpers/memory_helpers.hpp ${CMAKE_CURRENT_SOURCE_DIR}/ur_level_zero.cpp ${CMAKE_CURRENT_SOURCE_DIR}/common.cpp ${CMAKE_CURRENT_SOURCE_DIR}/context.cpp @@ -136,6 +137,7 @@ if(UR_BUILD_ADAPTER_L0) ${CMAKE_CURRENT_SOURCE_DIR}/sampler.cpp ${CMAKE_CURRENT_SOURCE_DIR}/image.cpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/kernel_helpers.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/helpers/memory_helpers.cpp ${CMAKE_CURRENT_SOURCE_DIR}/../../ur/ur.cpp ) @@ -199,6 +201,7 @@ if(UR_BUILD_ADAPTER_L0_V2) ${CMAKE_CURRENT_SOURCE_DIR}/platform.hpp ${CMAKE_CURRENT_SOURCE_DIR}/program.hpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/kernel_helpers.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/helpers/memory_helpers.hpp ${CMAKE_CURRENT_SOURCE_DIR}/adapter.cpp ${CMAKE_CURRENT_SOURCE_DIR}/common.cpp ${CMAKE_CURRENT_SOURCE_DIR}/device.cpp @@ -206,6 +209,7 @@ if(UR_BUILD_ADAPTER_L0_V2) ${CMAKE_CURRENT_SOURCE_DIR}/platform.cpp ${CMAKE_CURRENT_SOURCE_DIR}/program.cpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/kernel_helpers.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/helpers/memory_helpers.cpp ${CMAKE_CURRENT_SOURCE_DIR}/../../ur/ur.cpp # v2-only sources ${CMAKE_CURRENT_SOURCE_DIR}/v2/command_list_cache.hpp @@ -217,6 +221,7 @@ if(UR_BUILD_ADAPTER_L0_V2) ${CMAKE_CURRENT_SOURCE_DIR}/v2/event_provider.hpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/event.hpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/kernel.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/v2/memory.hpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_api.hpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_immediate_in_order.hpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/usm.hpp @@ -229,6 +234,7 @@ if(UR_BUILD_ADAPTER_L0_V2) ${CMAKE_CURRENT_SOURCE_DIR}/v2/event_provider_normal.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/event.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/kernel.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/v2/memory.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_api.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_create.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_immediate_in_order.cpp diff --git a/source/adapters/level_zero/context.cpp b/source/adapters/level_zero/context.cpp index de75dc6126..a8eaaa2317 100644 --- a/source/adapters/level_zero/context.cpp +++ b/source/adapters/level_zero/context.cpp @@ -512,7 +512,7 @@ ur_result_t ur_context_handle_t_::getFreeSlotInExistingOrNewPool( // Create one event ZePool per MaxNumEventsPerPool events if (*ZePool == nullptr) { ze_event_pool_counter_based_exp_desc_t counterBasedExt = { - ZE_STRUCTURE_TYPE_COUNTER_BASED_EVENT_POOL_EXP_DESC}; + ZE_STRUCTURE_TYPE_COUNTER_BASED_EVENT_POOL_EXP_DESC, nullptr, 0}; ZeStruct ZeEventPoolDesc; ZeEventPoolDesc.count = MaxNumEventsPerPool; ZeEventPoolDesc.flags = 0; diff --git a/source/adapters/level_zero/event.cpp b/source/adapters/level_zero/event.cpp index 84a7c0b159..01ea6efb91 100644 --- a/source/adapters/level_zero/event.cpp +++ b/source/adapters/level_zero/event.cpp @@ -221,9 +221,8 @@ ur_result_t urEnqueueEventsWaitWithBarrier( return UR_RESULT_SUCCESS; } - ur_event_handle_t InternalEvent; + ur_event_handle_t ResultEvent = nullptr; bool IsInternal = OutEvent == nullptr; - ur_event_handle_t *Event = OutEvent ? OutEvent : &InternalEvent; // For in-order queue and wait-list which is empty or has events from // the same queue just use the last command event as the barrier event. @@ -234,7 +233,10 @@ ur_result_t urEnqueueEventsWaitWithBarrier( EventWaitList) && Queue->LastCommandEvent && !Queue->LastCommandEvent->IsDiscarded) { UR_CALL(ur::level_zero::urEventRetain(Queue->LastCommandEvent)); - *Event = Queue->LastCommandEvent; + ResultEvent = Queue->LastCommandEvent; + if (OutEvent) { + *OutEvent = ResultEvent; + } return UR_RESULT_SUCCESS; } @@ -264,16 +266,21 @@ ur_result_t urEnqueueEventsWaitWithBarrier( EventWaitList, OkToBatch)); // Insert the barrier into the command-list and execute. - UR_CALL(insertBarrierIntoCmdList(CmdList, TmpWaitList, *Event, IsInternal)); + UR_CALL(insertBarrierIntoCmdList(CmdList, TmpWaitList, ResultEvent, + IsInternal)); UR_CALL(Queue->executeCommandList(CmdList, false, OkToBatch)); // Because of the dependency between commands in the in-order queue we don't // need to keep track of any active barriers if we have in-order queue. if (UseMultipleCmdlistBarriers && !Queue->isInOrderQueue()) { - auto UREvent = reinterpret_cast(*Event); + auto UREvent = reinterpret_cast(ResultEvent); Queue->ActiveBarriers.add(UREvent); } + + if (OutEvent) { + *OutEvent = ResultEvent; + } return UR_RESULT_SUCCESS; } @@ -361,14 +368,14 @@ ur_result_t urEnqueueEventsWaitWithBarrier( // Insert a barrier with the events from each command-queue into the // convergence command list. The resulting event signals the convergence of // all barriers. - UR_CALL(insertBarrierIntoCmdList(ConvergenceCmdList, BaseWaitList, *Event, - IsInternal)); + UR_CALL(insertBarrierIntoCmdList(ConvergenceCmdList, BaseWaitList, + ResultEvent, IsInternal)); } else { // If there is only a single queue then insert a barrier and the single // result event can be used as our active barrier and used as the return // event. Take into account whether output event is discarded or not. - UR_CALL(insertBarrierIntoCmdList(CmdLists[0], _ur_ze_event_list_t{}, *Event, - IsInternal)); + UR_CALL(insertBarrierIntoCmdList(CmdLists[0], _ur_ze_event_list_t{}, + ResultEvent, IsInternal)); } // Execute each command list so the barriers can be encountered. @@ -384,8 +391,10 @@ ur_result_t urEnqueueEventsWaitWithBarrier( } UR_CALL(Queue->ActiveBarriers.clear()); - auto UREvent = reinterpret_cast(*Event); - Queue->ActiveBarriers.add(UREvent); + Queue->ActiveBarriers.add(ResultEvent); + if (OutEvent) { + *OutEvent = ResultEvent; + } return UR_RESULT_SUCCESS; } @@ -1508,8 +1517,8 @@ ur_result_t _ur_ze_event_list_t::createAndRetainUrZeEventList( std::shared_lock Lock(EventList[I]->Mutex); - ur_device_handle_t QueueRootDevice; - ur_device_handle_t CurrentQueueRootDevice; + ur_device_handle_t QueueRootDevice = nullptr; + ur_device_handle_t CurrentQueueRootDevice = nullptr; if (Queue) { QueueRootDevice = Queue->Device; CurrentQueueRootDevice = CurQueueDevice; diff --git a/source/adapters/level_zero/helpers/memory_helpers.cpp b/source/adapters/level_zero/helpers/memory_helpers.cpp new file mode 100644 index 0000000000..aea32795ab --- /dev/null +++ b/source/adapters/level_zero/helpers/memory_helpers.cpp @@ -0,0 +1,33 @@ +//===--------- memory_helpers.cpp - Level Zero Adapter -------------------===// +// +// Copyright (C) 2024 Intel Corporation +// +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM +// Exceptions. See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "memory_helpers.hpp" +#include "../common.hpp" + +ze_memory_type_t getMemoryType(ze_context_handle_t hContext, void *ptr) { + // TODO: use UMF once + // https://github.com/oneapi-src/unified-memory-framework/issues/687 is + // implemented + ZeStruct zeMemoryAllocationProperties; + ZE2UR_CALL_THROWS(zeMemGetAllocProperties, + (hContext, ptr, &zeMemoryAllocationProperties, nullptr)); + return zeMemoryAllocationProperties.type; +} + +bool maybeImportUSM(ze_driver_handle_t hTranslatedDriver, + ze_context_handle_t hContext, void *ptr, size_t size) { + if (ZeUSMImport.Enabled && ptr != nullptr && + getMemoryType(hContext, ptr) == ZE_MEMORY_TYPE_UNKNOWN) { + // Promote the host ptr to USM host memory + ZeUSMImport.doZeUSMImport(hTranslatedDriver, ptr, size); + return true; + } + return false; +} diff --git a/source/adapters/level_zero/helpers/memory_helpers.hpp b/source/adapters/level_zero/helpers/memory_helpers.hpp new file mode 100644 index 0000000000..ad50be992c --- /dev/null +++ b/source/adapters/level_zero/helpers/memory_helpers.hpp @@ -0,0 +1,23 @@ +//===--------- memory_helpers.hpp - Level Zero Adapter -------------------===// +// +// Copyright (C) 2024 Intel Corporation +// +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM +// Exceptions. See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +#pragma once + +#include +#include + +// If USM Import feature is enabled and hostptr is supplied, +// import the hostptr if not already imported into USM. +// Data transfer rate is maximized when both source and destination +// are USM pointers. Promotion of the host pointer to USM thus +// optimizes data transfer performance. +bool maybeImportUSM(ze_driver_handle_t hTranslatedDriver, + ze_context_handle_t hContext, void *ptr, size_t size); + +ze_memory_type_t getMemoryType(ze_context_handle_t hContext, void *ptr); diff --git a/source/adapters/level_zero/image.cpp b/source/adapters/level_zero/image.cpp index fc623e7e74..a717597623 100644 --- a/source/adapters/level_zero/image.cpp +++ b/source/adapters/level_zero/image.cpp @@ -16,6 +16,7 @@ #include "sampler.hpp" #include "ur_interface_loader.hpp" #include "ur_level_zero.hpp" +#include "ze_api.h" typedef ze_result_t(ZE_APICALL *zeImageGetDeviceOffsetExp_pfn)( ze_image_handle_t hImage, uint64_t *pDeviceOffset); @@ -445,7 +446,8 @@ ur_result_t bindlessImagesCreateImpl(ur_context_handle_t hContext, ze_image_handle_t ZeImage; ze_memory_allocation_properties_t MemAllocProperties{ - ZE_STRUCTURE_TYPE_MEMORY_ALLOCATION_PROPERTIES}; + ZE_STRUCTURE_TYPE_MEMORY_ALLOCATION_PROPERTIES, nullptr, + ZE_MEMORY_TYPE_UNKNOWN, 0, 0}; ZE2UR_CALL(zeMemGetAllocProperties, (hContext->ZeContext, reinterpret_cast(hImageMem), &MemAllocProperties, nullptr)); diff --git a/source/adapters/level_zero/memory.cpp b/source/adapters/level_zero/memory.cpp index 9786092073..69edf83a78 100644 --- a/source/adapters/level_zero/memory.cpp +++ b/source/adapters/level_zero/memory.cpp @@ -15,6 +15,7 @@ #include "context.hpp" #include "event.hpp" +#include "helpers/memory_helpers.hpp" #include "image.hpp" #include "logger/ur_logger.hpp" #include "queue.hpp" @@ -1599,30 +1600,11 @@ ur_result_t urMemBufferCreate( Host = Properties->pHost; } - // If USM Import feature is enabled and hostptr is supplied, - // import the hostptr if not already imported into USM. - // Data transfer rate is maximized when both source and destination - // are USM pointers. Promotion of the host pointer to USM thus - // optimizes data transfer performance. bool HostPtrImported = false; - if (ZeUSMImport.Enabled && Host != nullptr && - (Flags & UR_MEM_FLAG_USE_HOST_POINTER) != 0) { - // Query memory type of the host pointer - ze_device_handle_t ZeDeviceHandle; - ZeStruct ZeMemoryAllocationProperties; - ZE2UR_CALL(zeMemGetAllocProperties, - (Context->ZeContext, Host, &ZeMemoryAllocationProperties, - &ZeDeviceHandle)); - - // If not shared of any type, we can import the ptr - if (ZeMemoryAllocationProperties.type == ZE_MEMORY_TYPE_UNKNOWN) { - // Promote the host ptr to USM host memory - ze_driver_handle_t driverHandle = - Context->getPlatform()->ZeDriverHandleExpTranslated; - ZeUSMImport.doZeUSMImport(driverHandle, Host, Size); - HostPtrImported = true; - } - } + if (Flags & UR_MEM_FLAG_USE_HOST_POINTER) + HostPtrImported = + maybeImportUSM(Context->getPlatform()->ZeDriverHandleExpTranslated, + Context->ZeContext, Host, Size); _ur_buffer *Buffer = nullptr; auto HostPtrOrNull = (Flags & UR_MEM_FLAG_USE_HOST_POINTER) diff --git a/source/adapters/level_zero/v2/api.cpp b/source/adapters/level_zero/v2/api.cpp index eba7359379..cd25f838fe 100644 --- a/source/adapters/level_zero/v2/api.cpp +++ b/source/adapters/level_zero/v2/api.cpp @@ -49,32 +49,6 @@ ur_result_t urMemImageCreate(ur_context_handle_t hContext, ur_mem_flags_t flags, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t urMemBufferCreate(ur_context_handle_t hContext, - ur_mem_flags_t flags, size_t size, - const ur_buffer_properties_t *pProperties, - ur_mem_handle_t *phBuffer) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -ur_result_t urMemRetain(ur_mem_handle_t hMem) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -ur_result_t urMemRelease(ur_mem_handle_t hMem) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -ur_result_t urMemBufferPartition(ur_mem_handle_t hBuffer, ur_mem_flags_t flags, - ur_buffer_create_type_t bufferCreateType, - const ur_buffer_region_t *pRegion, - ur_mem_handle_t *phMem) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urMemGetNativeHandle(ur_mem_handle_t hMem, ur_device_handle_t hDevice, ur_native_handle_t *phNativeMem) { @@ -82,13 +56,6 @@ ur_result_t urMemGetNativeHandle(ur_mem_handle_t hMem, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t urMemBufferCreateWithNativeHandle( - ur_native_handle_t hNativeMem, ur_context_handle_t hContext, - const ur_mem_native_properties_t *pProperties, ur_mem_handle_t *phMem) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urMemImageCreateWithNativeHandle( ur_native_handle_t hNativeMem, ur_context_handle_t hContext, const ur_image_format_t *pImageFormat, const ur_image_desc_t *pImageDesc, @@ -217,14 +184,6 @@ ur_result_t urPhysicalMemRelease(ur_physical_mem_handle_t hPhysicalMem) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t -urKernelSetArgLocal(ur_kernel_handle_t hKernel, uint32_t argIndex, - size_t argSize, - const ur_kernel_arg_local_properties_t *pProperties) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urKernelGetInfo(ur_kernel_handle_t hKernel, ur_kernel_info_t propName, size_t propSize, void *pPropValue, size_t *pPropSizeRet) { @@ -232,33 +191,6 @@ ur_result_t urKernelGetInfo(ur_kernel_handle_t hKernel, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t urKernelGetGroupInfo(ur_kernel_handle_t hKernel, - ur_device_handle_t hDevice, - ur_kernel_group_info_t propName, - size_t propSize, void *pPropValue, - size_t *pPropSizeRet) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -ur_result_t urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, - ur_device_handle_t hDevice, - ur_kernel_sub_group_info_t propName, - size_t propSize, void *pPropValue, - size_t *pPropSizeRet) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -ur_result_t -urKernelSetExecInfo(ur_kernel_handle_t hKernel, ur_kernel_exec_info_t propName, - size_t propSize, - const ur_kernel_exec_info_properties_t *pProperties, - const void *pPropValue) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urKernelSetArgSampler(ur_kernel_handle_t hKernel, uint32_t argIndex, const ur_kernel_arg_sampler_properties_t *pProperties, @@ -267,14 +199,6 @@ urKernelSetArgSampler(ur_kernel_handle_t hKernel, uint32_t argIndex, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t -urKernelSetArgMemObj(ur_kernel_handle_t hKernel, uint32_t argIndex, - const ur_kernel_arg_mem_obj_properties_t *pProperties, - ur_mem_handle_t hArgValue) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urKernelSetSpecializationConstants( ur_kernel_handle_t hKernel, uint32_t count, const ur_specialization_constant_info_t *pSpecConstants) { @@ -308,13 +232,6 @@ ur_result_t urKernelGetSuggestedLocalWorkSize(ur_kernel_handle_t hKernel, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t urEventGetInfo(ur_event_handle_t hEvent, ur_event_info_t propName, - size_t propSize, void *pPropValue, - size_t *pPropSizeRet) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urEventGetProfilingInfo(ur_event_handle_t hEvent, ur_profiling_info_t propName, size_t propSize, void *pPropValue, @@ -323,12 +240,6 @@ ur_result_t urEventGetProfilingInfo(ur_event_handle_t hEvent, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -ur_result_t urEventWait(uint32_t numEvents, - const ur_event_handle_t *phEventWaitList) { - logger::error("{} function not implemented!", __FUNCTION__); - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - ur_result_t urEventGetNativeHandle(ur_event_handle_t hEvent, ur_native_handle_t *phNativeEvent) { logger::error("{} function not implemented!", __FUNCTION__); diff --git a/source/adapters/level_zero/v2/command_list_cache.cpp b/source/adapters/level_zero/v2/command_list_cache.cpp index eee6555f87..651cb5944a 100644 --- a/source/adapters/level_zero/v2/command_list_cache.cpp +++ b/source/adapters/level_zero/v2/command_list_cache.cpp @@ -43,7 +43,7 @@ inline size_t command_list_descriptor_hash_t::operator()( command_list_cache_t::command_list_cache_t(ze_context_handle_t ZeContext) : ZeContext{ZeContext} {} -raii::ze_command_list_t +raii::ze_command_list_handle_t command_list_cache_t::createCommandList(const command_list_descriptor_t &desc) { if (auto ImmCmdDesc = std::get_if(&desc)) { @@ -61,7 +61,7 @@ command_list_cache_t::createCommandList(const command_list_descriptor_t &desc) { ZE2UR_CALL_THROWS( zeCommandListCreateImmediate, (ZeContext, ImmCmdDesc->ZeDevice, &QueueDesc, &ZeCommandList)); - return raii::ze_command_list_t(ZeCommandList, &zeCommandListDestroy); + return raii::ze_command_list_handle_t(ZeCommandList); } else { auto RegCmdDesc = std::get(desc); ZeStruct CmdListDesc; @@ -72,7 +72,7 @@ command_list_cache_t::createCommandList(const command_list_descriptor_t &desc) { ze_command_list_handle_t ZeCommandList; ZE2UR_CALL_THROWS(zeCommandListCreate, (ZeContext, RegCmdDesc.ZeDevice, &CmdListDesc, &ZeCommandList)); - return raii::ze_command_list_t(ZeCommandList, &zeCommandListDestroy); + return raii::ze_command_list_handle_t(ZeCommandList); } } @@ -94,8 +94,7 @@ command_list_cache_t::getImmediateCommandList( auto CommandList = getCommandList(Desc).release(); return raii::cache_borrowed_command_list_t( CommandList, [Cache = this, Desc](ze_command_list_handle_t CmdList) { - Cache->addCommandList( - Desc, raii::ze_command_list_t(CmdList, &zeCommandListDestroy)); + Cache->addCommandList(Desc, raii::ze_command_list_handle_t(CmdList)); }); } @@ -113,12 +112,11 @@ command_list_cache_t::getRegularCommandList(ze_device_handle_t ZeDevice, return raii::cache_borrowed_command_list_t( CommandList, [Cache = this, Desc](ze_command_list_handle_t CmdList) { - Cache->addCommandList( - Desc, raii::ze_command_list_t(CmdList, &zeCommandListDestroy)); + Cache->addCommandList(Desc, raii::ze_command_list_handle_t(CmdList)); }); } -raii::ze_command_list_t +raii::ze_command_list_handle_t command_list_cache_t::getCommandList(const command_list_descriptor_t &desc) { std::unique_lock Lock(ZeCommandListCacheMutex); auto it = ZeCommandListCache.find(desc); @@ -129,7 +127,8 @@ command_list_cache_t::getCommandList(const command_list_descriptor_t &desc) { assert(!it->second.empty()); - raii::ze_command_list_t CommandListHandle = std::move(it->second.top()); + raii::ze_command_list_handle_t CommandListHandle = + std::move(it->second.top()); it->second.pop(); if (it->second.empty()) @@ -138,8 +137,9 @@ command_list_cache_t::getCommandList(const command_list_descriptor_t &desc) { return CommandListHandle; } -void command_list_cache_t::addCommandList(const command_list_descriptor_t &desc, - raii::ze_command_list_t cmdList) { +void command_list_cache_t::addCommandList( + const command_list_descriptor_t &desc, + raii::ze_command_list_handle_t cmdList) { // TODO: add a limit? std::unique_lock Lock(ZeCommandListCacheMutex); auto [it, _] = ZeCommandListCache.try_emplace(desc); diff --git a/source/adapters/level_zero/v2/command_list_cache.hpp b/source/adapters/level_zero/v2/command_list_cache.hpp index bb32e0e64a..9884e16dc4 100644 --- a/source/adapters/level_zero/v2/command_list_cache.hpp +++ b/source/adapters/level_zero/v2/command_list_cache.hpp @@ -17,15 +17,13 @@ #include #include -#include "../common.hpp" +#include "common.hpp" namespace v2 { namespace raii { -using ze_command_list_t = std::unique_ptr<::_ze_command_list_handle_t, - decltype(&zeCommandListDestroy)>; using cache_borrowed_command_list_t = std::unique_ptr<::_ze_command_list_handle_t, - std::function>; + std::function>; } // namespace raii struct immediate_command_list_descriptor_t { @@ -72,15 +70,16 @@ struct command_list_cache_t { private: ze_context_handle_t ZeContext; std::unordered_map, + std::stack, command_list_descriptor_hash_t> ZeCommandListCache; ur_mutex ZeCommandListCacheMutex; - raii::ze_command_list_t getCommandList(const command_list_descriptor_t &desc); + raii::ze_command_list_handle_t + getCommandList(const command_list_descriptor_t &desc); void addCommandList(const command_list_descriptor_t &desc, - raii::ze_command_list_t cmdList); - raii::ze_command_list_t + raii::ze_command_list_handle_t cmdList); + raii::ze_command_list_handle_t createCommandList(const command_list_descriptor_t &desc); }; } // namespace v2 diff --git a/source/adapters/level_zero/v2/common.hpp b/source/adapters/level_zero/v2/common.hpp index ffef317ae8..4fb851bad8 100644 --- a/source/adapters/level_zero/v2/common.hpp +++ b/source/adapters/level_zero/v2/common.hpp @@ -54,6 +54,8 @@ struct ze_handle_wrapper { try { reset(); } catch (...) { + // TODO: add appropriate logging or pass the error + // to the caller (make the dtor noexcept(false) or use tls?) } } @@ -94,5 +96,11 @@ using ze_event_handle_t = using ze_event_pool_handle_t = ze_handle_wrapper<::ze_event_pool_handle_t, zeEventPoolDestroy>; +using ze_context_handle_t = + ze_handle_wrapper<::ze_context_handle_t, zeContextDestroy>; + +using ze_command_list_handle_t = + ze_handle_wrapper<::ze_command_list_handle_t, zeCommandListDestroy>; + } // namespace raii } // namespace v2 diff --git a/source/adapters/level_zero/v2/context.cpp b/source/adapters/level_zero/v2/context.cpp index 84e3d96b88..abb8a13538 100644 --- a/source/adapters/level_zero/v2/context.cpp +++ b/source/adapters/level_zero/v2/context.cpp @@ -17,8 +17,8 @@ ur_context_handle_t_::ur_context_handle_t_(ze_context_handle_t hContext, uint32_t numDevices, const ur_device_handle_t *phDevices, bool ownZeContext) - : hContext(hContext), hDevices(phDevices, phDevices + numDevices), - commandListCache(hContext), + : hContext(hContext, ownZeContext), + hDevices(phDevices, phDevices + numDevices), commandListCache(hContext), eventPoolCache(phDevices[0]->Platform->getNumDevices(), [context = this, platform = phDevices[0]->Platform](DeviceId deviceId) { @@ -27,19 +27,7 @@ ur_context_handle_t_::ur_context_handle_t_(ze_context_handle_t hContext, return std::make_unique( context, device, v2::EVENT_COUNTER, v2::QUEUE_IMMEDIATE); - }) { - std::ignore = ownZeContext; -} - -ur_context_handle_t_::~ur_context_handle_t_() noexcept(false) { - // ur_context_handle_t_ is only created/destroyed through urContextCreate - // and urContextRelease so it's safe to throw here - ZE2UR_CALL_THROWS(zeContextDestroy, (hContext)); -} - -ze_context_handle_t ur_context_handle_t_::getZeHandle() const { - return hContext; -} + }) {} ur_result_t ur_context_handle_t_::retain() { RefCount.increment(); @@ -115,6 +103,12 @@ ur_result_t urContextGetInfo(ur_context_handle_t hContext, return ReturnValue(uint32_t(hContext->getDevices().size())); case UR_CONTEXT_INFO_REFERENCE_COUNT: return ReturnValue(uint32_t{hContext->RefCount.load()}); + case UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT: + // TODO: this is currently not implemented + return ReturnValue(uint8_t{false}); + case UR_CONTEXT_INFO_USM_FILL2D_SUPPORT: + // 2D USM fill is not supported. + return ReturnValue(uint8_t{false}); default: return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/level_zero/v2/context.hpp b/source/adapters/level_zero/v2/context.hpp index 69bf406594..0ed701400d 100644 --- a/source/adapters/level_zero/v2/context.hpp +++ b/source/adapters/level_zero/v2/context.hpp @@ -13,17 +13,17 @@ #include #include "command_list_cache.hpp" +#include "common.hpp" #include "event_pool_cache.hpp" struct ur_context_handle_t_ : _ur_object { ur_context_handle_t_(ze_context_handle_t hContext, uint32_t numDevices, const ur_device_handle_t *phDevices, bool ownZeContext); - ~ur_context_handle_t_() noexcept(false); ur_result_t retain(); ur_result_t release(); - ze_context_handle_t getZeHandle() const; + inline ze_context_handle_t getZeHandle() const { return hContext.get(); } ur_platform_handle_t getPlatform() const; const std::vector &getDevices() const; @@ -31,7 +31,7 @@ struct ur_context_handle_t_ : _ur_object { // For that the Device or its root devices need to be in the context. bool isValidDevice(ur_device_handle_t Device) const; - const ze_context_handle_t hContext; + const v2::raii::ze_context_handle_t hContext; const std::vector hDevices; v2::command_list_cache_t commandListCache; v2::event_pool_cache eventPoolCache; diff --git a/source/adapters/level_zero/v2/event.cpp b/source/adapters/level_zero/v2/event.cpp index 8654e0b25a..df99c83b53 100644 --- a/source/adapters/level_zero/v2/event.cpp +++ b/source/adapters/level_zero/v2/event.cpp @@ -51,4 +51,41 @@ ur_result_t urEventRetain(ur_event_handle_t hEvent) { return hEvent->retain(); } ur_result_t urEventRelease(ur_event_handle_t hEvent) { return hEvent->release(); } + +ur_result_t urEventWait(uint32_t numEvents, + const ur_event_handle_t *phEventWaitList) { + for (uint32_t i = 0; i < numEvents; ++i) { + ZE2UR_CALL(zeEventHostSynchronize, + (phEventWaitList[i]->getZeEvent(), UINT64_MAX)); + } + return UR_RESULT_SUCCESS; +} + +ur_result_t urEventGetInfo(ur_event_handle_t hEvent, ur_event_info_t propName, + size_t propValueSize, void *pPropValue, + size_t *pPropValueSizeRet) { + UrReturnHelper returnValue(propValueSize, pPropValue, pPropValueSizeRet); + + switch (propName) { + case UR_EVENT_INFO_COMMAND_EXECUTION_STATUS: { + auto zeStatus = ZE_CALL_NOCHECK(zeEventQueryStatus, (hEvent->getZeEvent())); + + if (zeStatus == ZE_RESULT_NOT_READY) { + return returnValue(UR_EVENT_STATUS_SUBMITTED); + } else { + return returnValue(UR_EVENT_STATUS_COMPLETE); + } + } + case UR_EVENT_INFO_REFERENCE_COUNT: { + return returnValue(hEvent->RefCount.load()); + } + default: + logger::error( + "Unsupported ParamName in urEventGetInfo: ParamName=ParamName={}(0x{})", + propName, logger::toHex(propName)); + return UR_RESULT_ERROR_INVALID_VALUE; + } + + return UR_RESULT_SUCCESS; +} } // namespace ur::level_zero diff --git a/source/adapters/level_zero/v2/event_provider_counter.cpp b/source/adapters/level_zero/v2/event_provider_counter.cpp index 5334b2f888..76caea4c58 100644 --- a/source/adapters/level_zero/v2/event_provider_counter.cpp +++ b/source/adapters/level_zero/v2/event_provider_counter.cpp @@ -27,9 +27,9 @@ provider_counter::provider_counter(ur_platform_handle_t platform, ZE2UR_CALL_THROWS(zeDriverGetExtensionFunctionAddress, (platform->ZeDriver, "zexCounterBasedEventCreate", (void **)&this->eventCreateFunc)); - ZE2UR_CALL_THROWS( - zelLoaderTranslateHandle, - (ZEL_HANDLE_CONTEXT, context->hContext, (void **)&translatedContext)); + ZE2UR_CALL_THROWS(zelLoaderTranslateHandle, + (ZEL_HANDLE_CONTEXT, context->getZeHandle(), + (void **)&translatedContext)); ZE2UR_CALL_THROWS( zelLoaderTranslateHandle, (ZEL_HANDLE_DEVICE, device->ZeDevice, (void **)&translatedDevice)); @@ -39,7 +39,7 @@ event_allocation provider_counter::allocate() { if (freelist.empty()) { ZeStruct desc; desc.index = 0; - desc.signal = 0; + desc.signal = ZE_EVENT_SCOPE_FLAG_HOST; desc.wait = 0; ze_event_handle_t handle; diff --git a/source/adapters/level_zero/v2/event_provider_normal.cpp b/source/adapters/level_zero/v2/event_provider_normal.cpp index f5a1c940c6..4df05c12ed 100644 --- a/source/adapters/level_zero/v2/event_provider_normal.cpp +++ b/source/adapters/level_zero/v2/event_provider_normal.cpp @@ -32,7 +32,7 @@ provider_pool::provider_pool(ur_context_handle_t context, desc.flags = ZE_EVENT_POOL_FLAG_HOST_VISIBLE; ze_event_pool_counter_based_exp_desc_t counterBasedExt = { - ZE_STRUCTURE_TYPE_COUNTER_BASED_EVENT_POOL_EXP_DESC, nullptr}; + ZE_STRUCTURE_TYPE_COUNTER_BASED_EVENT_POOL_EXP_DESC, nullptr, 0}; if (events == event_type::EVENT_COUNTER) { counterBasedExt.flags = @@ -43,7 +43,7 @@ provider_pool::provider_pool(ur_context_handle_t context, } ZE2UR_CALL_THROWS(zeEventPoolCreate, - (context->hContext, &desc, 1, + (context->getZeHandle(), &desc, 1, const_cast(&device->ZeDevice), pool.ptr())); @@ -51,7 +51,7 @@ provider_pool::provider_pool(ur_context_handle_t context, for (int i = 0; i < EVENTS_BURST; ++i) { ZeStruct desc; desc.index = i; - desc.signal = 0; + desc.signal = ZE_EVENT_SCOPE_FLAG_HOST; desc.wait = 0; ZE2UR_CALL_THROWS(zeEventCreate, (pool.get(), &desc, freelist[i].ptr())); } diff --git a/source/adapters/level_zero/v2/kernel.cpp b/source/adapters/level_zero/v2/kernel.cpp index e6a37af814..8bfad2d2ad 100644 --- a/source/adapters/level_zero/v2/kernel.cpp +++ b/source/adapters/level_zero/v2/kernel.cpp @@ -12,13 +12,14 @@ #include "context.hpp" #include "kernel.hpp" +#include "memory.hpp" #include "../device.hpp" #include "../platform.hpp" #include "../program.hpp" #include "../ur_interface_loader.hpp" -ur_single_device_kernel_t::ur_single_device_kernel_t(ze_device_handle_t hDevice, +ur_single_device_kernel_t::ur_single_device_kernel_t(ur_device_handle_t hDevice, ze_kernel_handle_t hKernel, bool ownZeHandle) : hDevice(hDevice), hKernel(hKernel, ownZeHandle) { @@ -54,7 +55,7 @@ ur_kernel_handle_t_::ur_kernel_handle_t_(ur_program_handle_t hProgram, assert(urDevice != hProgram->Context->getDevices().end()); auto deviceId = (*urDevice)->Id.value(); - deviceKernels[deviceId].emplace(zeDevice, zeKernel, true); + deviceKernels[deviceId].emplace(*urDevice, zeKernel, true); } completeInitialization(); } @@ -118,7 +119,7 @@ ur_kernel_handle_t_::getZeHandle(ur_device_handle_t hDevice) { auto &kernel = deviceKernels[0].value(); // hDevice is nullptr for native handle - if ((kernel.hDevice != nullptr && kernel.hDevice != hDevice->ZeDevice)) { + if ((kernel.hDevice != nullptr && kernel.hDevice != hDevice)) { throw UR_RESULT_ERROR_INVALID_DEVICE; } @@ -197,6 +198,58 @@ ur_program_handle_t ur_kernel_handle_t_::getProgramHandle() const { return hProgram; } +ur_result_t ur_kernel_handle_t_::setExecInfo(ur_kernel_exec_info_t propName, + const void *pPropValue) { + std::scoped_lock Guard(Mutex); + + for (auto &kernel : deviceKernels) { + if (!kernel.has_value()) + continue; + if (propName == UR_KERNEL_EXEC_INFO_USM_INDIRECT_ACCESS && + *(static_cast(pPropValue)) == true) { + // The whole point for users really was to not need to know anything + // about the types of allocations kernel uses. So in DPC++ we always + // just set all 3 modes for each kernel. + ze_kernel_indirect_access_flags_t indirectFlags = + ZE_KERNEL_INDIRECT_ACCESS_FLAG_HOST | + ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE | + ZE_KERNEL_INDIRECT_ACCESS_FLAG_SHARED; + ZE2UR_CALL(zeKernelSetIndirectAccess, + (kernel->hKernel.get(), indirectFlags)); + } else if (propName == UR_KERNEL_EXEC_INFO_CACHE_CONFIG) { + ze_cache_config_flag_t zeCacheConfig{}; + auto cacheConfig = + *(static_cast(pPropValue)); + if (cacheConfig == UR_KERNEL_CACHE_CONFIG_LARGE_SLM) + zeCacheConfig = ZE_CACHE_CONFIG_FLAG_LARGE_SLM; + else if (cacheConfig == UR_KERNEL_CACHE_CONFIG_LARGE_DATA) + zeCacheConfig = ZE_CACHE_CONFIG_FLAG_LARGE_DATA; + else if (cacheConfig == UR_KERNEL_CACHE_CONFIG_DEFAULT) + zeCacheConfig = static_cast(0); + else + // Unexpected cache configuration value. + return UR_RESULT_ERROR_INVALID_VALUE; + ZE2UR_CALL(zeKernelSetCacheConfig, + (kernel->hKernel.get(), zeCacheConfig);); + } else { + logger::error("urKernelSetExecInfo: unsupported ParamName"); + return UR_RESULT_ERROR_INVALID_VALUE; + } + } + + return UR_RESULT_SUCCESS; +} + +std::vector ur_kernel_handle_t_::getDevices() const { + std::vector devices; + for (size_t i = 0; i < deviceKernels.size(); ++i) { + if (deviceKernels[i].has_value()) { + devices.push_back(deviceKernels[i].value().hDevice); + } + } + return devices; +} + namespace ur::level_zero { ur_result_t urKernelCreate(ur_program_handle_t hProgram, const char *pKernelName, @@ -248,4 +301,162 @@ ur_result_t urKernelSetArgPointer( TRACK_SCOPE_LATENCY("ur_kernel_handle_t_::setArgPointer"); return hKernel->setArgPointer(argIndex, pProperties, pArgValue); } + +ur_result_t +urKernelSetArgMemObj(ur_kernel_handle_t hKernel, uint32_t argIndex, + const ur_kernel_arg_mem_obj_properties_t *pProperties, + ur_mem_handle_t hArgValue) { + TRACK_SCOPE_LATENCY("ur_kernel_handle_t_::setArgMemObj"); + + // TODO: support properties + std::ignore = pProperties; + + auto kernelDevices = hKernel->getDevices(); + if (kernelDevices.size() == 1) { + auto zePtr = hArgValue->getPtr(kernelDevices.front()); + return hKernel->setArgPointer(argIndex, nullptr, zePtr); + } else { + // TODO: Implement this for multi-device kernels. + // Do this the same way as in legacy (keep a pending Args vector and + // do actual allocation on kernel submission) or allocate the memory + // immediately (only for small allocations?) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } +} + +ur_result_t +urKernelSetArgLocal(ur_kernel_handle_t hKernel, uint32_t argIndex, + size_t argSize, + const ur_kernel_arg_local_properties_t *pProperties) { + TRACK_SCOPE_LATENCY("ur_kernel_handle_t_::setArgLocal"); + + std::ignore = pProperties; + + return hKernel->setArgValue(argIndex, argSize, nullptr, nullptr); +} + +ur_result_t urKernelSetExecInfo( + ur_kernel_handle_t hKernel, ///< [in] handle of the kernel object + ur_kernel_exec_info_t propName, ///< [in] name of the execution attribute + size_t propSize, ///< [in] size in byte the attribute value + const ur_kernel_exec_info_properties_t + *pProperties, ///< [in][optional] pointer to execution info properties + const void *pPropValue ///< [in][range(0, propSize)] pointer to memory + ///< location holding the property value. +) { + std::ignore = propSize; + std::ignore = pProperties; + + return hKernel->setExecInfo(propName, pPropValue); +} + +ur_result_t urKernelGetGroupInfo( + ur_kernel_handle_t hKernel, ///< [in] handle of the Kernel object + ur_device_handle_t hDevice, ///< [in] handle of the Device object + ur_kernel_group_info_t + paramName, ///< [in] name of the work Group property to query + size_t + paramValueSize, ///< [in] size of the Kernel Work Group property value + void *pParamValue, ///< [in,out][optional][range(0, propSize)] value of the + ///< Kernel Work Group property. + size_t *pParamValueSizeRet ///< [out][optional] pointer to the actual size + ///< in bytes of data being queried by propName. +) { + UrReturnHelper returnValue(paramValueSize, pParamValue, pParamValueSizeRet); + + std::shared_lock Guard(hKernel->Mutex); + switch (paramName) { + case UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE: { + // TODO: To revisit after level_zero/issues/262 is resolved + struct { + size_t Arr[3]; + } GlobalWorkSize = {{(hDevice->ZeDeviceComputeProperties->maxGroupSizeX * + hDevice->ZeDeviceComputeProperties->maxGroupCountX), + (hDevice->ZeDeviceComputeProperties->maxGroupSizeY * + hDevice->ZeDeviceComputeProperties->maxGroupCountY), + (hDevice->ZeDeviceComputeProperties->maxGroupSizeZ * + hDevice->ZeDeviceComputeProperties->maxGroupCountZ)}}; + return returnValue(GlobalWorkSize); + } + case UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE: { + ZeStruct workGroupProperties; + workGroupProperties.maxGroupSize = 0; + + ZeStruct kernelProperties; + kernelProperties.pNext = &workGroupProperties; + + auto zeDevice = hKernel->getZeHandle(hDevice); + if (zeDevice) { + auto zeResult = + ZE_CALL_NOCHECK(zeKernelGetProperties, (zeDevice, &kernelProperties)); + if (zeResult == ZE_RESULT_SUCCESS && + workGroupProperties.maxGroupSize != 0) { + return returnValue(workGroupProperties.maxGroupSize); + } + return returnValue( + uint64_t{hDevice->ZeDeviceComputeProperties->maxTotalGroupSize}); + } + } + case UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE: { + auto props = hKernel->getProperties(hDevice); + struct { + size_t Arr[3]; + } WgSize = {{props.requiredGroupSizeX, props.requiredGroupSizeY, + props.requiredGroupSizeZ}}; + return returnValue(WgSize); + } + case UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE: { + auto props = hKernel->getProperties(hDevice); + return returnValue(uint32_t{props.localMemSize}); + } + case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: { + return returnValue( + size_t{hDevice->ZeDeviceProperties->physicalEUSimdWidth}); + } + case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: { + auto props = hKernel->getProperties(hDevice); + return returnValue(uint32_t{props.privateMemSize}); + } + default: { + logger::error( + "Unknown ParamName in urKernelGetGroupInfo: ParamName={}(0x{})", + paramName, logger::toHex(paramName)); + return UR_RESULT_ERROR_INVALID_VALUE; + } + } + return UR_RESULT_SUCCESS; +} + +ur_result_t urKernelGetSubGroupInfo( + ur_kernel_handle_t hKernel, ///< [in] handle of the Kernel object + ur_device_handle_t hDevice, ///< [in] handle of the Device object + ur_kernel_sub_group_info_t + propName, ///< [in] name of the SubGroup property to query + size_t propSize, ///< [in] size of the Kernel SubGroup property value + void *pPropValue, ///< [in,out][range(0, propSize)][optional] value of the + ///< Kernel SubGroup property. + size_t *pPropSizeRet ///< [out][optional] pointer to the actual size in + ///< bytes of data being queried by propName. +) { + std::ignore = hDevice; + + UrReturnHelper returnValue(propSize, pPropValue, pPropSizeRet); + + auto props = hKernel->getProperties(hDevice); + + std::shared_lock Guard(hKernel->Mutex); + if (propName == UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE) { + returnValue(uint32_t{props.maxSubgroupSize}); + } else if (propName == UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS) { + returnValue(uint32_t{props.maxNumSubgroups}); + } else if (propName == UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS) { + returnValue(uint32_t{props.requiredNumSubGroups}); + } else if (propName == UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL) { + returnValue(uint32_t{props.requiredSubgroupSize}); + } else { + die("urKernelGetSubGroupInfo: parameter not implemented"); + return {}; + } + return UR_RESULT_SUCCESS; +} } // namespace ur::level_zero diff --git a/source/adapters/level_zero/v2/kernel.hpp b/source/adapters/level_zero/v2/kernel.hpp index d4765ae9fc..2d3a891826 100644 --- a/source/adapters/level_zero/v2/kernel.hpp +++ b/source/adapters/level_zero/v2/kernel.hpp @@ -15,11 +15,11 @@ #include "common.hpp" struct ur_single_device_kernel_t { - ur_single_device_kernel_t(ze_device_handle_t hDevice, + ur_single_device_kernel_t(ur_device_handle_t hDevice, ze_kernel_handle_t hKernel, bool ownZeHandle); ur_result_t release(); - ze_device_handle_t hDevice; + ur_device_handle_t hDevice; v2::raii::ze_kernel_handle_t hKernel; mutable ZeCache> zeKernelProperties; }; @@ -40,6 +40,9 @@ struct ur_kernel_handle_t_ : _ur_object { // Get program handle of the kernel. ur_program_handle_t getProgramHandle() const; + // Get devices the kernel is built for. + std::vector getDevices() const; + // Get name of the kernel. const std::string &getName() const; @@ -57,6 +60,10 @@ struct ur_kernel_handle_t_ : _ur_object { const ur_kernel_arg_pointer_properties_t *pProperties, const void *pArgValue); + // Implementation of urKernelSetExecInfo. + ur_result_t setExecInfo(ur_kernel_exec_info_t propName, + const void *pPropValue); + // Perform cleanup. ur_result_t release(); diff --git a/source/adapters/level_zero/v2/memory.cpp b/source/adapters/level_zero/v2/memory.cpp new file mode 100644 index 0000000000..fc9a7522a4 --- /dev/null +++ b/source/adapters/level_zero/v2/memory.cpp @@ -0,0 +1,180 @@ +//===--------- memory.cpp - Level Zero Adapter ---------------------------===// +// +// Copyright (C) 2024 Intel Corporation +// +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM +// Exceptions. See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "memory.hpp" +#include "context.hpp" + +#include "../helpers/memory_helpers.hpp" + +ur_mem_handle_t_::ur_mem_handle_t_(ur_context_handle_t hContext, size_t size) + : hContext(hContext), size(size) {} + +ur_host_mem_handle_t::ur_host_mem_handle_t(ur_context_handle_t hContext, + void *hostPtr, size_t size, + host_ptr_action_t hostPtrAction) + : ur_mem_handle_t_(hContext, size) { + bool hostPtrImported = false; + if (hostPtrAction == host_ptr_action_t::import) { + hostPtrImported = + maybeImportUSM(hContext->getPlatform()->ZeDriverHandleExpTranslated, + hContext->getZeHandle(), hostPtr, size); + } + + if (!hostPtrImported) { + // TODO: use UMF + ZeStruct hostDesc; + ZE2UR_CALL_THROWS(zeMemAllocHost, (hContext->getZeHandle(), &hostDesc, size, + 0, &this->ptr)); + + if (hostPtr) { + std::memcpy(this->ptr, hostPtr, size); + } + } +} + +ur_host_mem_handle_t::~ur_host_mem_handle_t() { + // TODO: use UMF API here + if (ptr) { + ZE_CALL_NOCHECK(zeMemFree, (hContext->getZeHandle(), ptr)); + } +} + +void *ur_host_mem_handle_t::getPtr(ur_device_handle_t hDevice) { + std::ignore = hDevice; + return ptr; +} + +ur_device_mem_handle_t::ur_device_mem_handle_t(ur_context_handle_t hContext, + void *hostPtr, size_t size) + : ur_mem_handle_t_(hContext, size), + deviceAllocations(hContext->getPlatform()->getNumDevices()) { + // Legacy adapter allocated the memory directly on a device (first on the + // contxt) and if the buffer is used on another device, memory is migrated + // (depending on an env var setting). + // + // TODO: port this behavior or figure out if it makes sense to keep the memory + // in a host buffer (e.g. for smaller sizes). + if (hostPtr) { + buffer.assign(reinterpret_cast(hostPtr), + reinterpret_cast(hostPtr) + size); + } +} + +ur_device_mem_handle_t::~ur_device_mem_handle_t() { + // TODO: use UMF API here + for (auto &ptr : deviceAllocations) { + if (ptr) { + ZE_CALL_NOCHECK(zeMemFree, (hContext->getZeHandle(), ptr)); + } + } +} + +void *ur_device_mem_handle_t::getPtr(ur_device_handle_t hDevice) { + std::lock_guard lock(this->Mutex); + + auto &ptr = deviceAllocations[hDevice->Id.value()]; + if (!ptr) { + ZeStruct deviceDesc; + ZE2UR_CALL_THROWS(zeMemAllocDevice, (hContext->getZeHandle(), &deviceDesc, + size, 0, hDevice->ZeDevice, &ptr)); + + if (!buffer.empty()) { + auto commandList = hContext->commandListCache.getImmediateCommandList( + hDevice->ZeDevice, true, + hDevice + ->QueueGroup + [ur_device_handle_t_::queue_group_info_t::type::Compute] + .ZeOrdinal, + ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS, ZE_COMMAND_QUEUE_PRIORITY_NORMAL, + std::nullopt); + ZE2UR_CALL_THROWS( + zeCommandListAppendMemoryCopy, + (commandList.get(), ptr, buffer.data(), size, nullptr, 0, nullptr)); + } + } + return ptr; +} + +namespace ur::level_zero { +ur_result_t urMemBufferCreate(ur_context_handle_t hContext, + ur_mem_flags_t flags, size_t size, + const ur_buffer_properties_t *pProperties, + ur_mem_handle_t *phBuffer) { + if (flags & UR_MEM_FLAG_ALLOC_HOST_POINTER) { + // TODO: + // Having PI_MEM_FLAGS_HOST_PTR_ALLOC for buffer requires allocation of + // pinned host memory, see: + // sycl/doc/extensions/supported/sycl_ext_oneapi_use_pinned_host_memory_property.asciidoc + // We are however missing such functionality in Level Zero, so we just + // ignore the flag for now. + } + + void *hostPtr = pProperties ? pProperties->pHost : nullptr; + + // We treat integrated devices (physical memory shared with the CPU) + // differently from discrete devices (those with distinct memories). + // For integrated devices, allocating the buffer in the host memory + // enables automatic access from the device, and makes copying + // unnecessary in the map/unmap operations. This improves performance. + bool useHostBuffer = hContext->getDevices().size() == 1 && + hContext->getDevices()[0]->ZeDeviceProperties->flags & + ZE_DEVICE_PROPERTY_FLAG_INTEGRATED; + + if (useHostBuffer) { + // TODO: assert that if hostPtr is set, either UR_MEM_FLAG_USE_HOST_POINTER + // or UR_MEM_FLAG_ALLOC_COPY_HOST_POINTER is set? + auto hostPtrAction = flags & UR_MEM_FLAG_USE_HOST_POINTER + ? ur_host_mem_handle_t::host_ptr_action_t::import + : ur_host_mem_handle_t::host_ptr_action_t::copy; + *phBuffer = + new ur_host_mem_handle_t(hContext, hostPtr, size, hostPtrAction); + } else { + *phBuffer = new ur_device_mem_handle_t(hContext, hostPtr, size); + } + + return UR_RESULT_SUCCESS; +} + +ur_result_t urMemBufferPartition(ur_mem_handle_t hBuffer, ur_mem_flags_t flags, + ur_buffer_create_type_t bufferCreateType, + const ur_buffer_region_t *pRegion, + ur_mem_handle_t *phMem) { + std::ignore = hBuffer; + std::ignore = flags; + std::ignore = bufferCreateType; + std::ignore = pRegion; + std::ignore = phMem; + logger::error("{} function not implemented!", __FUNCTION__); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urMemBufferCreateWithNativeHandle( + ur_native_handle_t hNativeMem, ur_context_handle_t hContext, + const ur_mem_native_properties_t *pProperties, ur_mem_handle_t *phMem) { + std::ignore = hNativeMem; + std::ignore = hContext; + std::ignore = pProperties; + std::ignore = phMem; + logger::error("{} function not implemented!", __FUNCTION__); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urMemRetain(ur_mem_handle_t hMem) { + hMem->RefCount.increment(); + return UR_RESULT_SUCCESS; +} + +ur_result_t urMemRelease(ur_mem_handle_t hMem) { + if (hMem->RefCount.decrementAndTest()) { + delete hMem; + } + return UR_RESULT_SUCCESS; +} +} // namespace ur::level_zero diff --git a/source/adapters/level_zero/v2/memory.hpp b/source/adapters/level_zero/v2/memory.hpp new file mode 100644 index 0000000000..be4456d728 --- /dev/null +++ b/source/adapters/level_zero/v2/memory.hpp @@ -0,0 +1,55 @@ +//===--------- memory.hpp - Level Zero Adapter ---------------------------===// +// +// Copyright (C) 2024 Intel Corporation +// +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM +// Exceptions. See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include "common.hpp" + +struct ur_mem_handle_t_ : _ur_object { + ur_mem_handle_t_(ur_context_handle_t hContext, size_t size); + virtual ~ur_mem_handle_t_() = default; + + virtual void *getPtr(ur_device_handle_t) = 0; + + inline size_t getSize() { return size; } + +protected: + const ur_context_handle_t hContext; + const size_t size; +}; + +struct ur_host_mem_handle_t : public ur_mem_handle_t_ { + enum class host_ptr_action_t { import, copy }; + + ur_host_mem_handle_t(ur_context_handle_t hContext, void *hostPtr, size_t size, + host_ptr_action_t useHostPtr); + ~ur_host_mem_handle_t(); + + void *getPtr(ur_device_handle_t) override; + +private: + void *ptr; +}; + +struct ur_device_mem_handle_t : public ur_mem_handle_t_ { + ur_device_mem_handle_t(ur_context_handle_t hContext, void *hostPtr, + size_t size); + ~ur_device_mem_handle_t(); + + void *getPtr(ur_device_handle_t) override; + +private: + std::vector buffer; + + // Vector of per-device allocations indexed by device->Id + std::vector deviceAllocations; +}; diff --git a/source/adapters/level_zero/v2/queue_immediate_in_order.cpp b/source/adapters/level_zero/v2/queue_immediate_in_order.cpp index d6afd2eef3..3882a69135 100644 --- a/source/adapters/level_zero/v2/queue_immediate_in_order.cpp +++ b/source/adapters/level_zero/v2/queue_immediate_in_order.cpp @@ -10,6 +10,7 @@ #include "queue_immediate_in_order.hpp" #include "kernel.hpp" +#include "memory.hpp" #include "ur.hpp" #include "../helpers/kernel_helpers.hpp" @@ -146,6 +147,7 @@ ur_queue_immediate_in_order_t::queueGetInfo(ur_queue_info_t propName, // We can exit early if we have in-order queue. if (!lastHandler) return ReturnValue(true); + [[fallthrough]]; } default: logger::error( @@ -251,49 +253,59 @@ ur_result_t ur_queue_immediate_in_order_t::enqueueKernelLaunch( ur_result_t ur_queue_immediate_in_order_t::enqueueEventsWait( uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + TRACK_SCOPE_LATENCY("ur_queue_immediate_in_order_t::enqueueEventsWait"); + + std::unique_lock lock(this->Mutex); + + auto handler = getCommandListHandlerForCompute(); + auto signalEvent = getSignalEvent(handler, phEvent); + auto [pWaitEvents, numWaitEvents] = + getWaitListView(phEventWaitList, numEventsInWaitList, handler); + + ZE2UR_CALL(zeCommandListAppendWaitOnEvents, + (handler->commandList.get(), numWaitEvents, pWaitEvents)); + ZE2UR_CALL(zeCommandListAppendSignalEvent, + (handler->commandList.get(), signalEvent)); + + lastHandler = handler; + + return UR_RESULT_SUCCESS; } ur_result_t ur_queue_immediate_in_order_t::enqueueEventsWaitWithBarrier( uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + // For in-order queue we don't need a real barrier, just wait for + // requested events in potentially different queues and add a "barrier" + // event signal because it is already guaranteed that previous commands + // in this queue are completed when the signal is started. + return enqueueEventsWait(numEventsInWaitList, phEventWaitList, phEvent); } ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferRead( ur_mem_handle_t hBuffer, bool blockingRead, size_t offset, size_t size, void *pDst, uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = hBuffer; - std::ignore = blockingRead; - std::ignore = offset; - std::ignore = size; - std::ignore = pDst; - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + TRACK_SCOPE_LATENCY("ur_queue_immediate_in_order_t::enqueueMemBufferRead"); + + UR_ASSERT(offset + size <= hBuffer->getSize(), UR_RESULT_ERROR_INVALID_SIZE); + + auto ptr = ur_cast(hBuffer->getPtr(hDevice)); + return enqueueUSMMemcpy(blockingRead, pDst, ptr + offset, size, + numEventsInWaitList, phEventWaitList, phEvent); } ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferWrite( ur_mem_handle_t hBuffer, bool blockingWrite, size_t offset, size_t size, const void *pSrc, uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = hBuffer; - std::ignore = blockingWrite; - std::ignore = offset; - std::ignore = size; - std::ignore = pSrc; - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + TRACK_SCOPE_LATENCY("ur_queue_immediate_in_order_t::enqueueMemBufferWrite"); + + UR_ASSERT(offset + size <= hBuffer->getSize(), UR_RESULT_ERROR_INVALID_SIZE); + + auto ptr = ur_cast(hBuffer->getPtr(hDevice)); + return enqueueUSMMemcpy(blockingWrite, ptr + offset, pSrc, size, + numEventsInWaitList, phEventWaitList, phEvent); } ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferReadRect( @@ -344,15 +356,18 @@ ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferCopy( ur_mem_handle_t hBufferSrc, ur_mem_handle_t hBufferDst, size_t srcOffset, size_t dstOffset, size_t size, uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = hBufferSrc; - std::ignore = hBufferDst; - std::ignore = srcOffset; - std::ignore = dstOffset; - std::ignore = size; - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + TRACK_SCOPE_LATENCY("ur_queue_immediate_in_order_t::enqueueMemBufferCopy"); + + UR_ASSERT(srcOffset + size <= hBufferSrc->getSize(), + UR_RESULT_ERROR_INVALID_SIZE); + UR_ASSERT(dstOffset + size <= hBufferDst->getSize(), + UR_RESULT_ERROR_INVALID_SIZE); + + auto srcPtr = ur_cast(hBufferSrc->getPtr(hDevice)); + auto dstPtr = ur_cast(hBufferDst->getPtr(hDevice)); + + return enqueueUSMMemcpy(false, dstPtr + dstOffset, srcPtr + srcOffset, size, + numEventsInWaitList, phEventWaitList, phEvent); } ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferCopyRect( @@ -380,15 +395,13 @@ ur_result_t ur_queue_immediate_in_order_t::enqueueMemBufferFill( ur_mem_handle_t hBuffer, const void *pPattern, size_t patternSize, size_t offset, size_t size, uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { - std::ignore = hBuffer; - std::ignore = pPattern; - std::ignore = patternSize; - std::ignore = offset; - std::ignore = size; - std::ignore = numEventsInWaitList; - std::ignore = phEventWaitList; - std::ignore = phEvent; - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + TRACK_SCOPE_LATENCY("ur_queue_immediate_in_order_t::enqueueMemBufferFill"); + + UR_ASSERT(offset + size <= hBuffer->getSize(), UR_RESULT_ERROR_INVALID_SIZE); + + auto ptr = ur_cast(hBuffer->getPtr(hDevice)); + return enqueueUSMFill(ptr + offset, patternSize, pPattern, size, + numEventsInWaitList, phEventWaitList, phEvent); } ur_result_t ur_queue_immediate_in_order_t::enqueueMemImageRead( diff --git a/source/adapters/opencl/device.cpp b/source/adapters/opencl/device.cpp index 071a3a7c5a..36d08548c3 100644 --- a/source/adapters/opencl/device.cpp +++ b/source/adapters/opencl/device.cpp @@ -32,6 +32,17 @@ ur_result_t cl_adapter::getDeviceVersion(cl_device_id Dev, return UR_RESULT_SUCCESS; } +static bool isIntelFPGAEmuDevice(cl_device_id Dev) { + size_t NameSize = 0; + CL_RETURN_ON_FAILURE( + clGetDeviceInfo(Dev, CL_DEVICE_NAME, 0, nullptr, &NameSize)); + std::string NameStr(NameSize, '\0'); + CL_RETURN_ON_FAILURE( + clGetDeviceInfo(Dev, CL_DEVICE_NAME, NameSize, NameStr.data(), nullptr)); + + return NameStr.find("Intel(R) FPGA Emulation Device") != std::string::npos; +} + ur_result_t cl_adapter::checkDeviceExtensions( cl_device_id Dev, const std::vector &Exts, bool &Supported) { size_t ExtSize = 0; @@ -46,6 +57,14 @@ ur_result_t cl_adapter::checkDeviceExtensions( Supported = true; for (const std::string &Ext : Exts) { if (!(Supported = (ExtStr.find(Ext) != std::string::npos))) { + // The Intel FPGA emulation device does actually support these, even if it + // doesn't report them. + if (isIntelFPGAEmuDevice(Dev) && + (Ext == "cl_intel_device_attribute_query" || + Ext == "cl_intel_required_subgroup_size")) { + Supported = true; + continue; + } break; } } diff --git a/source/common/latency_tracker.hpp b/source/common/latency_tracker.hpp index 03ed6b6443..bf20e3819f 100644 --- a/source/common/latency_tracker.hpp +++ b/source/common/latency_tracker.hpp @@ -42,6 +42,7 @@ static constexpr double percentiles[numPercentiles] = { 50.0, 90.0, 99.0, 99.9, 99.99, 99.999, 99.9999}; struct latencyValues { + int64_t count; int64_t min; int64_t max; int64_t mean; @@ -54,6 +55,7 @@ using histogram_ptr = static inline latencyValues getValues(const struct hdr_histogram *histogram) { latencyValues values; + values.count = histogram->total_count; values.max = hdr_max(histogram); values.min = hdr_min(histogram); values.mean = static_cast(hdr_mean(histogram)); @@ -92,13 +94,16 @@ class latency_printer { for (auto &[name, histogram] : values) { auto value = getValues(histogram.get()); - logger.log(logger::Level::INFO, - "{},{},{},{},{},{},{},{},{},{},{},{},ns", name, - value.min, value.max, value.mean, value.stddev, - value.percentileValues[0], value.percentileValues[1], - value.percentileValues[2], value.percentileValues[3], - value.percentileValues[4], value.percentileValues[5], - value.percentileValues[6]); + auto f = groupDigits; + logger.log( + logger::Level::INFO, + "{},{},{},{},{},{},{},{},{},{},{},{},{},{},ns", name, + f(value.mean), f(value.percentileValues[0]), + f(value.percentileValues[1]), f(value.percentileValues[2]), + f(value.percentileValues[3]), f(value.percentileValues[4]), + f(value.percentileValues[5]), f(value.percentileValues[6]), + f(value.count), f(value.count * value.mean), f(value.min), + f(value.max), value.stddev); } } @@ -106,7 +111,8 @@ class latency_printer { inline void printHeader() { logger.log(logger::Level::INFO, "Latency histogram:"); logger.log(logger::Level::INFO, - "name,min,max,mean,stdev,p{},p{},p{},p{},p{},p{},p{},unit", + "name,mean,p{},p{},p{},p{},p{},p{}" + ",p{},count,sum,min,max,stdev,unit", percentiles[0], percentiles[1], percentiles[2], percentiles[3], percentiles[4], percentiles[5], percentiles[6]); diff --git a/source/common/ur_util.hpp b/source/common/ur_util.hpp index d66612b917..0ede3c93dc 100644 --- a/source/common/ur_util.hpp +++ b/source/common/ur_util.hpp @@ -480,6 +480,25 @@ template class AtomicSingleton { } }; +template +static inline std::string groupDigits(Numeric numeric) { + auto number = std::to_string(numeric); + std::string sign = numeric >= 0 ? "" : "-"; + auto digits = number.substr(sign.size(), number.size() - sign.size()); + + std::string separated; + + for (size_t i = 0; i < digits.size(); i++) { + separated.push_back(digits[i]); + + if (i != digits.size() - 1 && (digits.size() - i - 1) % 3 == 0) { + separated.push_back('\''); + } + } + + return sign + separated; +} + template Spinlock> AtomicSingleton::instance; #endif /* UR_UTIL_H */ diff --git a/source/loader/layers/sanitizer/asan_options.hpp b/source/loader/layers/sanitizer/asan_options.hpp index 298639b73c..eb3f6bb03d 100644 --- a/source/loader/layers/sanitizer/asan_options.hpp +++ b/source/loader/layers/sanitizer/asan_options.hpp @@ -42,7 +42,17 @@ struct AsanOptions { private: AsanOptions(logger::Logger &logger) { - auto OptionsEnvMap = getenv_to_map("UR_LAYER_ASAN_OPTIONS"); + std::optional OptionsEnvMap; + try { + OptionsEnvMap = getenv_to_map("UR_LAYER_ASAN_OPTIONS"); + } catch (const std::invalid_argument &e) { + std::stringstream SS; + SS << "[ERROR]: "; + SS << e.what(); + logger.always(SS.str().c_str()); + die("Sanitizer failed to parse options.\n"); + } + if (!OptionsEnvMap.has_value()) { return; } @@ -74,8 +84,8 @@ struct AsanOptions { Opt = false; } else { std::stringstream SS; - SS << "[ERROR]: \"" << Name << "\" is set to \"" - << Value << "\", which is not an valid setting. "; + SS << "\"" << Name << "\" is set to \"" << Value + << "\", which is not an valid setting. "; SS << "Acceptable input are: for enable, use:"; for (auto &S : TrueStrings) { SS << " \"" << S << "\""; @@ -86,7 +96,8 @@ struct AsanOptions { SS << " \"" << S << "\""; } SS << "."; - die(SS.str().c_str()); + logger.error(SS.str().c_str()); + die("Sanitizer failed to parse options.\n"); } } }; @@ -106,9 +117,10 @@ struct AsanOptions { } MaxQuarantineSizeMB = temp_long; } catch (...) { - die("[ERROR]: \"quarantine_size_mb\" should be " - "an positive integer that smaller than or equal to " - "4294967295."); + logger.error("\"quarantine_size_mb\" should be " + "an integer in range[0, {}].", + UINT32_MAX); + die("Sanitizer failed to parse options.\n"); } } @@ -120,10 +132,12 @@ struct AsanOptions { if (MinRZSize < 16) { MinRZSize = 16; logger.warning("Trying to set redzone size to a " - "value less than 16 is ignored"); + "value less than 16 is ignored."); } } catch (...) { - die("[ERROR]: \"redzone\" should be an integer"); + logger.error( + "\"redzone\" should be an integer in range[0, 16]."); + die("Sanitizer failed to parse options.\n"); } } @@ -135,10 +149,12 @@ struct AsanOptions { if (MaxRZSize > 2048) { MaxRZSize = 2048; logger.warning("Trying to set max redzone size to a " - "value greater than 2048 is ignored"); + "value greater than 2048 is ignored."); } } catch (...) { - die("[ERROR]: \"max_redzone\" should be an integer"); + logger.error( + "\"max_redzone\" should be an integer in range[0, 2048]."); + die("Sanitizer failed to parse options.\n"); } } } diff --git a/source/loader/layers/tracing/ur_tracing_layer.cpp b/source/loader/layers/tracing/ur_tracing_layer.cpp index 722ee77faa..c6fd4ca40d 100644 --- a/source/loader/layers/tracing/ur_tracing_layer.cpp +++ b/source/loader/layers/tracing/ur_tracing_layer.cpp @@ -34,14 +34,15 @@ struct XptiContextManager { ~XptiContextManager() { xptiFrameworkFinalize(); } }; -static std::shared_ptr xptiContextManagerGlobal = [] { - return std::make_shared(); -}(); +static std::shared_ptr xptiContextManagerGet() { + static auto contextManager = std::make_shared(); + return contextManager; +}; static thread_local xpti_td *activeEvent; /////////////////////////////////////////////////////////////////////////////// context_t::context_t() : logger(logger::create_logger("tracing", true, true)) { - this->xptiContextManager = xptiContextManagerGlobal; + this->xptiContextManager = xptiContextManagerGet(); call_stream_id = xptiRegisterStream(CALL_STREAM_NAME); std::ostringstream streamv; diff --git a/source/loader/ur_adapter_registry.hpp b/source/loader/ur_adapter_registry.hpp index f718f35490..7df799ab1e 100644 --- a/source/loader/ur_adapter_registry.hpp +++ b/source/loader/ur_adapter_registry.hpp @@ -163,10 +163,123 @@ class AdapterRegistry { return paths.empty() ? std::nullopt : std::optional(paths); } + ur_result_t readPreFilterODS(std::string platformBackendName) { + // TODO: Refactor this to the common code such that both the prefilter and urDeviceGetSelected use the same functionality. + bool acceptLibrary = true; + std::optional odsEnvMap; + try { + odsEnvMap = getenv_to_map("ONEAPI_DEVICE_SELECTOR", false); + + } catch (...) { + // If the selector is malformed, then we ignore selector and return success. + logger::error("ERROR: missing backend, format of filter = " + "'[!]backend:filterStrings'"); + return UR_RESULT_SUCCESS; + } + logger::debug( + "getenv_to_map parsed env var and {} a map", + (odsEnvMap.has_value() ? "produced" : "failed to produce")); + + // if the ODS env var is not set at all, then pretend it was set to the default + using EnvVarMap = std::map>; + EnvVarMap mapODS = + odsEnvMap.has_value() ? odsEnvMap.value() : EnvVarMap{{"*", {"*"}}}; + for (auto &termPair : mapODS) { + std::string backend = termPair.first; + // TODO: Figure out how to process all ODS errors rather than returning + // on the first error. + if (backend.empty()) { + // FIXME: never true because getenv_to_map rejects this case + // malformed term: missing backend -- output ERROR, then continue + logger::error("ERROR: missing backend, format of filter = " + "'[!]backend:filterStrings'"); + continue; + } + logger::debug("ONEAPI_DEVICE_SELECTOR Pre-Filter with backend '{}' " + "and platform library name '{}'", + backend, platformBackendName); + enum FilterType { + AcceptFilter, + DiscardFilter, + } termType = + (backend.front() != '!') ? AcceptFilter : DiscardFilter; + logger::debug( + "termType is {}", + (termType != AcceptFilter ? "DiscardFilter" : "AcceptFilter")); + if (termType != AcceptFilter) { + logger::debug("DEBUG: backend was '{}'", backend); + backend.erase(backend.cbegin()); + logger::debug("DEBUG: backend now '{}'", backend); + } + + // Verify that the backend string is valid, otherwise ignore the backend. + if ((strcmp(backend.c_str(), "*") != 0) && + (strcmp(backend.c_str(), "level_zero") != 0) && + (strcmp(backend.c_str(), "opencl") != 0) && + (strcmp(backend.c_str(), "cuda") != 0) && + (strcmp(backend.c_str(), "hip") != 0)) { + logger::debug("ONEAPI_DEVICE_SELECTOR Pre-Filter with illegal " + "backend '{}' ", + backend); + continue; + } + + // case-insensitive comparison by converting both tolower + std::transform(platformBackendName.begin(), + platformBackendName.end(), + platformBackendName.begin(), + [](unsigned char c) { return std::tolower(c); }); + std::transform(backend.begin(), backend.end(), backend.begin(), + [](unsigned char c) { return std::tolower(c); }); + std::size_t nameFound = platformBackendName.find(backend); + + bool backendFound = nameFound != std::string::npos; + if (termType == AcceptFilter) { + if (backend.front() != '*' && !backendFound) { + logger::debug( + "The ONEAPI_DEVICE_SELECTOR backend name '{}' was not " + "found in the platform library name '{}'", + backend, platformBackendName); + acceptLibrary = false; + continue; + } else if (backend.front() == '*' || backendFound) { + return UR_RESULT_SUCCESS; + } + } else { + if (backendFound || backend.front() == '*') { + acceptLibrary = false; + logger::debug( + "The ONEAPI_DEVICE_SELECTOR backend name for discard " + "'{}' was found in the platform library name '{}'", + backend, platformBackendName); + continue; + } + } + } + if (acceptLibrary) { + return UR_RESULT_SUCCESS; + } + return UR_RESULT_ERROR_INVALID_VALUE; + } + void discoverKnownAdapters() { auto searchPathsEnvOpt = getEnvAdapterSearchPaths(); auto loaderLibPathOpt = getLoaderLibPath(); +#if defined(_WIN32) + bool loaderPreFilter = getenv_tobool("UR_LOADER_PRELOAD_FILTER", false); +#else + bool loaderPreFilter = getenv_tobool("UR_LOADER_PRELOAD_FILTER", true); +#endif for (const auto &adapterName : knownAdapterNames) { + + if (loaderPreFilter) { + if (readPreFilterODS(adapterName) != UR_RESULT_SUCCESS) { + logger::debug("The adapter '{}' was removed based on the " + "pre-filter from ONEAPI_DEVICE_SELECTOR.", + adapterName); + continue; + } + } std::vector loadPaths; // Adapter search order: diff --git a/source/loader/ur_lib.cpp b/source/loader/ur_lib.cpp index 9aad7159c3..7f2d1baa13 100644 --- a/source/loader/ur_lib.cpp +++ b/source/loader/ur_lib.cpp @@ -560,19 +560,20 @@ ur_result_t urDeviceGetSelected(ur_platform_handle_t hPlatform, const auto thirdDeviceId = getDeviceId(thirdPart); deviceList.push_back(DeviceSpec{ DevicePartLevel::SUBSUB, hardwareType, firstDeviceId, - secondDeviceId, thirdDeviceId}); + secondDeviceId, thirdDeviceId, nullptr}); } else { // second dot not found, this is a subdevice - deviceList.push_back(DeviceSpec{DevicePartLevel::SUB, - hardwareType, firstDeviceId, - secondDeviceId}); + deviceList.push_back( + DeviceSpec{DevicePartLevel::SUB, hardwareType, + firstDeviceId, secondDeviceId, 0, nullptr}); } } else { // first dot not found, this is a root device const auto hardwareType = getRootHardwareType(filterString); const auto firstDeviceId = getDeviceId(filterString); deviceList.push_back(DeviceSpec{DevicePartLevel::ROOT, - hardwareType, firstDeviceId}); + hardwareType, firstDeviceId, 0, + 0, nullptr}); } } } @@ -587,8 +588,9 @@ ur_result_t urDeviceGetSelected(ur_platform_handle_t hPlatform, // for example, we pretend that "garbage:0;!cuda:*" was just "!cuda:*" // so we add an implicit accept-all term (equivalent to prepending "*:*;") // as we would have done if the user had given us the corrected string - acceptDeviceList.push_back(DeviceSpec{ - DevicePartLevel::ROOT, ::UR_DEVICE_TYPE_ALL, DeviceIdTypeALL}); + acceptDeviceList.push_back(DeviceSpec{DevicePartLevel::ROOT, + ::UR_DEVICE_TYPE_ALL, + DeviceIdTypeALL, 0, 0, nullptr}); } logger::debug("DEBUG: size of acceptDeviceList = {}", diff --git a/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp b/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp index 6ee49dbbfb..b3918c7818 100644 --- a/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp +++ b/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp @@ -24,7 +24,8 @@ TEST_P(urLevelZeroKernelNativeHandleTest, OwnedHandleRelease) { auto kernel_name = uur::KernelsEnvironment::instance->GetEntryPointNames("foo")[0]; - ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC}; + ze_module_desc_t moduleDesc{}; + moduleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC; moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; moduleDesc.inputSize = il_binary->size(); moduleDesc.pInputModule = @@ -36,7 +37,8 @@ TEST_P(urLevelZeroKernelNativeHandleTest, OwnedHandleRelease) { &module, NULL), ZE_RESULT_SUCCESS); - ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC}; + ze_kernel_desc_t kernelDesc{}; + kernelDesc.stype = ZE_STRUCTURE_TYPE_KERNEL_DESC; kernelDesc.pKernelName = kernel_name.c_str(); ze_kernel_handle_t native_kernel; @@ -75,7 +77,8 @@ TEST_P(urLevelZeroKernelNativeHandleTest, NullProgram) { auto kernel_name = uur::KernelsEnvironment::instance->GetEntryPointNames("foo")[0]; - ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC}; + ze_module_desc_t moduleDesc{}; + moduleDesc.stype = ZE_STRUCTURE_TYPE_MODULE_DESC; moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; moduleDesc.inputSize = il_binary->size(); moduleDesc.pInputModule = @@ -87,7 +90,8 @@ TEST_P(urLevelZeroKernelNativeHandleTest, NullProgram) { &module, NULL), ZE_RESULT_SUCCESS); - ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC}; + ze_kernel_desc_t kernelDesc{}; + kernelDesc.stype = ZE_STRUCTURE_TYPE_KERNEL_DESC; kernelDesc.pKernelName = kernel_name.c_str(); ze_kernel_handle_t native_kernel; diff --git a/test/adapters/level_zero/v2/command_list_cache_test.cpp b/test/adapters/level_zero/v2/command_list_cache_test.cpp index 74bcbf4634..44755b699e 100644 --- a/test/adapters/level_zero/v2/command_list_cache_test.cpp +++ b/test/adapters/level_zero/v2/command_list_cache_test.cpp @@ -23,7 +23,7 @@ struct CommandListCacheTest : public uur::urContextTest {}; UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(CommandListCacheTest); TEST_P(CommandListCacheTest, CanStoreAndRetriveImmediateAndRegularCmdLists) { - v2::command_list_cache_t cache(context->hContext); + v2::command_list_cache_t cache(context->getZeHandle()); bool IsInOrder = false; uint32_t Ordinal = 0; @@ -75,7 +75,7 @@ TEST_P(CommandListCacheTest, CanStoreAndRetriveImmediateAndRegularCmdLists) { } TEST_P(CommandListCacheTest, ImmediateCommandListsHaveProperAttributes) { - v2::command_list_cache_t cache(context->hContext); + v2::command_list_cache_t cache(context->getZeHandle()); uint32_t numQueueGroups = 0; ASSERT_EQ(zeDeviceGetCommandQueueGroupProperties(device->ZeDevice, diff --git a/test/conformance/context/context_adapter_level_zero_v2.match b/test/conformance/context/context_adapter_level_zero_v2.match index fc2d1b8324..3dea8da6e5 100644 --- a/test/conformance/context/context_adapter_level_zero_v2.match +++ b/test/conformance/context/context_adapter_level_zero_v2.match @@ -1,5 +1,3 @@ urContextCreateWithNativeHandleTest.InvalidNullHandleAdapter/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ urContextCreateWithNativeHandleTest.InvalidNullPointerContext/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urContextGetInfoTestWithInfoParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_CONTEXT_INFO_USM_MEMCPY2D_SUPPORT -urContextGetInfoTestWithInfoParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_CONTEXT_INFO_USM_FILL2D_SUPPORT urContextSetExtendedDeleterTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ diff --git a/test/conformance/context/urContextCreate.cpp b/test/conformance/context/urContextCreate.cpp index d5fb59389f..0f268a3992 100644 --- a/test/conformance/context/urContextCreate.cpp +++ b/test/conformance/context/urContextCreate.cpp @@ -17,7 +17,8 @@ TEST_P(urContextCreateTest, Success) { } TEST_P(urContextCreateTest, SuccessWithProperties) { - ur_context_properties_t properties{UR_STRUCTURE_TYPE_CONTEXT_PROPERTIES}; + ur_context_properties_t properties{UR_STRUCTURE_TYPE_CONTEXT_PROPERTIES, + nullptr, 0}; uur::raii::Context context = nullptr; ASSERT_SUCCESS(urContextCreate(1, &device, &properties, context.ptr())); ASSERT_NE(nullptr, context); diff --git a/test/conformance/device_code/CMakeLists.txt b/test/conformance/device_code/CMakeLists.txt index af0bc83d8a..5445531961 100644 --- a/test/conformance/device_code/CMakeLists.txt +++ b/test/conformance/device_code/CMakeLists.txt @@ -144,6 +144,7 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fill_usm.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/foo.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/image_copy.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/inc.cpp) +add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/increment.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/mean.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/cpy_and_mult.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/cpy_and_mult_usm.cpp) diff --git a/test/conformance/device_code/increment.cpp b/test/conformance/device_code/increment.cpp new file mode 100644 index 0000000000..14094c4963 --- /dev/null +++ b/test/conformance/device_code/increment.cpp @@ -0,0 +1,20 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include + +int main() { + + const size_t inputSize = 1; + sycl::queue sycl_queue; + uint32_t *inputArray = sycl::malloc_shared(inputSize, sycl_queue); + + sycl_queue.submit([&](sycl::handler &cgh) { + cgh.parallel_for( + sycl::range<1>(inputSize), + [=](sycl::id<1> itemID) { inputArray[itemID] += 1; }); + }); + return 0; +} diff --git a/test/conformance/device_code/indexers_usm.cpp b/test/conformance/device_code/indexers_usm.cpp index e055fa47cc..cd3b56bf0c 100644 --- a/test/conformance/device_code/indexers_usm.cpp +++ b/test/conformance/device_code/indexers_usm.cpp @@ -3,6 +3,9 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// Offsets are deprecated, but we should still test that they work +#pragma clang diagnostic ignored "-Wdeprecated-declarations" + #include int main() { diff --git a/test/conformance/device_code/linker_error.cpp b/test/conformance/device_code/linker_error.cpp index 5fc7eebf6f..8afa369bb2 100644 --- a/test/conformance/device_code/linker_error.cpp +++ b/test/conformance/device_code/linker_error.cpp @@ -3,17 +3,17 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include +#include SYCL_EXTERNAL void this_function_does_not_exist(); int main() { - cl::sycl::queue deviceQueue; - cl::sycl::range<1> numOfItems{1}; + sycl::queue deviceQueue; + sycl::range<1> numOfItems{1}; try { - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto kern = [=](cl::sycl::id<1>) { + deviceQueue.submit([&](sycl::handler &cgh) { + auto kern = [=](sycl::id<1>) { #ifdef __SYCL_DEVICE_ONLY__ this_function_does_not_exist(); #endif diff --git a/test/conformance/enqueue/CMakeLists.txt b/test/conformance/enqueue/CMakeLists.txt index 7cc68203a0..1e19658dac 100644 --- a/test/conformance/enqueue/CMakeLists.txt +++ b/test/conformance/enqueue/CMakeLists.txt @@ -9,6 +9,7 @@ add_conformance_test_with_kernels_environment(enqueue urEnqueueEventsWait.cpp urEnqueueEventsWaitWithBarrier.cpp urEnqueueKernelLaunch.cpp + urEnqueueKernelLaunchAndMemcpyInOrder.cpp urEnqueueMemBufferCopyRect.cpp urEnqueueMemBufferCopy.cpp urEnqueueMemBufferFill.cpp diff --git a/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match b/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match index bf20fe2e0f..e48c5175b4 100644 --- a/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match +++ b/test/conformance/enqueue/enqueue_adapter_level_zero_v2.match @@ -1,37 +1,14 @@ urEnqueueDeviceGetGlobalVariableReadTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitWithBarrierTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitWithBarrierTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueEventsWaitWithBarrierTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchTest.InvalidWorkGroupSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchTest.InvalidKernelArgs/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchKernelWgSizeTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -{{OPT}}urEnqueueKernelLaunchKernelWgSizeTest.SuccessWithExplicitLocalSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchKernelSubGroupTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchKernelStandardTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D_1 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D_31 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D_1027 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D_32 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D_256 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D_1_1 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D_31_7 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D_1027_1 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D_1_32 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D_256_79 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___3D_1_1_1 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___3D_31_7_1 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___3D_1027_1_19 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___3D_1_53_19 -urEnqueueKernelLaunchTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___3D_256_79_8 urEnqueueKernelLaunchWithVirtualMemory.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueKernelLaunchMultiDeviceTest.KernelLaunchReadDifferentQueues/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UsePoolEnabled urEnqueueKernelLaunchUSMLinkedList.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UsePoolDisabled +{{OPT}}urEnqueueKernelLaunchIncrementTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UseEventsEnabled +{{OPT}}urEnqueueKernelLaunchIncrementTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UseEventsDisabled +{{OPT}}urEnqueueKernelLaunchIncrementMultiDeviceMultiThreadTest.Success/UseEventsNoQueuePerThread +{{OPT}}urEnqueueKernelLaunchIncrementMultiDeviceMultiThreadTest.Success/NoUseEventsNoQueuePerThread urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_whole_buffer_2D urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_non_zero_offsets_2D urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_different_buffer_sizes_2D @@ -41,61 +18,7 @@ urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_ urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_3d_with_offsets urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_2d_3d urEnqueueMemBufferCopyRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___copy_3d_2d -urEnqueueMemBufferCopyRectTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferCopyRectTest.InvalidNullHandleBufferSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferCopyRectTest.InvalidNullHandleBufferDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferCopyRectTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueMemBufferCopyRectTest.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferCopyTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.InvalidNullHandleBufferDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferCopyTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferCopyTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024 -urEnqueueMemBufferCopyTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500 -urEnqueueMemBufferCopyTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096 -urEnqueueMemBufferCopyTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1__patternSize__1 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__256 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1024__patternSize__256 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__4 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__8 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__16 -urEnqueueMemBufferFillTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__32 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1__patternSize__1 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__256 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1024__patternSize__256 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__4 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__8 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__16 -urEnqueueMemBufferFillTest.SuccessPartialFill/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__32 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1__patternSize__1 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__256 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1024__patternSize__256 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__4 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__8 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__16 -urEnqueueMemBufferFillTest.SuccessOffset/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__32 -urEnqueueMemBufferFillNegativeTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferFillNegativeTest.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferFillNegativeTest.InvalidNullHandlePointerPattern/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferFillNegativeTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferFillNegativeTest.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueMemBufferMapTestWithParam.SuccessRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE urEnqueueMemBufferMapTestWithParam.SuccessRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE urEnqueueMemBufferMapTestWithParam.SuccessRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE @@ -192,86 +115,6 @@ urEnqueueMemBufferMapTestWithParam.SuccessMultiMaps/Intel_R__oneAPI_Unified_Runt urEnqueueMemBufferMapTestWithParam.SuccessMultiMaps/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferMapTestWithParam.SuccessMultiMaps/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferMapTestWithParam.SuccessMultiMaps/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidEnumerationMapFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPointerRetMap/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferMapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferMapTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE urEnqueueMemBufferMapTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE urEnqueueMemBufferMapTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE @@ -290,134 +133,6 @@ urEnqueueMemBufferMapTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_o urEnqueueMemBufferMapTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferMapTestWithWriteFlagParam.SuccessWrite/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MAP_FLAG_WRITE urEnqueueMemBufferMapTestWithWriteFlagParam.SuccessWrite/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MAP_FLAG_WRITE_INVALIDATE_REGION -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferReadTestWithParam.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_whole_buffer_2D urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_non_zero_offsets_2D urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_different_buffer_sizes_2D @@ -427,123 +142,7 @@ urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_ urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_3d_with_offsets urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_2d_3d urEnqueueMemBufferReadRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_3d_2d -urEnqueueMemBufferReadRectTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferReadRectTest.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferReadRectTest.InvalidNullPointerDst/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferReadRectTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueMemBufferReadRectTest.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.SuccessWriteRead/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_WRITE -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_WRITE_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_READ_ONLY -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1024_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueMemBufferWriteTestWithParam.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_whole_buffer_2D urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_non_zero_offsets_2D urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_different_buffer_sizes_2D @@ -553,10 +152,6 @@ urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_3d_with_offsets urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_2d_3d urEnqueueMemBufferWriteRectTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___write_3d_2d -urEnqueueMemBufferWriteRectTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferWriteRectTest.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferWriteRectTest.InvalidNullPointerSrc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueMemBufferWriteRectTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueMemBufferWriteRectTest.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueMemImageCopyTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___1D urEnqueueMemImageCopyTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2D @@ -691,159 +286,11 @@ urEnqueueMemUnmapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unifi urEnqueueMemUnmapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___2500_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemUnmapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___4096_UR_MEM_FLAG_ALLOC_HOST_POINTER urEnqueueMemUnmapTestWithParam.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___6000_UR_MEM_FLAG_ALLOC_HOST_POINTER -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1__patternSize__1 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__256 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__1024__patternSize__256 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__4 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__8 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__16 -urEnqueueUSMFillTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___size__256__patternSize__32 -urEnqueueUSMFillNegativeTest.InvalidEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__patternSize__1 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__patternSize__256 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__patternSize__4 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__57__height__1__patternSize__1 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__patternSize__256 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__patternSize__1024 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__patternSize__1 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__patternSize__256 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__patternSize__65536 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__patternSize__1 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__35__patternSize__1 -urEnqueueUSMFill2DTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__35__patternSize__128 -urEnqueueUSMFill2DNegativeTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueUSMAdviseWithParamTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_USM_ADVICE_FLAG_DEFAULT urEnqueueUSMAdviseTest.MultipleParamsSuccess/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMAdviseTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMAdviseTest.InvalidNullPointerMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMAdviseTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMAdviseTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMAdviseTest.InvalidSizeTooLarge/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueUSMAdviseTest.NonCoherentDeviceMemorySuccessOrWarning/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMMemcpyTest.Blocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMMemcpyTest.BlockingWithEvent/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMMemcpyTest.NonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMMemcpyTest.WaitForDependencies/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__1024__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1024__width__256__height__256__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__23__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_HOST__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_HOST -urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__234__width__233__height__1__src__UR_USM_TYPE_SHARED__dst__UR_USM_TYPE_SHARED -urEnqueueUSMMemcpy2DNegativeTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DNegativeTest.InvalidNullPointer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DNegativeTest.InvalidSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE -urEnqueueUSMMemcpy2DNegativeTest.InvalidEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___pitch__1__width__1__height__1__src__UR_USM_TYPE_DEVICE__dst__UR_USM_TYPE_DEVICE urEnqueueUSMPrefetchWithParamTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT urEnqueueUSMPrefetchWithParamTest.CheckWaitEvent/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT -urEnqueueUSMPrefetchTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMPrefetchTest.InvalidNullPointerMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMPrefetchTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMPrefetchTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMPrefetchTest.InvalidSizeTooLarge/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueUSMPrefetchTest.InvalidEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueReadHostPipeTest.InvalidNullHandleQueue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueReadHostPipeTest.InvalidNullHandleProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueReadHostPipeTest.InvalidNullPointerPipeSymbol/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ @@ -856,4 +303,3 @@ urEnqueueWriteHostPipeTest.InvalidNullPointerBuffer/Intel_R__oneAPI_Unified_Runt urEnqueueWriteHostPipeTest.InvalidEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueTimestampRecordingExpTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEnqueueTimestampRecordingExpTest.SuccessBlocking/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEnqueueTimestampRecordingExpTest.InvalidNullPtrEventWaitList/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ diff --git a/test/conformance/enqueue/enqueue_adapter_native_cpu.match b/test/conformance/enqueue/enqueue_adapter_native_cpu.match index 83e9f2391e..fc3cf2d975 100644 --- a/test/conformance/enqueue/enqueue_adapter_native_cpu.match +++ b/test/conformance/enqueue/enqueue_adapter_native_cpu.match @@ -45,6 +45,7 @@ {{OPT}}urEnqueueKernelLaunchTestWithParam.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__3D_1_53_19 {{OPT}}urEnqueueKernelLaunchTestWithParam.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__3D_256_79_8 {{OPT}}urEnqueueKernelLaunchWithVirtualMemory.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} +{{OPT}}urEnqueueKernelLaunchWithUSM.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} {{OPT}}urEnqueueKernelLaunchMultiDeviceTest.KernelLaunchReadDifferentQueues/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} {{OPT}}urEnqueueKernelLaunchUSMLinkedList.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__UsePoolEnabled {{OPT}}urEnqueueKernelLaunchUSMLinkedList.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__UsePoolDisabled diff --git a/test/conformance/enqueue/urEnqueueKernelLaunchAndMemcpyInOrder.cpp b/test/conformance/enqueue/urEnqueueKernelLaunchAndMemcpyInOrder.cpp new file mode 100644 index 0000000000..2aee087c73 --- /dev/null +++ b/test/conformance/enqueue/urEnqueueKernelLaunchAndMemcpyInOrder.cpp @@ -0,0 +1,392 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include +#include + +#include +#include + +// There was a bug in previous L0 drivers that caused the test to fail +std::tuple minL0DriverVersion = {1, 3, 29534}; + +template +struct urMultiQueueLaunchMemcpyTest : uur::urMultiDeviceContextTestTemplate<1>, + testing::WithParamInterface { + std::string KernelName; + std::vector programs; + std::vector kernels; + std::vector SharedMem; + + std::vector queues; + std::vector devices; + + std::function createQueues; + + static constexpr char ProgramName[] = "increment"; + static constexpr size_t ArraySize = 100; + static constexpr size_t InitialValue = 1; + + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE( + uur::urMultiDeviceContextTestTemplate<1>::SetUp()); + + createQueues(); + + for (auto &device : devices) { + SKIP_IF_DRIVER_TOO_OLD("Level-Zero", minL0DriverVersion, platform, + device); + } + + programs.resize(devices.size()); + kernels.resize(devices.size()); + SharedMem.resize(devices.size()); + + KernelName = uur::KernelsEnvironment::instance->GetEntryPointNames( + ProgramName)[0]; + + std::shared_ptr> il_binary; + std::vector metadatas{}; + + uur::KernelsEnvironment::instance->LoadSource(ProgramName, il_binary); + + for (size_t i = 0; i < devices.size(); i++) { + const ur_program_properties_t properties = { + UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES, nullptr, + static_cast(metadatas.size()), + metadatas.empty() ? nullptr : metadatas.data()}; + + uur::raii::Program program; + ASSERT_SUCCESS(uur::KernelsEnvironment::instance->CreateProgram( + platform, context, devices[i], *il_binary, &properties, + &programs[i])); + + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED( + urProgramBuild(context, programs[i], nullptr)); + ASSERT_SUCCESS( + urKernelCreate(programs[i], KernelName.data(), &kernels[i])); + + ASSERT_SUCCESS( + urUSMSharedAlloc(context, devices[i], nullptr, nullptr, + ArraySize * sizeof(uint32_t), &SharedMem[i])); + + ASSERT_SUCCESS(urEnqueueUSMFill(queues[i], SharedMem[i], + sizeof(uint32_t), &InitialValue, + ArraySize * sizeof(uint32_t), 0, + nullptr, nullptr /* &Event */)); + ASSERT_SUCCESS(urQueueFinish(queues[i])); + + ASSERT_SUCCESS( + urKernelSetArgPointer(kernels[i], 0, nullptr, SharedMem[i])); + } + } + + void TearDown() override { + for (auto &Ptr : SharedMem) { + urUSMFree(context, Ptr); + } + for (const auto &queue : queues) { + EXPECT_SUCCESS(urQueueRelease(queue)); + } + for (const auto &kernel : kernels) { + urKernelRelease(kernel); + } + for (const auto &program : programs) { + urProgramRelease(program); + } + UUR_RETURN_ON_FATAL_FAILURE( + uur::urMultiDeviceContextTestTemplate<1>::TearDown()); + } +}; + +template +struct urEnqueueKernelLaunchIncrementMultiDeviceTestWithParam + : public urMultiQueueLaunchMemcpyTest { + static constexpr size_t duplicateDevices = 8; + + using urMultiQueueLaunchMemcpyTest::context; + using urMultiQueueLaunchMemcpyTest::queues; + using urMultiQueueLaunchMemcpyTest::devices; + using urMultiQueueLaunchMemcpyTest::kernels; + using urMultiQueueLaunchMemcpyTest::SharedMem; + + void SetUp() override { + this->createQueues = [&] { + for (size_t i = 0; i < duplicateDevices; i++) { + devices.insert( + devices.end(), + uur::KernelsEnvironment::instance->devices.begin(), + uur::KernelsEnvironment::instance->devices.end()); + } + + for (auto &device : devices) { + ur_queue_handle_t queue = nullptr; + ASSERT_SUCCESS(urQueueCreate(context, device, 0, &queue)); + queues.push_back(queue); + } + }; + + UUR_RETURN_ON_FATAL_FAILURE( + urMultiQueueLaunchMemcpyTest::SetUp()); + } + + void TearDown() override { + UUR_RETURN_ON_FATAL_FAILURE( + urMultiQueueLaunchMemcpyTest::TearDown()); + } +}; + +struct urEnqueueKernelLaunchIncrementTest + : urMultiQueueLaunchMemcpyTest< + std::tuple> { + static constexpr size_t numOps = 50; + + ur_queue_handle_t queue; + + using Param = std::tuple; + using urMultiQueueLaunchMemcpyTest::context; + using urMultiQueueLaunchMemcpyTest::queues; + using urMultiQueueLaunchMemcpyTest::devices; + using urMultiQueueLaunchMemcpyTest::kernels; + using urMultiQueueLaunchMemcpyTest::SharedMem; + + void SetUp() override { + auto device = std::get<0>(GetParam()); + + this->createQueues = [&] { + ASSERT_SUCCESS(urQueueCreate(context, device, 0, &queue)); + + // use the same queue and device for all operations + for (size_t i = 0; i < numOps; i++) { + urQueueRetain(queue); + + queues.push_back(queue); + devices.push_back(device); + } + }; + + UUR_RETURN_ON_FATAL_FAILURE( + urMultiQueueLaunchMemcpyTest::SetUp()); + } + + void TearDown() override { + urQueueRelease(queue); + UUR_RETURN_ON_FATAL_FAILURE( + urMultiQueueLaunchMemcpyTest::TearDown()); + } +}; + +UUR_TEST_SUITE_P( + urEnqueueKernelLaunchIncrementTest, + testing::ValuesIn(uur::BoolTestParam::makeBoolParam("UseEvents")), + uur::deviceTestWithParamPrinter); + +TEST_P(urEnqueueKernelLaunchIncrementTest, Success) { + constexpr size_t global_offset = 0; + constexpr size_t n_dimensions = 1; + + auto useEvents = std::get<1>(GetParam()).value; + + std::vector Events(numOps * 2); + for (size_t i = 0; i < numOps; i++) { + size_t waitNum = 0; + ur_event_handle_t *lastEvent = nullptr; + ur_event_handle_t *kernelEvent = nullptr; + ur_event_handle_t *memcpyEvent = nullptr; + + if (useEvents) { + // Events are: kernelEvent0, memcpyEvent0, kernelEvent1, ... + waitNum = i > 0 ? 1 : 0; + lastEvent = i > 0 ? Events[i * 2 - 1].ptr() : nullptr; + + kernelEvent = Events[i * 2].ptr(); + memcpyEvent = Events[i * 2 + 1].ptr(); + } + + // execute kernel that increments each element by 1 + ASSERT_SUCCESS(urEnqueueKernelLaunch( + queue, kernels[i], n_dimensions, &global_offset, &ArraySize, + nullptr, waitNum, lastEvent, kernelEvent)); + + // copy the memory (input for the next kernel) + if (i < numOps - 1) { + ASSERT_SUCCESS( + urEnqueueUSMMemcpy(queue, false, SharedMem[i + 1], SharedMem[i], + ArraySize * sizeof(uint32_t), useEvents, + kernelEvent, memcpyEvent)); + } + } + + if (useEvents) { + // TODO: just wait on the last event, once urEventWait is implemented + // by V2 L0 adapter + urQueueFinish(queue); + } else { + urQueueFinish(queue); + } + + size_t ExpectedValue = InitialValue; + for (size_t i = 0; i < numOps; i++) { + ExpectedValue++; + for (uint32_t j = 0; j < ArraySize; ++j) { + ASSERT_EQ(reinterpret_cast(SharedMem[i])[j], + ExpectedValue); + } + } +} + +struct VoidParam {}; +using urEnqueueKernelLaunchIncrementMultiDeviceTest = + urEnqueueKernelLaunchIncrementMultiDeviceTestWithParam; + +// Do a chain of kernelLaunch(dev0) -> memcpy(dev0, dev1) -> kernelLaunch(dev1) ... ops +TEST_F(urEnqueueKernelLaunchIncrementMultiDeviceTest, Success) { + size_t returned_size; + ASSERT_SUCCESS(urDeviceGetInfo(devices[0], UR_DEVICE_INFO_EXTENSIONS, 0, + nullptr, &returned_size)); + + std::unique_ptr returned_extensions(new char[returned_size]); + + ASSERT_SUCCESS(urDeviceGetInfo(devices[0], UR_DEVICE_INFO_EXTENSIONS, + returned_size, returned_extensions.get(), + nullptr)); + + std::string_view extensions_string(returned_extensions.get()); + const bool usm_p2p_support = + extensions_string.find(UR_USM_P2P_EXTENSION_STRING_EXP) != + std::string::npos; + + if (!usm_p2p_support) { + GTEST_SKIP() << "EXP usm p2p feature is not supported."; + } + + constexpr size_t global_offset = 0; + constexpr size_t n_dimensions = 1; + + std::vector Events(devices.size() * 2); + for (size_t i = 0; i < devices.size(); i++) { + // Events are: kernelEvent0, memcpyEvent0, kernelEvent1, ... + size_t waitNum = i > 0 ? 1 : 0; + ur_event_handle_t *lastEvent = + i > 0 ? Events[i * 2 - 1].ptr() : nullptr; + ur_event_handle_t *kernelEvent = Events[i * 2].ptr(); + ur_event_handle_t *memcpyEvent = Events[i * 2 + 1].ptr(); + + // execute kernel that increments each element by 1 + ASSERT_SUCCESS(urEnqueueKernelLaunch( + queues[i], kernels[i], n_dimensions, &global_offset, &ArraySize, + nullptr, waitNum, lastEvent, kernelEvent)); + + // copy the memory to next device + if (i < devices.size() - 1) { + ASSERT_SUCCESS(urEnqueueUSMMemcpy( + queues[i], false, SharedMem[i + 1], SharedMem[i], + ArraySize * sizeof(uint32_t), 1, kernelEvent, memcpyEvent)); + } + } + + // synchronize on the last queue only, this has to ensure all the operations + // are completed + urQueueFinish(queues.back()); + + size_t ExpectedValue = InitialValue; + for (size_t i = 0; i < devices.size(); i++) { + ExpectedValue++; + for (uint32_t j = 0; j < ArraySize; ++j) { + ASSERT_EQ(reinterpret_cast(SharedMem[i])[j], + ExpectedValue); + } + } +} + +template +inline std::string +printParams(const testing::TestParamInfo &info) { + std::stringstream ss; + + auto param1 = std::get<0>(info.param); + auto param2 = std::get<1>(info.param); + + ss << (param1.value ? "" : "No") << param1.name; + ss << (param2.value ? "" : "No") << param2.name; + + return ss.str(); +} + +using urEnqueueKernelLaunchIncrementMultiDeviceMultiThreadTest = + urEnqueueKernelLaunchIncrementMultiDeviceTestWithParam< + std::tuple>; + +INSTANTIATE_TEST_SUITE_P( + , urEnqueueKernelLaunchIncrementMultiDeviceMultiThreadTest, + testing::Combine( + testing::ValuesIn(uur::BoolTestParam::makeBoolParam("UseEvents")), + testing::ValuesIn(uur::BoolTestParam::makeBoolParam("QueuePerThread"))), + printParams); + +// Enqueue kernelLaunch concurrently from multiple threads +// With !queuePerThread this becomes a test on a single device +TEST_P(urEnqueueKernelLaunchIncrementMultiDeviceMultiThreadTest, Success) { + size_t numThreads = devices.size(); + std::vector threads; + + static constexpr size_t numOpsPerThread = 6; + + auto useEvents = std::get<0>(GetParam()).value; + auto queuePerThread = std::get<1>(GetParam()).value; + + for (size_t i = 0; i < numThreads; i++) { + threads.emplace_back([this, i, queuePerThread, useEvents]() { + constexpr size_t global_offset = 0; + constexpr size_t n_dimensions = 1; + + auto queue = queuePerThread ? queues[i] : queues.back(); + auto kernel = kernels[i]; + auto sharedPtr = SharedMem[i]; + + std::vector Events(numOpsPerThread + 1); + for (size_t j = 0; j < numOpsPerThread; j++) { + size_t waitNum = 0; + ur_event_handle_t *lastEvent = nullptr; + ur_event_handle_t *signalEvent = nullptr; + + if (useEvents) { + waitNum = j > 0 ? 1 : 0; + lastEvent = j > 0 ? Events[j - 1].ptr() : nullptr; + signalEvent = Events[j].ptr(); + } + + // execute kernel that increments each element by 1 + ASSERT_SUCCESS(urEnqueueKernelLaunch( + queue, kernel, n_dimensions, &global_offset, &ArraySize, + nullptr, waitNum, lastEvent, signalEvent)); + } + + std::vector data(ArraySize); + + auto lastEvent = + useEvents ? Events[numOpsPerThread - 1].ptr() : nullptr; + auto signalEvent = useEvents ? Events.back().ptr() : nullptr; + ASSERT_SUCCESS( + urEnqueueUSMMemcpy(queue, false, data.data(), sharedPtr, + ArraySize * sizeof(uint32_t), useEvents, + lastEvent, signalEvent)); + + urQueueFinish(queue); + // TODO: when useEvents is implemented for L0 v2 adapter + // wait on event instead + + size_t ExpectedValue = InitialValue; + ExpectedValue += numOpsPerThread; + for (uint32_t j = 0; j < ArraySize; ++j) { + ASSERT_EQ(data[j], ExpectedValue); + } + }); + } + + for (auto &thread : threads) { + thread.join(); + } +} diff --git a/test/conformance/event/event_adapter_level_zero_v2.match b/test/conformance/event/event_adapter_level_zero_v2.match index a9d97d5044..e3f93c54c3 100644 --- a/test/conformance/event/event_adapter_level_zero_v2.match +++ b/test/conformance/event/event_adapter_level_zero_v2.match @@ -1,15 +1,9 @@ - urEventGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_EVENT_INFO_COMMAND_QUEUE urEventGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_EVENT_INFO_CONTEXT urEventGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_EVENT_INFO_COMMAND_TYPE -urEventGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_EVENT_INFO_COMMAND_EXECUTION_STATUS -urEventGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_EVENT_INFO_REFERENCE_COUNT urEventGetInfoNegativeTest.InvalidNullHandle/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetInfoNegativeTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventGetInfoNegativeTest.InvalidSizePropSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventGetInfoNegativeTest.InvalidSizePropSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetInfoNegativeTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetInfoNegativeTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventGetProfilingInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_PROFILING_INFO_COMMAND_QUEUED urEventGetProfilingInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_PROFILING_INFO_COMMAND_SUBMIT urEventGetProfilingInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_PROFILING_INFO_COMMAND_START @@ -17,22 +11,8 @@ urEventGetProfilingInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Z urEventGetProfilingInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_PROFILING_INFO_COMMAND_COMPLETE urEventGetProfilingInfoWithTimingComparisonTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventGetProfilingInfoNegativeTest.InvalidNullHandle/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetProfilingInfoNegativeTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventGetProfilingInfoNegativeTest.InvalidValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventWaitTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventRetainTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventRetainTest.InvalidNullHandle/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventReleaseTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventReleaseTest.InvalidNullHandle/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetNativeHandleTest.InvalidNullHandleEvent/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventGetNativeHandleTest.InvalidNullPointerNativeEvent/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventCreateWithNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventSetCallbackTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventSetCallbackTest.ValidateParameters/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventSetCallbackTest.AllStates/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urEventSetCallbackTest.EventAlreadyCompleted/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventSetCallbackNegativeTest.InvalidNullHandleEvent/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventSetCallbackNegativeTest.InvalidNullPointerCallback/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urEventSetCallbackNegativeTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -{{OPT}}{{Segmentation fault|Aborted}} diff --git a/test/conformance/exp_command_buffer/fixtures.h b/test/conformance/exp_command_buffer/fixtures.h index c144ac5fa2..9a38772eb7 100644 --- a/test/conformance/exp_command_buffer/fixtures.h +++ b/test/conformance/exp_command_buffer/fixtures.h @@ -154,7 +154,8 @@ struct urUpdatableCommandBufferExpExecutionTest // Create a command-buffer with update enabled. ur_exp_command_buffer_desc_t desc{ - UR_STRUCTURE_TYPE_EXP_COMMAND_BUFFER_DESC, nullptr, true}; + UR_STRUCTURE_TYPE_EXP_COMMAND_BUFFER_DESC, nullptr, true, false, + false}; ASSERT_SUCCESS(urCommandBufferCreateExp(context, device, &desc, &updatable_cmd_buf_handle)); diff --git a/test/conformance/kernel/kernel_adapter_level_zero_v2.match b/test/conformance/kernel/kernel_adapter_level_zero_v2.match index dca885cc4d..074a58720a 100644 --- a/test/conformance/kernel/kernel_adapter_level_zero_v2.match +++ b/test/conformance/kernel/kernel_adapter_level_zero_v2.match @@ -1,85 +1,63 @@ -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE -urKernelGetGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE -urKernelGetGroupInfoSingleTest.CompileWorkGroupSizeEmpty/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelGetGroupInfoWgSizeTest.CompileWorkGroupSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_FUNCTION_NAME -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_ARGS -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_REFERENCE_COUNT -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_CONTEXT -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_PROGRAM -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_ATTRIBUTES -urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_REGS -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_FUNCTION_NAME -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_ARGS -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_REFERENCE_COUNT -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_CONTEXT -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_PROGRAM -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_ATTRIBUTES -urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_REGS -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_FUNCTION_NAME -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_ARGS -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_REFERENCE_COUNT -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_CONTEXT -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_PROGRAM -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_ATTRIBUTES -urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_REGS -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_FUNCTION_NAME -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_ARGS -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_REFERENCE_COUNT -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_CONTEXT -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_PROGRAM -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_ATTRIBUTES -urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_INFO_NUM_REGS -urKernelGetInfoSingleTest.KernelNameCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelGetInfoSingleTest.KernelContextCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelGetSubGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE -urKernelGetSubGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS -urKernelGetSubGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS -urKernelGetSubGroupInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL -urKernelGetSubGroupInfoSingleTest.CompileNumSubgroupsIsZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgLocalTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgLocalTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgMemObjTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgMemObjTest.InvalidNullHandleKernel/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgMemObjTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -{{OPT}}urKernelSetArgPointerTest.SuccessHost/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -{{OPT}}urKernelSetArgPointerTest.SuccessDevice/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -{{OPT}}urKernelSetArgPointerTest.SuccessShared/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgPointerNegativeTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST -urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR -urKernelSetArgSamplerTest.SuccessWithProps/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgSamplerTest.InvalidNullHandleKernel/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgSamplerTest.InvalidNullHandleArgValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgSamplerTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgValueTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetArgValueTest.InvalidKernelArgumentSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetExecInfoTest.SuccessIndirectAccess/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetExecInfoUSMPointersTest.SuccessHost/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetExecInfoUSMPointersTest.SuccessDevice/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetExecInfoUSMPointersTest.SuccessShared/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urKernelSetExecInfoCacheConfigTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_CACHE_CONFIG_DEFAULT -urKernelSetExecInfoCacheConfigTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_CACHE_CONFIG_LARGE_SLM -urKernelSetExecInfoCacheConfigTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_KERNEL_CACHE_CONFIG_LARGE_DATA +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_ARGS +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_CONTEXT +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_PROGRAM +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES +urKernelGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_ARGS +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_CONTEXT +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_PROGRAM +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES +urKernelGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_ARGS +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_CONTEXT +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_PROGRAM +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES +urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_ARGS +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_CONTEXT +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_PROGRAM +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES +urKernelGetInfoTest.InvalidNullPointerPropValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoSingleTest.KernelNameCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelGetInfoSingleTest.KernelContextCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgLocalTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgMemObjTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgPointerNegativeTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___NORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_NONE_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_CLAMP_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_NEAREST +urKernelSetArgSamplerTestWithParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UNNORMALIZED_UR_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT_UR_SAMPLER_FILTER_MODE_LINEAR +urKernelSetArgSamplerTest.SuccessWithProps/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgSamplerTest.InvalidNullHandleKernel/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgSamplerTest.InvalidNullHandleArgValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgSamplerTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgValueTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetArgValueTest.InvalidKernelArgumentSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetExecInfoTest.SuccessIndirectAccess/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetExecInfoUSMPointersTest.SuccessHost/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetExecInfoUSMPointersTest.SuccessDevice/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelSetExecInfoUSMPointersTest.SuccessShared/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ diff --git a/test/conformance/memory/memory_adapter_level_zero_v2.match b/test/conformance/memory/memory_adapter_level_zero_v2.match index 7e92763def..e6639680ed 100644 --- a/test/conformance/memory/memory_adapter_level_zero_v2.match +++ b/test/conformance/memory/memory_adapter_level_zero_v2.match @@ -1,38 +1,12 @@ -urMemBufferCreateWithFlagsTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_READ_WRITE -urMemBufferCreateWithFlagsTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_WRITE_ONLY -urMemBufferCreateWithFlagsTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_READ_ONLY -urMemBufferCreateWithFlagsTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_ALLOC_HOST_POINTER -urMemBufferCreateWithHostPtrFlagsTest.SUCCESS/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_ALLOC_COPY_HOST_POINTER -urMemBufferCreateWithHostPtrFlagsTest.SUCCESS/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_FLAG_USE_HOST_POINTER -urMemBufferCreateWithNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urMemBufferPartitionTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidNullHandleBuffer/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidEnumerationFlags/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidEnumerationBufferCreateType/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidNullPointerBufferCreateInfo/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidNullPointerMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemBufferPartitionTest.InvalidBufferSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urMemBufferPartitionTest.InvalidValueCreateType/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urMemBufferPartitionTest.InvalidValueBufferCreateInfoOutOfBounds/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urMemGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE urMemGetInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetInfoTest.InvalidNullHandleMemory/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE -urMemGetInfoTest.InvalidNullHandleMemory/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetInfoTest.InvalidEnumerationMemInfoType/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE -urMemGetInfoTest.InvalidEnumerationMemInfoType/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE -urMemGetInfoTest.InvalidSizeZero/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT urMemGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE urMemGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetInfoTest.InvalidNullPointerParamValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE -urMemGetInfoTest.InvalidNullPointerParamValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE -urMemGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT urMemGetInfoImageTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_SIZE urMemGetInfoImageTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_MEM_INFO_CONTEXT -urMemGetNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemGetNativeHandleTest.InvalidNullHandleMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemGetNativeHandleTest.InvalidNullPointerNativeMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urMemImageCreateTestWithImageFormatParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_CHANNEL_ORDER_A__UR_IMAGE_CHANNEL_TYPE_SNORM_INT8 urMemImageCreateTestWithImageFormatParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_CHANNEL_ORDER_A__UR_IMAGE_CHANNEL_TYPE_SNORM_INT16 urMemImageCreateTestWithImageFormatParam.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_CHANNEL_ORDER_A__UR_IMAGE_CHANNEL_TYPE_UNORM_INT8 @@ -307,7 +281,3 @@ urMemImageGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runt urMemImageGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_INFO_WIDTH urMemImageGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_INFO_HEIGHT urMemImageGetInfoTest.InvalidNullPointerPropSizeRet/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_IMAGE_INFO_DEPTH -urMemReleaseTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemReleaseTest.InvalidNullHandleMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemRetainTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urMemRetainTest.InvalidNullHandleMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ diff --git a/test/conformance/program/program_adapter_level_zero_v2.match b/test/conformance/program/program_adapter_level_zero_v2.match index 7c53e24502..70e0a12609 100644 --- a/test/conformance/program/program_adapter_level_zero_v2.match +++ b/test/conformance/program/program_adapter_level_zero_v2.match @@ -1,12 +1,8 @@ -urProgramCreateWithNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramGetBuildInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_PROGRAM_BUILD_INFO_STATUS -urProgramGetFunctionPointerTest.InvalidKernelName/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramGetNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ +urProgramCreateWithNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramGetBuildInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS +urProgramGetFunctionPointerTest.InvalidKernelName/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramGetNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ {{OPT}}urProgramLinkErrorTest.LinkFailure/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ {{OPT}}urProgramLinkErrorTest.SetOutputOnLinkError/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ -urProgramSetSpecializationConstantsTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramSetSpecializationConstantsTest.UseDefaultValue/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramSetMultipleSpecializationConstantsTest.MultipleCalls/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urProgramSetMultipleSpecializationConstantsTest.SingleCall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ diff --git a/test/conformance/program/urProgramCreateWithIL.cpp b/test/conformance/program/urProgramCreateWithIL.cpp index 3d81d14104..7c02c3c7b9 100644 --- a/test/conformance/program/urProgramCreateWithIL.cpp +++ b/test/conformance/program/urProgramCreateWithIL.cpp @@ -37,7 +37,8 @@ TEST_P(urProgramCreateWithILTest, Success) { } TEST_P(urProgramCreateWithILTest, SuccessWithProperties) { - ur_program_properties_t properties{UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES}; + ur_program_properties_t properties{UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES, + nullptr, 0, nullptr}; ur_program_handle_t program = nullptr; ASSERT_SUCCESS(urProgramCreateWithIL( context, il_binary->data(), il_binary->size(), &properties, &program)); diff --git a/test/conformance/queue/queue_adapter_level_zero_v2.match b/test/conformance/queue/queue_adapter_level_zero_v2.match index 524e04fa1d..e69de29bb2 100644 --- a/test/conformance/queue/queue_adapter_level_zero_v2.match +++ b/test/conformance/queue/queue_adapter_level_zero_v2.match @@ -1,2 +0,0 @@ -urQueueFinishTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urQueueFlushTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ diff --git a/test/conformance/testing/include/uur/fixtures.h b/test/conformance/testing/include/uur/fixtures.h index e57a31584a..568f700da1 100644 --- a/test/conformance/testing/include/uur/fixtures.h +++ b/test/conformance/testing/include/uur/fixtures.h @@ -513,11 +513,12 @@ struct urMultiQueueTest : urContextTest { ur_queue_handle_t queue2 = nullptr; }; -struct urMultiDeviceContextTest : urPlatformTest { +template +struct urMultiDeviceContextTestTemplate : urPlatformTest { void SetUp() override { UUR_RETURN_ON_FATAL_FAILURE(urPlatformTest::SetUp()); auto &devices = DevicesEnvironment::instance->devices; - if (devices.size() <= 1) { + if (devices.size() < MinDevices) { GTEST_SKIP(); } ASSERT_SUCCESS(urContextCreate(static_cast(devices.size()), @@ -534,6 +535,10 @@ struct urMultiDeviceContextTest : urPlatformTest { ur_context_handle_t context = nullptr; }; +struct urMultiDeviceContextTest : urMultiDeviceContextTestTemplate<> { + using urMultiDeviceContextTestTemplate::context; +}; + struct urMultiDeviceMemBufferTest : urMultiDeviceContextTest { void SetUp() override { UUR_RETURN_ON_FATAL_FAILURE(urMultiDeviceContextTest::SetUp()); diff --git a/test/conformance/testing/include/uur/utils.h b/test/conformance/testing/include/uur/utils.h index 7e23e55843..a1febf9769 100644 --- a/test/conformance/testing/include/uur/utils.h +++ b/test/conformance/testing/include/uur/utils.h @@ -424,6 +424,65 @@ ur_result_t MakeUSMAllocationByType(ur_usm_type_t type, ur_usm_pool_handle_t hPool, size_t size, void **ppMem); +inline std::tuple +decodeSemVersion(std::string version) { + auto posMajor = version.find('.'); + auto posMinor = version.find('.', posMajor + 1); + auto major = std::stoi(version.substr(0, posMajor)); + auto minor = + std::stoi(version.substr(posMajor + 1, posMinor - posMajor - 1)); + auto patch = std::stoi(version.substr(posMinor + 1)); + return std::make_tuple(major, minor, patch); +} + +inline bool isGivenAdapter(ur_platform_handle_t hPlatform, + std::string adapterName) { + size_t psize; + EXPECT_EQ( + urPlatformGetInfo(hPlatform, UR_PLATFORM_INFO_NAME, 0, nullptr, &psize), + UR_RESULT_SUCCESS); + std::string platform(psize, '\0'); + EXPECT_EQ(urPlatformGetInfo(hPlatform, UR_PLATFORM_INFO_NAME, psize, + platform.data(), nullptr), + UR_RESULT_SUCCESS); + + return platform.find(adapterName) != std::string::npos; +} + +inline std::tuple +getDriverVersion(ur_device_handle_t hDevice) { + size_t driverVersionSize = 0; + EXPECT_EQ(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_DRIVER_VERSION, 0, + nullptr, &driverVersionSize), + UR_RESULT_SUCCESS); + std::string driver(driverVersionSize, '\0'); + EXPECT_EQ(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_DRIVER_VERSION, + driverVersionSize, driver.data(), + &driverVersionSize), + UR_RESULT_SUCCESS); + + return decodeSemVersion(driver); +} + +#define SKIP_IF_DRIVER_TOO_OLD(adapterName, minDriverVersion, hPlatform, \ + hDevice) \ + do { \ + if (uur::isGivenAdapter(hPlatform, adapterName)) { \ + auto [major, minor, patch] = uur::getDriverVersion(hDevice); \ + auto [minMajor, minMinor, minPatch] = minL0DriverVersion; \ + if (major < minMajor || (major == minMajor && minor < minMinor) || \ + (major == minMajor && minor == minMinor && \ + patch < minPatch)) { \ + GTEST_SKIP() \ + << "Skipping test because driver version is too old for " \ + << adapterName << ". " \ + << "Driver version: " << major << "." << minor << "." \ + << patch << " Minimum required version: " << minMajor \ + << "." << minMinor << "." << minPatch; \ + } \ + } \ + } while (0) + } // namespace uur #endif // UR_CONFORMANCE_INCLUDE_UTILS_H_INCLUDED diff --git a/test/conformance/usm/usm_adapter_level_zero_v2.match b/test/conformance/usm/usm_adapter_level_zero_v2.match index 88501eea9d..0908da40da 100644 --- a/test/conformance/usm/usm_adapter_level_zero_v2.match +++ b/test/conformance/usm/usm_adapter_level_zero_v2.match @@ -1,122 +1,7 @@ -urUSMDeviceAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMDeviceAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMDeviceAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMDeviceAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled urUSMDeviceAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled urUSMDeviceAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_2048 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_8 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_512 -urUSMDeviceAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_2048 -urUSMFreeTest.SuccessDeviceAlloc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMFreeTest.SuccessHostAlloc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMFreeTest.SuccessSharedAlloc/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMGetMemAllocInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_USM_ALLOC_INFO_TYPE -urUSMGetMemAllocInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_USM_ALLOC_INFO_BASE_PTR -urUSMGetMemAllocInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_USM_ALLOC_INFO_SIZE -urUSMGetMemAllocInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_USM_ALLOC_INFO_DEVICE urUSMGetMemAllocInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UR_USM_ALLOC_INFO_POOL -urUSMGetMemAllocInfoNegativeTest.InvalidNullHandleContext/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMGetMemAllocInfoNegativeTest.InvalidNullPointerMem/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMGetMemAllocInfoNegativeTest.InvalidEnumeration/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMGetMemAllocInfoNegativeTest.InvalidValuePropSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}__ -urUSMHostAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMHostAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMHostAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMHostAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled urUSMHostAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled urUSMHostAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_2048 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_8 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_512 -urUSMHostAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_2048 -urUSMSharedAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMSharedAllocTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMSharedAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMSharedAllocTest.SuccessWithDescriptors/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMSharedAllocTest.SuccessWithMultipleAdvices/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled -urUSMSharedAllocTest.SuccessWithMultipleAdvices/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled urUSMSharedAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled urUSMSharedAllocTest.InvalidUSMSize/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_4_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_8_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_16_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_32_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolEnabled_64_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_4_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_8_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_16_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_32_2048 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_8 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_512 -urUSMSharedAllocAlignmentTest.SuccessAlignedAllocations/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}____UsePoolDisabled_64_2048 diff --git a/test/loader/adapter_registry/CMakeLists.txt b/test/loader/adapter_registry/CMakeLists.txt index 2778ad5c40..6d80430e6c 100644 --- a/test/loader/adapter_registry/CMakeLists.txt +++ b/test/loader/adapter_registry/CMakeLists.txt @@ -51,3 +51,7 @@ add_adapter_reg_search_test(search-order SEARCH_PATH ${TEST_SEARCH_PATH} ENVS "TEST_ADAPTER_SEARCH_PATH=\"${TEST_SEARCH_PATH}\"" "TEST_CUR_SEARCH_PATH=\"${TEST_BIN_PATH}\"" SOURCES search_order.cpp) + +add_adapter_reg_search_test(prefilter + SEARCH_PATH "" + SOURCES prefilter.cpp) diff --git a/test/loader/adapter_registry/fixtures.hpp b/test/loader/adapter_registry/fixtures.hpp index 79a831d40f..da5c963e8a 100644 --- a/test/loader/adapter_registry/fixtures.hpp +++ b/test/loader/adapter_registry/fixtures.hpp @@ -74,5 +74,49 @@ struct adapterRegSearchTest : ::testing::Test { } } }; +#ifndef _WIN32 +struct adapterPreFilterTest : ::testing::Test { + ur_loader::AdapterRegistry *registry; + const fs::path levelzeroLibName = + MAKE_LIBRARY_NAME("ur_adapter_level_zero", "0"); + std::function islevelzeroLibName = + [this](const fs::path &path) { return path == levelzeroLibName; }; + + std::function &)> haslevelzeroLibName = + [this](const std::vector &paths) { + return std::any_of(paths.cbegin(), paths.cend(), + islevelzeroLibName); + }; + + const fs::path openclLibName = MAKE_LIBRARY_NAME("ur_adapter_opencl", "0"); + std::function isOpenclLibName = + [this](const fs::path &path) { return path == openclLibName; }; + + std::function &)> hasOpenclLibName = + [this](const std::vector &paths) { + return std::any_of(paths.cbegin(), paths.cend(), isOpenclLibName); + }; + + const fs::path cudaLibName = MAKE_LIBRARY_NAME("ur_adapter_cuda", "0"); + std::function isCudaLibName = + [this](const fs::path &path) { return path == cudaLibName; }; + + std::function &)> hasCudaLibName = + [this](const std::vector &paths) { + return std::any_of(paths.cbegin(), paths.cend(), isCudaLibName); + }; + + void SetUp(std::string filter) { + try { + setenv("ONEAPI_DEVICE_SELECTOR", filter.c_str(), 1); + registry = new ur_loader::AdapterRegistry; + } catch (const std::invalid_argument &e) { + FAIL() << e.what(); + } + } + void SetUp() override {} + void TearDown() override { delete registry; } +}; +#endif #endif // UR_ADAPTER_REG_TEST_HELPERS_H diff --git a/test/loader/adapter_registry/prefilter.cpp b/test/loader/adapter_registry/prefilter.cpp new file mode 100644 index 0000000000..1d2b095da3 --- /dev/null +++ b/test/loader/adapter_registry/prefilter.cpp @@ -0,0 +1,140 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "fixtures.hpp" + +#ifndef _WIN32 + +TEST_F(adapterPreFilterTest, testPrefilterAcceptFilterSingleBackend) { + SetUp("level_zero:*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_FALSE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterAcceptFilterMultipleBackends) { + SetUp("level_zero:*;opencl:*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterDiscardFilterSingleBackend) { + SetUp("!level_zero:*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_FALSE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_TRUE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterDiscardFilterMultipleBackends) { + SetUp("!level_zero:*;!cuda:*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_FALSE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterAcceptAndDiscardFilter) { + SetUp("!cuda:*;level_zero:*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_FALSE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterDiscardFilterAll) { + SetUp("*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_TRUE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterWithInvalidMissingBackend) { + SetUp(":garbage"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_TRUE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterWithInvalidBackend) { + SetUp("garbage:0"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_TRUE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_TRUE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterWithNotAllAndAcceptFilter) { + SetUp("!*;level_zero"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_TRUE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_FALSE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +TEST_F(adapterPreFilterTest, testPrefilterWithNotAllFilter) { + SetUp("!*"); + auto levelZeroExists = + std::any_of(registry->cbegin(), registry->cend(), haslevelzeroLibName); + EXPECT_FALSE(levelZeroExists); + auto openclExists = + std::any_of(registry->cbegin(), registry->cend(), hasOpenclLibName); + EXPECT_FALSE(openclExists); + auto cudaExists = + std::any_of(registry->cbegin(), registry->cend(), hasCudaLibName); + EXPECT_FALSE(cudaExists); +} + +#endif diff --git a/test/unit/utils/CMakeLists.txt b/test/unit/utils/CMakeLists.txt index a0e0fd3ef7..62681b1032 100644 --- a/test/unit/utils/CMakeLists.txt +++ b/test/unit/utils/CMakeLists.txt @@ -13,3 +13,6 @@ add_unit_test(params add_unit_test(print print.cpp) + +add_unit_test(helpers + helpers.cpp) diff --git a/test/unit/utils/helpers.cpp b/test/unit/utils/helpers.cpp new file mode 100644 index 0000000000..87223b21cc --- /dev/null +++ b/test/unit/utils/helpers.cpp @@ -0,0 +1,30 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include +#include + +#include "ur_util.hpp" + +TEST(groupDigits, Success) { + EXPECT_EQ(groupDigits(-1), "-1"); + EXPECT_EQ(groupDigits(-12), "-12"); + EXPECT_EQ(groupDigits(-123), "-123"); + EXPECT_EQ(groupDigits(-1234), "-1'234"); + EXPECT_EQ(groupDigits(-12345), "-12'345"); + EXPECT_EQ(groupDigits(-123456), "-123'456"); + EXPECT_EQ(groupDigits(-1234567), "-1'234'567"); + EXPECT_EQ(groupDigits(-12345678), "-12'345'678"); + + EXPECT_EQ(groupDigits(0), "0"); + EXPECT_EQ(groupDigits(1), "1"); + EXPECT_EQ(groupDigits(12), "12"); + EXPECT_EQ(groupDigits(123), "123"); + EXPECT_EQ(groupDigits(1234), "1'234"); + EXPECT_EQ(groupDigits(12345), "12'345"); + EXPECT_EQ(groupDigits(123456), "123'456"); + EXPECT_EQ(groupDigits(1234567), "1'234'567"); + EXPECT_EQ(groupDigits(12345678), "12'345'678"); +} diff --git a/test/unit/utils/params.cpp b/test/unit/utils/params.cpp index c456f69795..e86181344c 100644 --- a/test/unit/utils/params.cpp +++ b/test/unit/utils/params.cpp @@ -27,3 +27,17 @@ TEST(PrintPtr, nested_void_ptrs) { ur::details::printPtr(out, pppreal); EXPECT_THAT(out.str(), MatchesRegex(".+ \\(.+ \\(.+ \\(.+\\)\\)\\)")); } + +TEST(PrintBool, False) { + ur_bool_t value = false; + std::ostringstream out; + out << value; + EXPECT_STREQ(out.str().data(), "false"); +} + +TEST(PrintBool, True) { + ur_bool_t value = 1; + std::ostringstream out; + out << value; + EXPECT_STREQ(out.str().data(), "true"); +}