From a3930168d6383c797a1f9ce727c7ef409ab36aa3 Mon Sep 17 00:00:00 2001 From: Kamil Diedrich Date: Fri, 25 Jun 2021 01:25:22 +0000 Subject: [PATCH] Extend image functionality - add imageView extension - add import win32 NT handle - add black box test with imageView usage example Signed-off-by: Kamil Diedrich --- level_zero/api/core/ze_core_loader.cpp | 1 + .../api/extensions/public/ze_exp_ext.cpp | 10 + level_zero/core/source/CMakeLists.txt | 1 + .../core/source/helpers/properties_parser.h | 113 +++++++ level_zero/core/source/image/image.h | 6 +- level_zero/core/source/image/image_hw.h | 1 - level_zero/core/source/image/image_hw.inl | 58 ++-- level_zero/core/source/image/image_imp.cpp | 25 +- level_zero/core/source/image/image_imp.h | 38 +-- .../core/test/black_box_tests/CMakeLists.txt | 2 + .../test/black_box_tests/zello_copy_image.cpp | 4 +- .../test/black_box_tests/zello_image_view.cpp | 301 +++++++++++++++++ .../sources/cmdlist/test_cmdlist_2.cpp | 15 + .../sources/cmdlist/test_cmdlist_blit.cpp | 1 + .../sources/context/test_context.cpp | 1 + .../unit_tests/sources/helper/CMakeLists.txt | 1 + .../helper/properties_parser_tests.cpp | 159 +++++++++ .../unit_tests/sources/image/test_image.cpp | 313 +++++++++++++----- .../unit_tests/sources/kernel/test_kernel.cpp | 4 + 19 files changed, 915 insertions(+), 139 deletions(-) create mode 100644 level_zero/core/source/helpers/properties_parser.h create mode 100644 level_zero/core/test/black_box_tests/zello_image_view.cpp create mode 100644 level_zero/core/test/unit_tests/sources/helper/properties_parser_tests.cpp diff --git a/level_zero/api/core/ze_core_loader.cpp b/level_zero/api/core/ze_core_loader.cpp index 8f4d6a1967..01fc3b04ac 100644 --- a/level_zero/api/core/ze_core_loader.cpp +++ b/level_zero/api/core/ze_core_loader.cpp @@ -599,6 +599,7 @@ zeGetImageExpProcAddrTable( ze_result_t result = ZE_RESULT_SUCCESS; pDdiTable->pfnGetMemoryPropertiesExp = zeImageGetMemoryPropertiesExp; + pDdiTable->pfnViewCreateExp = zeImageViewCreateExp; driver_ddiTable.core_ddiTable.ImageExp = *pDdiTable; return result; } diff --git a/level_zero/api/extensions/public/ze_exp_ext.cpp b/level_zero/api/extensions/public/ze_exp_ext.cpp index 9efa0599ea..e2e1a5965c 100644 --- a/level_zero/api/extensions/public/ze_exp_ext.cpp +++ b/level_zero/api/extensions/public/ze_exp_ext.cpp @@ -30,6 +30,16 @@ zeImageGetMemoryPropertiesExp( return L0::Image::fromHandle(hImage)->getMemoryProperties(pMemoryProperties); } +ZE_APIEXPORT ze_result_t ZE_APICALL +zeImageViewCreateExp( + ze_context_handle_t hContext, + ze_device_handle_t hDevice, + const ze_image_desc_t *desc, + ze_image_handle_t hImage, + ze_image_handle_t *phImageView) { + return L0::Image::fromHandle(hImage)->createView(L0::Device::fromHandle(hDevice), desc, phImageView); +} + #if defined(__cplusplus) } // extern "C" #endif diff --git a/level_zero/core/source/CMakeLists.txt b/level_zero/core/source/CMakeLists.txt index cff820b6f5..d1aca29d70 100644 --- a/level_zero/core/source/CMakeLists.txt +++ b/level_zero/core/source/CMakeLists.txt @@ -52,6 +52,7 @@ set(L0_RUNTIME_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/helpers/api_specific_config_l0.cpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/implicit_scaling_l0.cpp ${CMAKE_CURRENT_SOURCE_DIR}/helpers/l0_populate_factory.h + ${CMAKE_CURRENT_SOURCE_DIR}/helpers/properties_parser.h ${CMAKE_CURRENT_SOURCE_DIR}/hw_helpers${BRANCH_DIR_SUFFIX}/hw_helpers.h ${CMAKE_CURRENT_SOURCE_DIR}/hw_helpers/l0_hw_helper_base.inl ${CMAKE_CURRENT_SOURCE_DIR}/hw_helpers/l0_hw_helper_skl_plus.inl diff --git a/level_zero/core/source/helpers/properties_parser.h b/level_zero/core/source/helpers/properties_parser.h new file mode 100644 index 0000000000..a0a4e960d8 --- /dev/null +++ b/level_zero/core/source/helpers/properties_parser.h @@ -0,0 +1,113 @@ +/* + * Copyright (C) 2021 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#pragma once + +#include "shared/source/helpers/debug_helpers.h" +#include "shared/source/helpers/surface_format_info.h" + +#include + +#include + +namespace L0 { +inline NEO::ImageType convertType(const ze_image_type_t type) { + switch (type) { + case ZE_IMAGE_TYPE_2D: + return NEO::ImageType::Image2D; + case ZE_IMAGE_TYPE_3D: + return NEO::ImageType::Image3D; + case ZE_IMAGE_TYPE_2DARRAY: + return NEO::ImageType::Image2DArray; + case ZE_IMAGE_TYPE_1D: + return NEO::ImageType::Image1D; + case ZE_IMAGE_TYPE_1DARRAY: + return NEO::ImageType::Image1DArray; + case ZE_IMAGE_TYPE_BUFFER: + return NEO::ImageType::Image1DBuffer; + default: + break; + } + return NEO::ImageType::Invalid; +} + +inline NEO::ImageDescriptor convertDescriptor(const ze_image_desc_t &imageDesc) { + NEO::ImageDescriptor desc = {}; + desc.fromParent = false; + desc.imageArraySize = imageDesc.arraylevels; + desc.imageDepth = imageDesc.depth; + desc.imageHeight = imageDesc.height; + desc.imageRowPitch = 0u; + desc.imageSlicePitch = 0u; + desc.imageType = convertType(imageDesc.type); + desc.imageWidth = imageDesc.width; + desc.numMipLevels = imageDesc.miplevels; + desc.numSamples = 0u; + return desc; +} + +struct StructuresLookupTable { + bool isSharedHandle; + struct SharedHandleType { + bool isSupportedHandle; + bool isDMABUFHandle; + int fd; + bool isNTHandle; + void *ntHnadle; + } sharedHandleType; + bool areImageProperties; + struct ImageProperties { + bool isPlanarExtension; + uint32_t planeIndex; + NEO::ImageDescriptor imageDescriptor; + } imageProperties; +}; + +inline ze_result_t prepareL0StructuresLookupTable(StructuresLookupTable &lookupTable, const void *desc) { + const ze_base_desc_t *extendedDesc = reinterpret_cast(desc); + while (extendedDesc) { + if (extendedDesc->stype == ZE_STRUCTURE_TYPE_IMAGE_DESC) { + const ze_image_desc_t *imageDesc = reinterpret_cast(extendedDesc); + lookupTable.areImageProperties = true; + lookupTable.imageProperties.imageDescriptor = convertDescriptor(*imageDesc); + } else if (extendedDesc->stype == ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD) { + lookupTable.isSharedHandle = true; + const ze_external_memory_import_fd_t *linuxExternalMemoryImportDesc = reinterpret_cast(extendedDesc); + if (linuxExternalMemoryImportDesc->flags == ZE_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF) { + lookupTable.sharedHandleType.isSupportedHandle = true; + lookupTable.sharedHandleType.isDMABUFHandle = true; + lookupTable.sharedHandleType.fd = linuxExternalMemoryImportDesc->fd; + } else { + lookupTable.sharedHandleType.isSupportedHandle = false; + return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; + } + } else if (extendedDesc->stype == ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_WIN32) { + lookupTable.isSharedHandle = true; + const ze_external_memory_import_win32_handle_t *windowsExternalMemoryImportDesc = reinterpret_cast(extendedDesc); + if (windowsExternalMemoryImportDesc->flags == ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32) { + lookupTable.sharedHandleType.isSupportedHandle = true; + lookupTable.sharedHandleType.isNTHandle = true; + lookupTable.sharedHandleType.ntHnadle = windowsExternalMemoryImportDesc->handle; + } else { + lookupTable.sharedHandleType.isSupportedHandle = false; + return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; + } + } else if (extendedDesc->stype == ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC) { + const ze_image_view_planar_exp_desc_t *imageViewDesc = reinterpret_cast(extendedDesc); + lookupTable.areImageProperties = true; + lookupTable.imageProperties.isPlanarExtension = true; + lookupTable.imageProperties.planeIndex = imageViewDesc->planeIndex; + } else { + return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; + } + + extendedDesc = reinterpret_cast(extendedDesc->pNext); + } + + return ZE_RESULT_SUCCESS; +} +} // namespace L0 diff --git a/level_zero/core/source/image/image.h b/level_zero/core/source/image/image.h index 0fa5ef8e9a..90ff070a5c 100644 --- a/level_zero/core/source/image/image.h +++ b/level_zero/core/source/image/image.h @@ -16,10 +16,10 @@ struct _ze_image_handle_t {}; namespace NEO { struct ImageInfo; -} +struct ImageDescriptor; +} // namespace NEO namespace L0 { - struct Image : _ze_image_handle_t { template struct Allocator { @@ -31,6 +31,8 @@ struct Image : _ze_image_handle_t { static ze_result_t create(uint32_t productFamily, Device *device, const ze_image_desc_t *desc, Image **pImage); + virtual ze_result_t createView(Device *device, const ze_image_desc_t *desc, ze_image_handle_t *pImage) = 0; + virtual NEO::GraphicsAllocation *getAllocation() = 0; virtual void copySurfaceStateToSSH(void *surfaceStateHeap, const uint32_t surfaceStateOffset, diff --git a/level_zero/core/source/image/image_hw.h b/level_zero/core/source/image/image_hw.h index 6c7f23c4af..edefe2976e 100644 --- a/level_zero/core/source/image/image_hw.h +++ b/level_zero/core/source/image/image_hw.h @@ -15,7 +15,6 @@ #include "level_zero/core/source/image/image_imp.h" namespace L0 { - template struct ImageCoreFamily : public ImageImp { using GfxFamily = typename NEO::GfxFamilyMapper::GfxFamily; diff --git a/level_zero/core/source/image/image_hw.inl b/level_zero/core/source/image/image_hw.inl index 30f282b484..11e6aa8806 100644 --- a/level_zero/core/source/image/image_hw.inl +++ b/level_zero/core/source/image/image_hw.inl @@ -16,18 +16,30 @@ #include "shared/source/memory_manager/memory_manager.h" #include "shared/source/utilities/compiler_support.h" +#include "level_zero/core/source/helpers/properties_parser.h" #include "level_zero/core/source/image/image_formats.h" #include "level_zero/core/source/image/image_hw.h" namespace L0 { + template ze_result_t ImageCoreFamily::initialize(Device *device, const ze_image_desc_t *desc) { using RENDER_SURFACE_STATE = typename GfxFamily::RENDER_SURFACE_STATE; + StructuresLookupTable lookupTable = {}; + + lookupTable.areImageProperties = true; + lookupTable.imageProperties.imageDescriptor = convertDescriptor(*desc); + + auto parseResult = prepareL0StructuresLookupTable(lookupTable, desc->pNext); + + if (parseResult != ZE_RESULT_SUCCESS) { + return parseResult; + } + bool isMediaFormatLayout = isMediaFormat(desc->format.layout); - auto imageDescriptor = convertDescriptor(*desc); - imgInfo.imgDesc = imageDescriptor; + imgInfo.imgDesc = lookupTable.imageProperties.imageDescriptor; imgInfo.surfaceFormat = &ImageFormats::formats[desc->format.layout][desc->format.type]; imageFormatDesc = *const_cast(desc); @@ -57,34 +69,36 @@ ze_result_t ImageCoreFamily::initialize(Device *device, const ze_ } imgInfo.linearStorage = surfaceType == RENDER_SURFACE_STATE::SURFACE_TYPE_SURFTYPE_1D; - imgInfo.plane = GMM_NO_PLANE; + imgInfo.plane = lookupTable.imageProperties.isPlanarExtension ? static_cast(lookupTable.imageProperties.planeIndex + 1u) : GMM_NO_PLANE; imgInfo.useLocalMemory = false; imgInfo.preferRenderCompression = false; - if (desc->pNext) { - const ze_base_desc_t *extendedDesc = reinterpret_cast(desc->pNext); - if (extendedDesc->stype == ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_DESC) { - return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; - } else if (extendedDesc->stype == ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD) { - const ze_external_memory_import_fd_t *externalMemoryImportDesc = - reinterpret_cast(extendedDesc); - if (externalMemoryImportDesc->flags & ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD) { + if (!isImageView) { + if (lookupTable.isSharedHandle) { + if (!lookupTable.sharedHandleType.isSupportedHandle) { return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } - - NEO::AllocationProperties properties(device->getRootDeviceIndex(), true, imgInfo, NEO::GraphicsAllocation::AllocationType::SHARED_IMAGE, device->getNEODevice()->getDeviceBitfield()); - - allocation = device->getNEODevice()->getMemoryManager()->createGraphicsAllocationFromSharedHandle(externalMemoryImportDesc->fd, properties, false, false); - device->getNEODevice()->getMemoryManager()->closeSharedHandle(allocation); + if (lookupTable.sharedHandleType.isDMABUFHandle) { + NEO::AllocationProperties properties(device->getRootDeviceIndex(), true, imgInfo, NEO::GraphicsAllocation::AllocationType::SHARED_IMAGE, device->getNEODevice()->getDeviceBitfield()); + allocation = device->getNEODevice()->getMemoryManager()->createGraphicsAllocationFromSharedHandle(lookupTable.sharedHandleType.fd, properties, false, false); + device->getNEODevice()->getMemoryManager()->closeSharedHandle(allocation); + } else if (lookupTable.sharedHandleType.isNTHandle) { + auto verifyResult = device->getNEODevice()->getMemoryManager()->verifyHandle(NEO::toOsHandle(lookupTable.sharedHandleType.ntHnadle), device->getNEODevice()->getRootDeviceIndex(), true); + if (!verifyResult) { + return ZE_RESULT_ERROR_INVALID_ARGUMENT; + } + allocation = device->getNEODevice()->getMemoryManager()->createGraphicsAllocationFromNTHandle(lookupTable.sharedHandleType.ntHnadle, device->getNEODevice()->getRootDeviceIndex()); + } } else { - return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION; - } - } else { - NEO::AllocationProperties properties(device->getRootDeviceIndex(), true, imgInfo, NEO::GraphicsAllocation::AllocationType::IMAGE, device->getNEODevice()->getDeviceBitfield()); + NEO::AllocationProperties properties(device->getRootDeviceIndex(), true, imgInfo, NEO::GraphicsAllocation::AllocationType::IMAGE, device->getNEODevice()->getDeviceBitfield()); - allocation = device->getNEODevice()->getMemoryManager()->allocateGraphicsMemoryWithProperties(properties); + allocation = device->getNEODevice()->getMemoryManager()->allocateGraphicsMemoryWithProperties(properties); + } + if (allocation == nullptr) { + return ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY; + } } - UNRECOVERABLE_IF(allocation == nullptr); + auto gmm = this->allocation->getDefaultGmm(); auto gmmHelper = static_cast(device->getNEODevice()->getRootDeviceEnvironment()).getGmmHelper(); diff --git a/level_zero/core/source/image/image_imp.cpp b/level_zero/core/source/image/image_imp.cpp index db082740cf..ca8988f4e3 100644 --- a/level_zero/core/source/image/image_imp.cpp +++ b/level_zero/core/source/image/image_imp.cpp @@ -16,7 +16,7 @@ namespace L0 { ImageAllocatorFn imageFactory[IGFX_MAX_PRODUCT] = {}; ImageImp::~ImageImp() { - if (this->device != nullptr) { + if (!isImageView && this->device != nullptr) { this->device->getNEODevice()->getMemoryManager()->freeGraphicsMemory(this->allocation); } } @@ -26,6 +26,29 @@ ze_result_t ImageImp::destroy() { return ZE_RESULT_SUCCESS; } +ze_result_t ImageImp::createView(Device *device, const ze_image_desc_t *desc, ze_image_handle_t *pImage) { + auto productFamily = device->getNEODevice()->getHardwareInfo().platform.eProductFamily; + + ImageAllocatorFn allocator = nullptr; + allocator = imageFactory[productFamily]; + + ImageImp *image = nullptr; + + image = static_cast((*allocator)()); + image->isImageView = true; + image->allocation = allocation; + auto result = image->initialize(device, desc); + + if (result != ZE_RESULT_SUCCESS) { + image->destroy(); + image = nullptr; + } + + *pImage = image; + + return result; +} + ze_result_t Image::create(uint32_t productFamily, Device *device, const ze_image_desc_t *desc, Image **pImage) { ze_result_t result = ZE_RESULT_SUCCESS; ImageAllocatorFn allocator = nullptr; diff --git a/level_zero/core/source/image/image_imp.h b/level_zero/core/source/image/image_imp.h index b22d39593a..48df3153fe 100644 --- a/level_zero/core/source/image/image_imp.h +++ b/level_zero/core/source/image/image_imp.h @@ -26,6 +26,8 @@ struct ImageImp : public Image { return imageFormatDesc; } + ze_result_t createView(Device *device, const ze_image_desc_t *desc, ze_image_handle_t *pImage) override; + ze_result_t getMemoryProperties(ze_image_memory_properties_exp_t *pMemoryProperties) override { pMemoryProperties->rowPitch = imgInfo.rowPitch; pMemoryProperties->slicePitch = imgInfo.slicePitch; @@ -34,42 +36,8 @@ struct ImageImp : public Image { return ZE_RESULT_SUCCESS; } - static NEO::ImageType convertType(const ze_image_type_t type) { - switch (type) { - case ZE_IMAGE_TYPE_2D: - return NEO::ImageType::Image2D; - case ZE_IMAGE_TYPE_3D: - return NEO::ImageType::Image3D; - case ZE_IMAGE_TYPE_2DARRAY: - return NEO::ImageType::Image2DArray; - case ZE_IMAGE_TYPE_1D: - return NEO::ImageType::Image1D; - case ZE_IMAGE_TYPE_1DARRAY: - return NEO::ImageType::Image1DArray; - case ZE_IMAGE_TYPE_BUFFER: - return NEO::ImageType::Image1DBuffer; - default: - break; - } - return NEO::ImageType::Invalid; - } - - static NEO::ImageDescriptor convertDescriptor(const ze_image_desc_t &imageDesc) { - NEO::ImageDescriptor desc = {}; - desc.fromParent = false; - desc.imageArraySize = imageDesc.arraylevels; - desc.imageDepth = imageDesc.depth; - desc.imageHeight = imageDesc.height; - desc.imageRowPitch = 0u; - desc.imageSlicePitch = 0u; - desc.imageType = convertType(imageDesc.type); - desc.imageWidth = imageDesc.width; - desc.numMipLevels = imageDesc.miplevels; - desc.numSamples = 0u; - return desc; - } - protected: + bool isImageView = false; Device *device = nullptr; NEO::ImageInfo imgInfo = {}; NEO::GraphicsAllocation *allocation = nullptr; diff --git a/level_zero/core/test/black_box_tests/CMakeLists.txt b/level_zero/core/test/black_box_tests/CMakeLists.txt index 48665e3318..c4c18fb7d7 100644 --- a/level_zero/core/test/black_box_tests/CMakeLists.txt +++ b/level_zero/core/test/black_box_tests/CMakeLists.txt @@ -22,6 +22,7 @@ if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug") zello_fence zello_printf zello_image + zello_image_view ) include_directories(common) @@ -58,6 +59,7 @@ if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug") target_link_libraries(zello_scratch PUBLIC ocloc_lib) target_link_libraries(zello_fence PUBLIC ocloc_lib) target_link_libraries(zello_printf PUBLIC ocloc_lib) + target_link_libraries(zello_image_view PUBLIC ocloc_lib) if(UNIX) target_link_libraries(zello_world_global_work_offset PUBLIC ocloc_lib) endif() diff --git a/level_zero/core/test/black_box_tests/zello_copy_image.cpp b/level_zero/core/test/black_box_tests/zello_copy_image.cpp index 4bb2024dbf..3e2f4e3754 100644 --- a/level_zero/core/test/black_box_tests/zello_copy_image.cpp +++ b/level_zero/core/test/black_box_tests/zello_copy_image.cpp @@ -39,7 +39,7 @@ void testAppendImageCopy(ze_context_handle_t &context, ze_device_handle_t &devic cmdListDesc.flags = 0; SUCCESS_OR_TERMINATE(zeCommandListCreate(context, device, &cmdListDesc, &cmdList)); - ze_image_desc_t srcImgDesc = {ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES, + ze_image_desc_t srcImgDesc = {ZE_STRUCTURE_TYPE_IMAGE_DESC, nullptr, 0, ZE_IMAGE_TYPE_2D, @@ -58,7 +58,7 @@ void testAppendImageCopy(ze_context_handle_t &context, ze_device_handle_t &devic SUCCESS_OR_TERMINATE( zeImageCreate(context, device, const_cast(&srcImgDesc), &srcImg)); - ze_image_desc_t dstImgDesc = {ZE_STRUCTURE_TYPE_DEVICE_IMAGE_PROPERTIES, + ze_image_desc_t dstImgDesc = {ZE_STRUCTURE_TYPE_IMAGE_DESC, nullptr, ZE_IMAGE_FLAG_KERNEL_WRITE, ZE_IMAGE_TYPE_2D, diff --git a/level_zero/core/test/black_box_tests/zello_image_view.cpp b/level_zero/core/test/black_box_tests/zello_image_view.cpp new file mode 100644 index 0000000000..dc46e372aa --- /dev/null +++ b/level_zero/core/test/black_box_tests/zello_image_view.cpp @@ -0,0 +1,301 @@ +/* + * Copyright (C) 2021 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include "zello_common.h" +#include "zello_compile.h" + +extern bool verbose; +bool verbose = false; + +const char *readNV12Module = R"===( + __kernel void + ReadNV12Kernel( + read_only image2d_t nv12Img, + uint width, + uint height, + __global uchar *pDest) { + int tid_x = get_global_id(0); + int tid_y = get_global_id(1); + float4 colorY; + int2 coord; + const sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_NONE | + CLK_FILTER_NEAREST; + if (tid_x < width && tid_y < height) { + coord = (int2)(tid_x, tid_y); + if (((tid_y * width) + tid_x) < (width * height)) { + colorY = read_imagef(nv12Img, samplerA, coord); + pDest[(tid_y * width) + tid_x] = (uchar)(255.0f * colorY.y); + if ((tid_x % 2 == 0) && (tid_y % 2 == 0)) { + pDest[(width * height) + (tid_y / 2 * width) + (tid_x)] = (uchar)(255.0f * colorY.z); + pDest[(width * height) + (tid_y / 2 * width) + (tid_x) + 1] = (uchar)(255.0f * colorY.x); + } + } + } +} +)==="; + +void testAppendImageViewCopy(ze_context_handle_t &context, ze_device_handle_t &device, bool &validRet) { + std::string buildLog; + auto spirV = compileToSpirV(readNV12Module, "", buildLog); + if (buildLog.size() > 0) { + std::cout << "Build log " << buildLog; + } + SUCCESS_OR_TERMINATE((0 == spirV.size())); + + const size_t width = 32; + const size_t height = 32; + const size_t depth = 1; + const size_t size = 4 * width * height * depth; /* 4 channels per pixel */ + + ze_command_queue_handle_t cmdQueue = nullptr; + ze_command_list_handle_t cmdList = nullptr; + + ze_command_queue_desc_t cmdQueueDesc = {}; + cmdQueueDesc.pNext = nullptr; + cmdQueueDesc.flags = 0; + cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS; + cmdQueueDesc.priority = ZE_COMMAND_QUEUE_PRIORITY_NORMAL; + cmdQueueDesc.ordinal = 0; + cmdQueueDesc.index = 0; + SUCCESS_OR_TERMINATE(zeCommandQueueCreate(context, device, &cmdQueueDesc, &cmdQueue)); + + ze_command_list_desc_t cmdListDesc = {}; + cmdListDesc.stype = ZE_STRUCTURE_TYPE_COMMAND_LIST_DESC; + cmdListDesc.pNext = nullptr; + cmdListDesc.flags = 0; + SUCCESS_OR_TERMINATE(zeCommandListCreate(context, device, &cmdListDesc, &cmdList)); + + // create NV12 image + ze_image_desc_t srcImgDesc = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + nullptr, + (ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_NV12, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}, + width, + height, + depth, + 0, + 0}; + + ze_image_handle_t srcImg; + + SUCCESS_OR_TERMINATE( + zeImageCreate(context, device, const_cast(&srcImgDesc), &srcImg)); + + // create image_veiw for Y plane + ze_image_view_planar_exp_desc_t planeYdesc = {}; + planeYdesc.stype = ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC; + planeYdesc.planeIndex = 0u; // Y plane + + ze_image_desc_t imageViewDescPlaneY = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + &planeYdesc, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_8, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_A, ZE_IMAGE_FORMAT_SWIZZLE_B, + ZE_IMAGE_FORMAT_SWIZZLE_G, ZE_IMAGE_FORMAT_SWIZZLE_R}, + width, + height, + depth, + 0, + 0}; + ze_image_handle_t planeYImageView; + + SUCCESS_OR_TERMINATE( + zeImageViewCreateExp(context, device, &imageViewDescPlaneY, srcImg, &planeYImageView)); + + // create image_view for UV plane + ze_image_view_planar_exp_desc_t planeUVdesc = {}; + planeUVdesc.stype = ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC; + planeUVdesc.planeIndex = 1u; // UV plane + + ze_image_desc_t imageViewDescPlaneUV = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + &planeUVdesc, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_8_8, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_A, ZE_IMAGE_FORMAT_SWIZZLE_B, + ZE_IMAGE_FORMAT_SWIZZLE_G, ZE_IMAGE_FORMAT_SWIZZLE_R}, + width / 2, + height / 2, + depth, + 0, + 0}; + ze_image_handle_t planeUVImageView; + + SUCCESS_OR_TERMINATE( + zeImageViewCreateExp(context, device, &imageViewDescPlaneUV, srcImg, &planeUVImageView)); + + // prepare input data + std::vector srcVecY; + srcVecY.resize(width * height); + + for (size_t i = 0; i < width * height; ++i) { + srcVecY[i] = static_cast(i); + } + + std::vector srcVecUV; + srcVecUV.resize((width / 2) * (height)); + + for (size_t i = 0; i < (width / 2) * (height); ++i) { + if (i % 2 == 0) { + srcVecUV[i] = static_cast(0x33); + } else { + srcVecUV[i] = static_cast(0x55); + } + } + + // prepare destination buffer + uint8_t *dstMem; + + ze_host_mem_alloc_desc_t hostDesc = {}; + hostDesc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC; + hostDesc.pNext = nullptr; + hostDesc.flags = 0; + SUCCESS_OR_TERMINATE(zeMemAllocHost(context, &hostDesc, size, 1, (void **)(&dstMem))); + + for (size_t i = 0; i < size; ++i) { + dstMem[i] = 0xff; + } + + ze_image_region_t srcYRegion = {0, 0, 0, width, height, depth}; + ze_image_region_t srcUVRegion = {0, 0, 0, width / 2, height / 2, depth}; + + SUCCESS_OR_TERMINATE(zeCommandListAppendImageCopyFromMemory(cmdList, planeYImageView, srcVecY.data(), + &srcYRegion, nullptr, 0, nullptr)); + + SUCCESS_OR_TERMINATE(zeCommandListAppendImageCopyFromMemory(cmdList, planeUVImageView, srcVecUV.data(), + &srcUVRegion, nullptr, 0, nullptr)); + + SUCCESS_OR_TERMINATE(zeCommandListAppendBarrier(cmdList, nullptr, 0, nullptr)); + + { + // create kernel which reads NV12 surface + ze_module_handle_t module = nullptr; + ze_kernel_handle_t kernel = nullptr; + + ze_module_desc_t moduleDesc = {}; + ze_module_build_log_handle_t buildlog; + moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + moduleDesc.pInputModule = spirV.data(); + moduleDesc.inputSize = spirV.size(); + moduleDesc.pBuildFlags = ""; + + if (zeModuleCreate(context, device, &moduleDesc, &module, &buildlog) != ZE_RESULT_SUCCESS) { + size_t szLog = 0; + zeModuleBuildLogGetString(buildlog, &szLog, nullptr); + + char *strLog = (char *)malloc(szLog); + zeModuleBuildLogGetString(buildlog, &szLog, strLog); + std::cout << "Build log:" << strLog << std::endl; + + free(strLog); + } + SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog)); + + ze_kernel_desc_t kernelDesc = {}; + kernelDesc.pKernelName = "ReadNV12Kernel"; + SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel)); + + uint32_t groupSizeX = width; + uint32_t groupSizeY = height; + uint32_t groupSizeZ = 1u; + SUCCESS_OR_TERMINATE(zeKernelSuggestGroupSize(kernel, groupSizeX, 1U, 1U, &groupSizeX, &groupSizeY, &groupSizeZ)); + SUCCESS_OR_TERMINATE(zeKernelSetGroupSize(kernel, groupSizeX, groupSizeY, groupSizeZ)); + + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 0, sizeof(srcImg), &srcImg)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 1, sizeof(int), &width)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 2, sizeof(int), &height)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 3, sizeof(void *), &dstMem)); + + ze_group_count_t dispatchTraits; + dispatchTraits.groupCountX = width / groupSizeX; + dispatchTraits.groupCountY = height / groupSizeY; + dispatchTraits.groupCountZ = 1u; + + SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatchTraits, + nullptr, 0, nullptr)); + + SUCCESS_OR_TERMINATE(zeCommandListAppendBarrier(cmdList, nullptr, 0, nullptr)); + } + + // destination buffer for Y plane + std::vector dstVecY; + dstVecY.resize(width * height); + // destination buffer for UV plane + std::vector dstVecUV; + dstVecUV.resize((width / 2) * (height)); + + // read Y plane data + SUCCESS_OR_TERMINATE(zeCommandListAppendImageCopyToMemory(cmdList, dstVecY.data(), planeYImageView, + &srcYRegion, nullptr, 0, nullptr)); + + // read UV plane data + SUCCESS_OR_TERMINATE(zeCommandListAppendImageCopyToMemory(cmdList, dstVecUV.data(), planeUVImageView, + &srcUVRegion, nullptr, 0, nullptr)); + + SUCCESS_OR_TERMINATE(zeCommandListAppendBarrier(cmdList, nullptr, 0, nullptr)); + + SUCCESS_OR_TERMINATE(zeCommandListClose(cmdList)); + SUCCESS_OR_TERMINATE(zeCommandQueueExecuteCommandLists(cmdQueue, 1, &cmdList, nullptr)); + SUCCESS_OR_TERMINATE(zeCommandQueueSynchronize(cmdQueue, std::numeric_limits::max())); + + // validate Y plane data + auto result = memcmp(srcVecY.data(), dstVecY.data(), width * height); + + if (result != 0) { + std::cout << "Failed to validate data read for plane Y from Y-plane view" << std::endl; + } + + result = memcmp(dstVecY.data(), dstMem, width * height); + + if (result != 0) { + std::cout << "Failed to validate data read for plane Y from nv12 surface" << std::endl; + } + + // validate UV plane data + result = memcmp(srcVecUV.data(), dstVecUV.data(), (width / 2) * (height)); + + if (result != 0) { + std::cout << "Failed to validate data read for plane Y from Y-plane view" << std::endl; + } + + result = memcmp(dstVecUV.data(), (dstMem + (width * height)), (width / 2) * (height)); + + if (result != 0) { + std::cout << "Failed to validate data read for plane UV from nv12 surface" << std::endl; + } + + // cleanup + SUCCESS_OR_TERMINATE(zeImageDestroy(srcImg)); + SUCCESS_OR_TERMINATE(zeImageDestroy(planeYImageView)); + SUCCESS_OR_TERMINATE(zeImageDestroy(planeUVImageView)); + SUCCESS_OR_TERMINATE(zeCommandListDestroy(cmdList)); + SUCCESS_OR_TERMINATE(zeCommandQueueDestroy(cmdQueue)); +} + +int main(int argc, char *argv[]) { + ze_context_handle_t context = nullptr; + auto devices = zelloInitContextAndGetDevices(context); + auto device = devices[0]; + bool outputValidationSuccessful; + + ze_device_properties_t deviceProperties = {}; + SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties)); + std::cout << "Device : \n" + << " * name : " << deviceProperties.name << "\n" + << " * vendorId : " << std::hex << deviceProperties.vendorId << "\n"; + + testAppendImageViewCopy(context, device, outputValidationSuccessful); + + SUCCESS_OR_TERMINATE(zeContextDestroy(context)); + std::cout << "\nZello Copy Image Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n"; + return 0; +} \ No newline at end of file diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp index cf6c40b159..06488d3567 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp @@ -413,6 +413,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListWhenCopyFromMemoryToImageThenBl void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHW = std::make_unique>>(); imageHW->initialize(device, &zeDesc); @@ -427,6 +428,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhenIma void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHW = std::make_unique>>(); imageHW->initialize(device, &zeDesc); @@ -446,6 +448,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhenIma void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHW = std::make_unique>>(); imageHW->initialize(device, &zeDesc); @@ -465,6 +468,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen1DI void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_1D; zeDesc.height = 9; zeDesc.depth = 9; @@ -487,6 +491,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen1DI void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_1D; zeDesc.height = 9; zeDesc.depth = 9; @@ -509,6 +514,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen1DA void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_1DARRAY; zeDesc.height = 9; zeDesc.depth = 9; @@ -531,6 +537,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen1DA void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_1DARRAY; zeDesc.height = 9; zeDesc.depth = 9; @@ -553,6 +560,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen2DI void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_2D; zeDesc.height = 2; @@ -574,6 +582,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen2DI void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_2D; zeDesc.height = 2; auto imageHW = std::make_unique>>(); @@ -594,6 +603,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen2DI void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_2D; zeDesc.height = 2; zeDesc.depth = 9; @@ -616,6 +626,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen2DI void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_2D; zeDesc.height = 2; zeDesc.depth = 9; @@ -638,6 +649,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen3DI void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_3D; zeDesc.height = 2; zeDesc.depth = 2; @@ -657,6 +669,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen3DI void *srcPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.type = ZE_IMAGE_TYPE_3D; zeDesc.height = 2; zeDesc.depth = 2; @@ -676,6 +689,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListWhenCopyFromImageToMemoryThenBl void *dstPtr = reinterpret_cast(0x1234); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHW = std::make_unique>>(); imageHW->initialize(device, &zeDesc); @@ -688,6 +702,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListWhenCopyFromImageToImageThenBli MockCommandListHw cmdList; cmdList.initialize(device, NEO::EngineGroupType::Copy, 0u); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHWSrc = std::make_unique>>(); auto imageHWDst = std::make_unique>>(); imageHWSrc->initialize(device, &zeDesc); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_blit.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_blit.cpp index 75d0733a42..97a30b63c7 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_blit.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_blit.cpp @@ -278,6 +278,7 @@ HWTEST2_F(AppendMemoryCopy, givenCopyCommandListWhenCopyFromImagBlitThenCommandA ze_result_t returnValue; std::unique_ptr commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::Copy, 0u, returnValue)); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHWSrc = std::make_unique>>(); auto imageHWDst = std::make_unique>>(); imageHWSrc->initialize(device, &zeDesc); diff --git a/level_zero/core/test/unit_tests/sources/context/test_context.cpp b/level_zero/core/test/unit_tests/sources/context/test_context.cpp index 9e6a6235b2..187a41465d 100644 --- a/level_zero/core/test/unit_tests/sources/context/test_context.cpp +++ b/level_zero/core/test/unit_tests/sources/context/test_context.cpp @@ -639,6 +639,7 @@ HWTEST2_F(ContextTest, WhenCreatingImageThenSuccessIsReturned, IsAtMostProductDG ze_image_handle_t image = {}; ze_image_desc_t imageDesc = {}; + imageDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; res = contextImp->createImage(device, &imageDesc, &image); EXPECT_EQ(ZE_RESULT_SUCCESS, res); diff --git a/level_zero/core/test/unit_tests/sources/helper/CMakeLists.txt b/level_zero/core/test/unit_tests/sources/helper/CMakeLists.txt index 09c326b792..1c8888001a 100644 --- a/level_zero/core/test/unit_tests/sources/helper/CMakeLists.txt +++ b/level_zero/core/test/unit_tests/sources/helper/CMakeLists.txt @@ -8,5 +8,6 @@ target_sources(${TARGET_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt ${CMAKE_CURRENT_SOURCE_DIR}/api_specific_config_l0_tests.cpp ${CMAKE_CURRENT_SOURCE_DIR}/heap_assigner_l0_tests.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/properties_parser_tests.cpp ) add_subdirectories() diff --git a/level_zero/core/test/unit_tests/sources/helper/properties_parser_tests.cpp b/level_zero/core/test/unit_tests/sources/helper/properties_parser_tests.cpp new file mode 100644 index 0000000000..21fee06797 --- /dev/null +++ b/level_zero/core/test/unit_tests/sources/helper/properties_parser_tests.cpp @@ -0,0 +1,159 @@ +/* + * Copyright (C) 2021 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include "test.h" + +#include "level_zero/core/source/helpers/properties_parser.h" + +namespace L0 { +namespace ult { +struct ImageStaticFunctionConvertTypeTest : public testing::TestWithParam> { + void SetUp() override { + } + + void TearDown() override { + } +}; + +TEST_P(ImageStaticFunctionConvertTypeTest, givenZeImageFormatTypeWhenConvertTypeThenCorrectImageTypeReturned) { + auto params = GetParam(); + EXPECT_EQ(convertType(params.first), params.second); +} + +std::pair validTypes[] = { + {ZE_IMAGE_TYPE_2D, NEO::ImageType::Image2D}, + {ZE_IMAGE_TYPE_3D, NEO::ImageType::Image3D}, + {ZE_IMAGE_TYPE_2DARRAY, NEO::ImageType::Image2DArray}, + {ZE_IMAGE_TYPE_1D, NEO::ImageType::Image1D}, + {ZE_IMAGE_TYPE_1DARRAY, NEO::ImageType::Image1DArray}, + {ZE_IMAGE_TYPE_BUFFER, NEO::ImageType::Image1DBuffer}}; + +INSTANTIATE_TEST_CASE_P( + imageTypeFlags, + ImageStaticFunctionConvertTypeTest, + testing::ValuesIn(validTypes)); + +TEST(ImageStaticFunctionConvertInvalidType, givenInvalidZeImageFormatTypeWhenConvertTypeThenInvalidFormatIsRetrurned) { + EXPECT_EQ(convertType(ZE_IMAGE_TYPE_FORCE_UINT32), NEO::ImageType::Invalid); +} + +TEST(ConvertDescriptorTest, givenZeImageDescWhenConvertDescriptorThenCorrectImageDescriptorReturned) { + ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + zeDesc.arraylevels = 1u; + zeDesc.depth = 1u; + zeDesc.height = 1u; + zeDesc.width = 1u; + zeDesc.miplevels = 1u; + zeDesc.type = ZE_IMAGE_TYPE_2DARRAY; + + NEO::ImageDescriptor desc = convertDescriptor(zeDesc); + EXPECT_EQ(desc.fromParent, false); + EXPECT_EQ(desc.imageArraySize, zeDesc.arraylevels); + EXPECT_EQ(desc.imageDepth, zeDesc.depth); + EXPECT_EQ(desc.imageHeight, zeDesc.height); + EXPECT_EQ(desc.imageRowPitch, 0u); + EXPECT_EQ(desc.imageSlicePitch, 0u); + EXPECT_EQ(desc.imageType, NEO::ImageType::Image2DArray); + EXPECT_EQ(desc.imageWidth, zeDesc.width); + EXPECT_EQ(desc.numMipLevels, zeDesc.miplevels); + EXPECT_EQ(desc.numSamples, 0u); +} + +TEST(L0StructuresLookupTableTests, givenL0StructuresWithFDWhenPrepareLookupTableThenProperFieldsInLookupTableAreSet) { + ze_image_desc_t imageDesc = {}; + imageDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + imageDesc.height = 10; + imageDesc.width = 10; + imageDesc.depth = 10; + ze_external_memory_import_fd_t fdStructure = {}; + fdStructure.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; + fdStructure.fd = 1; + fdStructure.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF; + ze_image_view_planar_exp_desc_t imageView = {}; + imageView.stype = ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC; + imageView.planeIndex = 1u; + + imageDesc.pNext = &fdStructure; + fdStructure.pNext = &imageView; + imageView.pNext = nullptr; + + StructuresLookupTable l0LookupTable = {}; + auto result = prepareL0StructuresLookupTable(l0LookupTable, &imageDesc); + + EXPECT_EQ(result, ZE_RESULT_SUCCESS); + + EXPECT_TRUE(l0LookupTable.isSharedHandle); + + EXPECT_TRUE(l0LookupTable.sharedHandleType.isSupportedHandle); + EXPECT_TRUE(l0LookupTable.sharedHandleType.isDMABUFHandle); + EXPECT_EQ(l0LookupTable.sharedHandleType.fd, fdStructure.fd); + + EXPECT_TRUE(l0LookupTable.areImageProperties); + + EXPECT_EQ(l0LookupTable.imageProperties.planeIndex, imageView.planeIndex); + EXPECT_EQ(l0LookupTable.imageProperties.imageDescriptor.imageWidth, imageDesc.width); + EXPECT_EQ(l0LookupTable.imageProperties.imageDescriptor.imageHeight, imageDesc.height); + EXPECT_EQ(l0LookupTable.imageProperties.imageDescriptor.imageDepth, imageDesc.depth); +} + +TEST(L0StructuresLookupTableTests, givenL0StructuresWithNTHandleWhenPrepareLookupTableThenProperFieldsInLookupTableAreSet) { + uint64_t handle = 0x02; + ze_external_memory_import_win32_handle_t importNTHandle = {}; + importNTHandle.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_WIN32; + importNTHandle.handle = &handle; + importNTHandle.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32; + + StructuresLookupTable l0LookupTable = {}; + auto result = prepareL0StructuresLookupTable(l0LookupTable, &importNTHandle); + + EXPECT_EQ(result, ZE_RESULT_SUCCESS); + + EXPECT_TRUE(l0LookupTable.isSharedHandle); + EXPECT_TRUE(l0LookupTable.sharedHandleType.isSupportedHandle); + EXPECT_TRUE(l0LookupTable.sharedHandleType.isNTHandle); + EXPECT_EQ(l0LookupTable.sharedHandleType.ntHnadle, importNTHandle.handle); +} + +TEST(L0StructuresLookupTableTests, givenL0StructuresWithUnsuportedOptionsWhenPrepareLookupTableThenProperFieldsInLookupTableAreSet) { + uint64_t handle = 0x02; + ze_external_memory_import_win32_handle_t importNTHandle = {}; + importNTHandle.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_WIN32; + importNTHandle.handle = &handle; + importNTHandle.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32_KMT; + + StructuresLookupTable l0LookupTable = {}; + auto result = prepareL0StructuresLookupTable(l0LookupTable, &importNTHandle); + + EXPECT_EQ(result, ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION); + + EXPECT_TRUE(l0LookupTable.isSharedHandle); + EXPECT_FALSE(l0LookupTable.sharedHandleType.isSupportedHandle); + + ze_external_memory_import_fd_t fdStructure = {}; + fdStructure.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; + fdStructure.fd = 1; + fdStructure.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + + l0LookupTable = {}; + result = prepareL0StructuresLookupTable(l0LookupTable, &importNTHandle); + + EXPECT_EQ(result, ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION); + + EXPECT_TRUE(l0LookupTable.isSharedHandle); + EXPECT_FALSE(l0LookupTable.sharedHandleType.isSupportedHandle); + + l0LookupTable = {}; + + ze_device_module_properties_t moduleProperties = {}; + moduleProperties.stype = ZE_STRUCTURE_TYPE_DEVICE_MODULE_PROPERTIES; + result = prepareL0StructuresLookupTable(l0LookupTable, &importNTHandle); + + EXPECT_EQ(result, ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION); +} +} // namespace ult +} // namespace L0 diff --git a/level_zero/core/test/unit_tests/sources/image/test_image.cpp b/level_zero/core/test/unit_tests/sources/image/test_image.cpp index 056c7758db..dff03bf7af 100644 --- a/level_zero/core/test/unit_tests/sources/image/test_image.cpp +++ b/level_zero/core/test/unit_tests/sources/image/test_image.cpp @@ -10,6 +10,7 @@ #include "shared/source/helpers/surface_format_info.h" #include "shared/test/common/helpers/default_hw_info.h" #include "shared/test/common/mocks/mock_device.h" +#include "shared/test/common/mocks/mock_gmm_client_context.h" #include "test.h" @@ -22,59 +23,13 @@ namespace L0 { namespace ult { -struct ImageStaticFunctionConvertTypeTest : public testing::TestWithParam> { - void SetUp() override { - } - - void TearDown() override { - } -}; - -TEST_P(ImageStaticFunctionConvertTypeTest, givenZeImageFormatTypeWhenConvertTypeThenCorrectImageTypeReturned) { - auto params = GetParam(); - EXPECT_EQ(ImageImp::convertType(params.first), params.second); -} - -std::pair validTypes[] = { - {ZE_IMAGE_TYPE_2D, NEO::ImageType::Image2D}, - {ZE_IMAGE_TYPE_3D, NEO::ImageType::Image3D}, - {ZE_IMAGE_TYPE_2DARRAY, NEO::ImageType::Image2DArray}, - {ZE_IMAGE_TYPE_1D, NEO::ImageType::Image1D}, - {ZE_IMAGE_TYPE_1DARRAY, NEO::ImageType::Image1DArray}, - {ZE_IMAGE_TYPE_BUFFER, NEO::ImageType::Image1DBuffer}}; - -INSTANTIATE_TEST_CASE_P( - imageTypeFlags, - ImageStaticFunctionConvertTypeTest, - testing::ValuesIn(validTypes)); - -TEST(ImageStaticFunctionDescriptorTest, givenZeImageDescWhenConvertDescriptorThenCorrectImageDescriptorReturned) { - ze_image_desc_t zeDesc = {}; - zeDesc.arraylevels = 1u; - zeDesc.depth = 1u; - zeDesc.height = 1u; - zeDesc.width = 1u; - zeDesc.miplevels = 1u; - zeDesc.type = ZE_IMAGE_TYPE_2DARRAY; - - NEO::ImageDescriptor desc = ImageImp::convertDescriptor(zeDesc); - EXPECT_EQ(desc.fromParent, false); - EXPECT_EQ(desc.imageArraySize, zeDesc.arraylevels); - EXPECT_EQ(desc.imageDepth, zeDesc.depth); - EXPECT_EQ(desc.imageHeight, zeDesc.height); - EXPECT_EQ(desc.imageRowPitch, 0u); - EXPECT_EQ(desc.imageSlicePitch, 0u); - EXPECT_EQ(desc.imageType, NEO::ImageType::Image2DArray); - EXPECT_EQ(desc.imageWidth, zeDesc.width); - EXPECT_EQ(desc.numMipLevels, zeDesc.miplevels); - EXPECT_EQ(desc.numSamples, 0u); -} - using ImageSupport = IsAtLeastProduct; using ImageCreate = Test; +using ImageView = Test; HWTEST2_F(ImageCreate, givenValidImageDescriptionWhenImageCreateThenImageIsCreatedCorrectly, ImageSupport) { ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.arraylevels = 1u; zeDesc.depth = 1u; zeDesc.height = 1u; @@ -117,6 +72,7 @@ HWTEST2_F(ImageCreate, givenValidImageDescriptionWhenImageCreateThenImageIsCreat HWTEST2_F(ImageCreate, givenValidImageDescriptionWhenImageCreateWithUnsupportedImageThenNullPtrImageIsReturned, ImageSupport) { ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.arraylevels = 1u; zeDesc.depth = 1u; zeDesc.height = 1u; @@ -157,6 +113,7 @@ HWTEST2_F(ImageCreate, givenDifferentSwizzleFormatWhenImageInitializeThenCorrect ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -185,11 +142,133 @@ HWTEST2_F(ImageCreate, givenDifferentSwizzleFormatWhenImageInitializeThenCorrect RENDER_SURFACE_STATE::SHADER_CHANNEL_SELECT_ZERO); } +HWTEST2_F(ImageView, givenPlanarImageWhenCreateImageViewThenProperPlaneIsCreated, ImageSupport) { + using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; + const size_t width = 32; + const size_t height = 32; + const size_t depth = 1; + + ze_image_desc_t srcImgDesc = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + nullptr, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_NV12, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}, + width, + height, + depth, + 0, + 0}; + + auto imageHW = std::make_unique>>(); + auto ret = imageHW->initialize(device, &srcImgDesc); + ASSERT_EQ(ZE_RESULT_SUCCESS, ret); + + ze_image_view_planar_exp_desc_t planeYdesc = {}; + planeYdesc.stype = ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC; + planeYdesc.planeIndex = 0u; // Y plane + + ze_image_desc_t imageViewDescPlaneY = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + &planeYdesc, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_8, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_A, ZE_IMAGE_FORMAT_SWIZZLE_B, + ZE_IMAGE_FORMAT_SWIZZLE_G, ZE_IMAGE_FORMAT_SWIZZLE_R}, + width, + height, + depth, + 0, + 0}; + ze_image_handle_t planeY; + + ret = imageHW->createView(device, &imageViewDescPlaneY, &planeY); + ASSERT_EQ(ZE_RESULT_SUCCESS, ret); + + ze_image_view_planar_exp_desc_t planeUVdesc = {}; + planeUVdesc.stype = ZE_STRUCTURE_TYPE_IMAGE_VIEW_PLANAR_EXP_DESC; + planeUVdesc.planeIndex = 1u; // UV plane + + ze_image_desc_t imageViewDescPlaneUV = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + &planeUVdesc, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_8_8, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_A, ZE_IMAGE_FORMAT_SWIZZLE_B, + ZE_IMAGE_FORMAT_SWIZZLE_G, ZE_IMAGE_FORMAT_SWIZZLE_R}, + width / 2, + height / 2, + depth, + 0, + 0}; + ze_image_handle_t planeUV; + ret = imageHW->createView(device, &imageViewDescPlaneUV, &planeUV); + ASSERT_EQ(ZE_RESULT_SUCCESS, ret); + + auto nv12Allocation = imageHW->getAllocation(); + + auto planeYAllocation = Image::fromHandle(planeY)->getAllocation(); + auto planeUVAllocation = Image::fromHandle(planeUV)->getAllocation(); + + EXPECT_EQ(nv12Allocation->getGpuBaseAddress(), planeYAllocation->getGpuBaseAddress()); + EXPECT_EQ(nv12Allocation->getGpuBaseAddress(), planeUVAllocation->getGpuBaseAddress()); + + zeImageDestroy(planeY); + zeImageDestroy(planeUV); +} + +HWTEST2_F(ImageView, givenPlanarImageWhenCreateImageWithInvalidStructViewThenProperErrorIsReturned, ImageSupport) { + using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; + const size_t width = 32; + const size_t height = 32; + const size_t depth = 1; + + ze_image_desc_t srcImgDesc = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + nullptr, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_NV12, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}, + width, + height, + depth, + 0, + 0}; + + auto imageHW = std::make_unique>>(); + auto ret = imageHW->initialize(device, &srcImgDesc); + ASSERT_EQ(ZE_RESULT_SUCCESS, ret); + + ze_image_view_planar_exp_desc_t planeYdesc = {}; + planeYdesc.stype = ZE_STRUCTURE_TYPE_DEVICE_CACHE_PROPERTIES; + planeYdesc.planeIndex = 0u; // Y plane + + ze_image_desc_t imageViewDescPlaneY = {ZE_STRUCTURE_TYPE_IMAGE_DESC, + &planeYdesc, + (ZE_IMAGE_FLAG_KERNEL_WRITE | ZE_IMAGE_FLAG_BIAS_UNCACHED), + ZE_IMAGE_TYPE_2D, + {ZE_IMAGE_FORMAT_LAYOUT_8, ZE_IMAGE_FORMAT_TYPE_UINT, + ZE_IMAGE_FORMAT_SWIZZLE_A, ZE_IMAGE_FORMAT_SWIZZLE_B, + ZE_IMAGE_FORMAT_SWIZZLE_G, ZE_IMAGE_FORMAT_SWIZZLE_R}, + width, + height, + depth, + 0, + 0}; + ze_image_handle_t planeY; + + ret = imageHW->createView(device, &imageViewDescPlaneY, &planeY); + ASSERT_EQ(ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION, ret); +} + HWTEST2_F(ImageCreate, givenFDWhenCreatingImageThenSuccessIsReturned, ImageSupport) { using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -219,6 +298,7 @@ HWTEST2_F(ImageCreate, givenOpaqueFdWhenCreatingImageThenUnsuportedErrorIsReturn ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -242,39 +322,12 @@ HWTEST2_F(ImageCreate, givenOpaqueFdWhenCreatingImageThenUnsuportedErrorIsReturn ASSERT_EQ(ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION, ret); } -HWTEST2_F(ImageCreate, givenInvalidTypeStructWhenCreatingImageThenUnsuportedErrorIsReturned, ImageSupport) { - using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; - - ze_image_desc_t desc = {}; - - desc.type = ZE_IMAGE_TYPE_3D; - desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; - desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; - desc.width = 11; - desc.height = 13; - desc.depth = 17; - - desc.format.x = ZE_IMAGE_FORMAT_SWIZZLE_A; - desc.format.y = ZE_IMAGE_FORMAT_SWIZZLE_0; - desc.format.z = ZE_IMAGE_FORMAT_SWIZZLE_1; - desc.format.w = ZE_IMAGE_FORMAT_SWIZZLE_X; - - ze_external_memory_import_fd_t importFd = {}; - importFd.fd = 1; - importFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; - importFd.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC; - desc.pNext = &importFd; - - auto imageHW = std::make_unique>>(); - auto ret = imageHW->initialize(device, &desc); - ASSERT_EQ(ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION, ret); -} - HWTEST2_F(ImageCreate, givenInvalidExensionStructWhenCreatingImageThenUnsuportedErrorIsReturned, ImageSupport) { using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -297,11 +350,117 @@ HWTEST2_F(ImageCreate, givenInvalidExensionStructWhenCreatingImageThenUnsuported ASSERT_EQ(ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION, ret); } +class MemoryManagerNTHandleMock : public NEO::OsAgnosticMemoryManager { + public: + MemoryManagerNTHandleMock(NEO::ExecutionEnvironment &executionEnvironment) : NEO::OsAgnosticMemoryManager(executionEnvironment) {} + + NEO::GraphicsAllocation *createGraphicsAllocationFromNTHandle(void *handle, uint32_t rootDeviceIndex) override { + auto graphicsAllocation = createMemoryAllocation(GraphicsAllocation::AllocationType::INTERNAL_HOST_MEMORY, nullptr, reinterpret_cast(1), 1, + 4096u, reinterpret_cast(handle), MemoryPool::SystemCpuInaccessible, + rootDeviceIndex, false, false, false); + graphicsAllocation->setSharedHandle(static_cast(reinterpret_cast(handle))); + graphicsAllocation->set32BitAllocation(false); + graphicsAllocation->setDefaultGmm(new Gmm(executionEnvironment.rootDeviceEnvironments[0]->getGmmClientContext(), nullptr, 1, 0, false)); + return graphicsAllocation; + } +}; + +HWTEST2_F(ImageCreate, givenNTHandleWhenCreatingImageThenSuccessIsReturned, ImageSupport) { + using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; + + ze_image_desc_t desc = {}; + + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + desc.type = ZE_IMAGE_TYPE_3D; + desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; + desc.width = 11; + desc.height = 13; + desc.depth = 17; + + desc.format.x = ZE_IMAGE_FORMAT_SWIZZLE_A; + desc.format.y = ZE_IMAGE_FORMAT_SWIZZLE_0; + desc.format.z = ZE_IMAGE_FORMAT_SWIZZLE_1; + desc.format.w = ZE_IMAGE_FORMAT_SWIZZLE_X; + + uint64_t imageHandle = 0x1; + ze_external_memory_import_win32_handle_t importNTHandle = {}; + importNTHandle.handle = &imageHandle; + importNTHandle.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32; + importNTHandle.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_WIN32; + desc.pNext = &importNTHandle; + + NEO::MockDevice *neoDevice = nullptr; + neoDevice = NEO::MockDevice::createWithNewExecutionEnvironment(NEO::defaultHwInfo.get()); + NEO::MemoryManager *prevMemoryManager = driverHandle->getMemoryManager(); + NEO::MemoryManager *currMemoryManager = new MemoryManagerNTHandleMock(*neoDevice->executionEnvironment); + driverHandle->setMemoryManager(currMemoryManager); + neoDevice->injectMemoryManager(currMemoryManager); + + ze_result_t result = ZE_RESULT_SUCCESS; + auto device = L0::Device::create(driverHandle.get(), neoDevice, 0, false, &result); + + auto imageHW = std::make_unique>>(); + auto ret = imageHW->initialize(device, &desc); + ASSERT_EQ(ZE_RESULT_SUCCESS, ret); + ASSERT_EQ(imageHW->getAllocation()->peekSharedHandle(), NEO::toOsHandle(importNTHandle.handle)); + + imageHW.reset(nullptr); + delete device; + driverHandle->setMemoryManager(prevMemoryManager); +} + +class FailMemoryManagerMock : public NEO::OsAgnosticMemoryManager { + public: + FailMemoryManagerMock(NEO::ExecutionEnvironment &executionEnvironment) : NEO::OsAgnosticMemoryManager(executionEnvironment) {} + + NEO::GraphicsAllocation *allocateGraphicsMemoryWithProperties(const AllocationProperties &properties) override { + return nullptr; + } +}; + +HWTEST2_F(ImageCreate, givenImageDescWhenFailImageAllocationThenProperErrorIsReturned, ImageSupport) { + using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; + + ze_image_desc_t desc = {}; + + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + desc.type = ZE_IMAGE_TYPE_3D; + desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; + desc.width = 11; + desc.height = 13; + desc.depth = 17; + + desc.format.x = ZE_IMAGE_FORMAT_SWIZZLE_A; + desc.format.y = ZE_IMAGE_FORMAT_SWIZZLE_0; + desc.format.z = ZE_IMAGE_FORMAT_SWIZZLE_1; + desc.format.w = ZE_IMAGE_FORMAT_SWIZZLE_X; + + NEO::MockDevice *neoDevice = nullptr; + neoDevice = NEO::MockDevice::createWithNewExecutionEnvironment(NEO::defaultHwInfo.get()); + NEO::MemoryManager *currMemoryManager = new FailMemoryManagerMock(*neoDevice->executionEnvironment); + neoDevice->injectMemoryManager(currMemoryManager); + + ze_result_t result = ZE_RESULT_SUCCESS; + auto device = L0::Device::create(driverHandle.get(), neoDevice, 0, false, &result); + + L0::Image *imageHandle; + + auto ret = L0::Image::create(neoDevice->getHardwareInfo().platform.eProductFamily, device, &desc, &imageHandle); + + ASSERT_EQ(ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY, ret); + EXPECT_EQ(imageHandle, nullptr); + + delete device; +} + HWTEST2_F(ImageCreate, givenMediaBlockOptionWhenCopySurfaceStateThenSurfaceStateIsSet, ImageSupport) { using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE; ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -331,6 +490,7 @@ HWTEST2_P(TestImageFormats, givenValidLayoutAndTypeWhenCreateImageCoreFamilyThen auto params = GetParam(); ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.arraylevels = 1u; zeDesc.depth = 10u; zeDesc.height = 10u; @@ -623,6 +783,7 @@ using ImageGetMemoryProperties = Test; HWTEST2_F(ImageGetMemoryProperties, givenImageMemoryPropertiesExpStructureWhenGetMemroyPropertiesThenProperDataAreSet, ImageSupport) { ze_image_desc_t zeDesc = {}; + zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; zeDesc.arraylevels = 1u; zeDesc.depth = 1u; zeDesc.height = 1u; diff --git a/level_zero/core/test/unit_tests/sources/kernel/test_kernel.cpp b/level_zero/core/test/unit_tests/sources/kernel/test_kernel.cpp index f486f72d85..fc902121f5 100644 --- a/level_zero/core/test/unit_tests/sources/kernel/test_kernel.cpp +++ b/level_zero/core/test/unit_tests/sources/kernel/test_kernel.cpp @@ -156,6 +156,7 @@ HWTEST2_F(SetKernelArg, givenImageAndKernelWhenSetArgImageThenCrossThreadDataIsS ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; desc.type = ZE_IMAGE_TYPE_3D; desc.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; desc.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; @@ -1668,6 +1669,7 @@ HWTEST2_F(SetKernelArg, givenImageAndBindlessKernelWhenSetArgImageThenCopySurfac imageArg.bindless = 0x0; imageArg.bindful = undefined; ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto &hwHelper = NEO::HwHelper::get(neoDevice->getHardwareInfo().platform.eRenderCoreFamily); auto surfaceStateSize = hwHelper.getRenderSurfaceStateSize(); @@ -1693,6 +1695,7 @@ HWTEST2_F(SetKernelArg, givenImageAndBindfulKernelWhenSetArgImageThenCopySurface imageArg.bindless = undefined; imageArg.bindful = 0x40; ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto imageHW = std::make_unique>(); auto ret = imageHW->initialize(device, &desc); @@ -1720,6 +1723,7 @@ HWTEST2_F(SetKernelArg, givenSupportsMediaBlockAndIsMediaBlockImageWhenSetArgIma auto &arg = const_cast(kernel->kernelImmData->getDescriptor().payloadMappings.explicitArgs[argIndex]); auto imageHW = std::make_unique>(); ze_image_desc_t desc = {}; + desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; auto ret = imageHW->initialize(device, &desc); ASSERT_EQ(ZE_RESULT_SUCCESS, ret); auto handle = imageHW->toHandle();