Media
/
compute-runtime
mirror of https://github.com/intel/compute-runtime.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Extend image functionality 

- add imageView extension
- add import win32 NT handle
- add black box test with imageView usage example


Signed-off-by: Kamil Diedrich <kamil.diedrich@intel.com>
This commit is contained in:
Kamil Diedrich 2021-06-25 01:25:22 +00:00 committed by Compute-Runtime-Automation
parent afd5f766c2
commit af55117fa0
19 changed files with 909 additions and 139 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
level_zero/api/core/ze_core_loader.cpp
View File

@ -597,6 +597,7 @@ zeGetImageExpProcAddrTable(
ze_result_t result = ZE_RESULT_SUCCESS;
pDdiTable->pfnGetMemoryPropertiesExp = zeImageGetMemoryPropertiesExp;
pDdiTable->pfnViewCreateExp = zeImageViewCreateExp;
driver_ddiTable.core_ddiTable.ImageExp = *pDdiTable;
return result;
}

10
level_zero/api/extensions/public/ze_exp_ext.cpp
View File

@ -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

1
level_zero/core/source/CMakeLists.txt
View File

@ -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

111
level_zero/core/source/helpers/properties_parser.h Normal file
View File

@ -0,0 +1,111 @@
/*
* 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 <level_zero/ze_api.h>
#include <cstdint>
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 {
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<const ze_base_desc_t *>(desc);
while (extendedDesc) {
if (extendedDesc->stype == ZE_STRUCTURE_TYPE_IMAGE_DESC) {
const ze_image_desc_t *imageDesc = reinterpret_cast<const ze_image_desc_t *>(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<const ze_external_memory_import_fd_t *>(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<const ze_external_memory_import_win32_handle_t *>(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<const ze_image_view_planar_exp_desc_t *>(extendedDesc);
lookupTable.areImageProperties = true;
lookupTable.imageProperties.planeIndex = imageViewDesc->planeIndex;
} else {
return ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
}
extendedDesc = reinterpret_cast<const ze_base_desc_t *>(extendedDesc->pNext);
}
return ZE_RESULT_SUCCESS;
}
} // namespace L0

6
level_zero/core/source/image/image.h
View File

@ -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 <typename Type>
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,

1
level_zero/core/source/image/image_hw.h
View File

@ -15,7 +15,6 @@
#include "level_zero/core/source/image/image_imp.h"
namespace L0 {
template <GFXCORE_FAMILY gfxCoreFamily>
struct ImageCoreFamily : public ImageImp {
using GfxFamily = typename NEO::GfxFamilyMapper<gfxCoreFamily>::GfxFamily;

54
level_zero/core/source/image/image_hw.inl
View File

@ -16,18 +16,26 @@
#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 <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t ImageCoreFamily<gfxCoreFamily>::initialize(Device *device, const ze_image_desc_t *desc) {
using RENDER_SURFACE_STATE = typename GfxFamily::RENDER_SURFACE_STATE;
StructuresLookupTable lookupTable = {};
auto parseResult = prepareL0StructuresLookupTable(lookupTable, desc);
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<ze_image_desc_t *>(desc);
@ -57,34 +65,36 @@ ze_result_t ImageCoreFamily<gfxCoreFamily>::initialize(Device *device, const ze_
}
imgInfo.linearStorage = surfaceType == RENDER_SURFACE_STATE::SURFACE_TYPE_SURFTYPE_1D;
imgInfo.plane = GMM_NO_PLANE;
imgInfo.plane = isImageView ? static_cast<GMM_YUV_PLANE>(lookupTable.imageProperties.planeIndex + 1u) : GMM_NO_PLANE;
imgInfo.useLocalMemory = false;
imgInfo.preferRenderCompression = false;
if (desc->pNext) {
const ze_base_desc_t *extendedDesc = reinterpret_cast<const ze_base_desc_t *>(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<const ze_external_memory_import_fd_t *>(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<const NEO::RootDeviceEnvironment &>(device->getNEODevice()->getRootDeviceEnvironment()).getGmmHelper();

25
level_zero/core/source/image/image_imp.cpp
View File

@ -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<ImageImp *>((*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;

38
level_zero/core/source/image/image_imp.h
View File

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

2
level_zero/core/test/black_box_tests/CMakeLists.txt
View File

@ -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()

4
level_zero/core/test/black_box_tests/zello_copy_image.cpp
View File

@ -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<const ze_image_desc_t *>(&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,

301
level_zero/core/test/black_box_tests/zello_image_view.cpp Normal file
View File

@ -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<const ze_image_desc_t *>(&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<uint8_t> srcVecY;
srcVecY.resize(width * height);
for (size_t i = 0; i < width * height; ++i) {
srcVecY[i] = static_cast<uint8_t>(i);
}
std::vector<uint8_t> srcVecUV;
srcVecUV.resize((width / 2) * (height));
for (size_t i = 0; i < (width / 2) * (height); ++i) {
if (i % 2 == 0) {
srcVecUV[i] = static_cast<uint8_t>(0x33);
} else {
srcVecUV[i] = static_cast<uint8_t>(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<uint8_t> dstVecY;
dstVecY.resize(width * height);
// destination buffer for UV plane
std::vector<uint8_t> 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<uint32_t>::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;
}

15
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp
View File

@ -413,6 +413,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListWhenCopyFromMemoryToImageThenBl
void *srcPtr = reinterpret_cast<void *>(0x1234);
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHW = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHW->initialize(device, &zeDesc);
@ -427,6 +428,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhenIma
void *srcPtr = reinterpret_cast<void *>(0x1234);
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHW = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHW->initialize(device, &zeDesc);
@ -446,6 +448,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhenIma
void *dstPtr = reinterpret_cast<void *>(0x1234);
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHW = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHW->initialize(device, &zeDesc);
@ -465,6 +468,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen1DI
void *srcPtr = reinterpret_cast<void *>(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<void *>(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<void *>(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<void *>(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<void *>(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<void *>(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<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
@ -594,6 +603,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListAndNullDestinationRegionWhen2DI
void *dstPtr = reinterpret_cast<void *>(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<void *>(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<void *>(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<void *>(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<void *>(0x1234);
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHW = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHW->initialize(device, &zeDesc);
@ -688,6 +702,7 @@ HWTEST2_F(CommandListCreate, givenCopyCommandListWhenCopyFromImageToImageThenBli
MockCommandListHw<gfxCoreFamily> cmdList;
cmdList.initialize(device, NEO::EngineGroupType::Copy);
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHWSrc = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
auto imageHWDst = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHWSrc->initialize(device, &zeDesc);

1
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_blit.cpp
View File

@ -278,6 +278,7 @@ HWTEST2_F(AppendMemoryCopy, givenCopyCommandListWhenCopyFromImagBlitThenCommandA
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::Copy, returnValue));
ze_image_desc_t zeDesc = {};
zeDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHWSrc = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
auto imageHWDst = std::make_unique<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
imageHWSrc->initialize(device, &zeDesc);

1
level_zero/core/test/unit_tests/sources/context/test_context.cpp
View File

@ -637,6 +637,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);

1
level_zero/core/test/unit_tests/sources/helper/CMakeLists.txt
View File

@ -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()

159
level_zero/core/test/unit_tests/sources/helper/properties_parser_tests.cpp Normal file
View File

@ -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<std::pair<ze_image_type_t, NEO::ImageType>> {
void SetUp() override {
}
void TearDown() override {
}
};
TEST_P(ImageStaticFunctionConvertTypeTest, givenZeImageFormatTypeWhenConvertTypeThenCorrectImageTypeReturned) {
auto params = GetParam();
EXPECT_EQ(convertType(params.first), params.second);
}
std::pair<ze_image_type_t, NEO::ImageType> 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

313
level_zero/core/test/unit_tests/sources/image/test_image.cpp
View File

@ -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<std::pair<ze_image_type_t, NEO::ImageType>> {
void SetUp() override {
}
void TearDown() override {
}
};
TEST_P(ImageStaticFunctionConvertTypeTest, givenZeImageFormatTypeWhenConvertTypeThenCorrectImageTypeReturned) {
auto params = GetParam();
EXPECT_EQ(ImageImp::convertType(params.first), params.second);
}
std::pair<ze_image_type_t, NEO::ImageType> 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<IGFX_SKYLAKE>;
using ImageCreate = Test<DeviceFixture>;
using ImageView = Test<DeviceFixture>;
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<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
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<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
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<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
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<void *>(1), 1,
4096u, reinterpret_cast<uint64_t>(handle), MemoryPool::SystemCpuInaccessible,
rootDeviceIndex, false, false, false);
graphicsAllocation->setSharedHandle(static_cast<osHandle>(reinterpret_cast<uint64_t>(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::MockDevice>(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<WhiteBox<::L0::ImageCoreFamily<gfxCoreFamily>>>();
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::MockDevice>(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<DeviceFixture>;
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;

4
level_zero/core/test/unit_tests/sources/kernel/test_kernel.cpp
View File

@ -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;
@ -1603,6 +1604,7 @@ HWTEST2_F(SetKernelArg, givenImageAndBindlessKernelWhenSetArgImageThenCopySurfac
imageArg.bindless = 0x0;
imageArg.bindful = undefined<SurfaceStateHeapOffset>;
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();
@ -1628,6 +1630,7 @@ HWTEST2_F(SetKernelArg, givenImageAndBindfulKernelWhenSetArgImageThenCopySurface
imageArg.bindless = undefined<CrossThreadDataOffset>;
imageArg.bindful = 0x40;
ze_image_desc_t desc = {};
desc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
auto imageHW = std::make_unique<MyMockImage<gfxCoreFamily>>();
auto ret = imageHW->initialize(device, &desc);
@ -1655,6 +1658,7 @@ HWTEST2_F(SetKernelArg, givenSupportsMediaBlockAndIsMediaBlockImageWhenSetArgIma
auto &arg = const_cast<NEO::ArgDescriptor &>(kernel->kernelImmData->getDescriptor().payloadMappings.explicitArgs[argIndex]);
auto imageHW = std::make_unique<MyMockImageMediaBlock<gfxCoreFamily>>();
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();