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

AppendLaunchKernel tests 

RelatedTo: NEO-4515

Change-Id: I1f719ea1f60f313fba44d49f84fe2caf6ea4e5c4
Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>
This commit is contained in:
Mateusz Hoppe 2020-04-03 23:05:30 +02:00 committed by sys_ocldev
parent 884ae6fbab
commit bd247d725b
50 changed files with 353 additions and 108 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
level_zero/cmake/l0_tests.cmake
View File

@ -79,11 +79,11 @@ add_library(compute_runtime_mockable_extra
EXCLUDE_FROM_ALL
${CMAKE_CURRENT_LIST_DIR}/l0_tests.cmake
${NEO_SHARED_TEST_DIRECTORY}/unit_test/utilities/cpuintrinsics.cpp
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/test_files.cpp
${COMPUTE_RUNTIME_DIR}/opencl/source/aub/aub_stream_interface.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/abort.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/helpers/built_ins_helper.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/helpers/debug_helpers.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/helpers/test_files.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/libult/os_interface.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/libult/source_level_debugger_ult.cpp
${COMPUTE_RUNTIME_DIR}/opencl/test/unit_test/libult/source_level_debugger_library.cpp

20
level_zero/core/test/unit_tests/CMakeLists.txt
View File

@ -8,6 +8,8 @@ link_libraries(${ASAN_LIBS} ${TSAN_LIBS})
set(TARGET_NAME ${TARGET_NAME_L0}_core_tests)
include(gen_kernel.cmake)
append_sources_from_properties(L0_CORE_ENABLERS NEO_CORE_SRCS_LINK)
function(ADD_SUPPORTED_TEST_PRODUCT_FAMILIES_DEFINITION)
@ -28,6 +30,7 @@ add_executable(${TARGET_NAME}
)
target_sources(${TARGET_NAME} PRIVATE
${CMAKE_CURRENT_LIST_DIR}/gen_kernel.cmake
${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
${CMAKE_CURRENT_SOURCE_DIR}/main.cpp
${CMAKE_CURRENT_SOURCE_DIR}/mock.h
@ -118,6 +121,23 @@ target_sources(${TARGET_NAME} PRIVATE
$<TARGET_OBJECTS:${TARGET_NAME_L0}_mocks>
)
set(TEST_MODULES
${CMAKE_CURRENT_SOURCE_DIR}/test_modules/test_kernel.cl
)
macro(macro_for_each_gen)
foreach(PLATFORM_TYPE ${PLATFORM_TYPES})
if(${GEN_TYPE}_HAS_${PLATFORM_TYPE})
get_family_name_with_type(${GEN_TYPE} ${PLATFORM_TYPE})
string(TOLOWER ${PLATFORM_TYPE} PLATFORM_TYPE_LOWER)
set(PLATFORM_LOWER ${DEFAULT_SUPPORTED_${GEN_TYPE}_${PLATFORM_TYPE}_PLATFORM})
level_zero_gen_kernels(level_zero_test_kernels_${family_name_with_type} ${PLATFORM_LOWER} ${family_name_with_type} ${TEST_MODULES})
add_dependencies(${TARGET_NAME} level_zero_test_kernels_${family_name_with_type})
endif()
endforeach()
endmacro()
apply_macro_for_each_gen("TESTED")
option(L0_ULT_VERBOSE "Use the default/verbose test output" OFF)
if(NOT L0_ULT_VERBOSE)
set(L0_TESTS_LISTENER_OPTION "--disable_default_listener")

1
level_zero/core/test/unit_tests/fixtures/CMakeLists.txt
View File

@ -10,6 +10,7 @@ set(L0_FIXTURES_SOURCES
${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
${CMAKE_CURRENT_SOURCE_DIR}/device_fixture.h
${CMAKE_CURRENT_SOURCE_DIR}/device_fixture.cpp
${CMAKE_CURRENT_SOURCE_DIR}/module_fixture.h
)
add_library(${TARGET_NAME} OBJECT ${L0_FIXTURES_SOURCES})

11
level_zero/core/test/unit_tests/fixtures/device_fixture.h
View File

@ -11,6 +11,7 @@
#include "shared/test/unit_test/mocks/mock_device.h"
#include "level_zero/core/test/unit_tests/mocks/mock_device.h"
#include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
namespace L0 {
namespace ult {
@ -18,14 +19,18 @@ namespace ult {
struct DeviceFixture {
void SetUp() {
neoDevice = NEO::MockDevice::createWithNewExecutionEnvironment<NEO::MockDevice>(NEO::defaultHwInfo.get());
device = std::make_unique<Mock<L0::DeviceImp>>(neoDevice, neoDevice->getExecutionEnvironment());
NEO::DeviceVector devices;
devices.push_back(std::unique_ptr<NEO::Device>(neoDevice));
driverHandle = std::make_unique<Mock<L0::DriverHandleImp>>();
driverHandle->initialize(std::move(devices));
device = driverHandle->devices[0];
}
void TearDown() {
}
std::unique_ptr<Mock<L0::DriverHandleImp>> driverHandle;
NEO::MockDevice *neoDevice = nullptr;
std::unique_ptr<Mock<L0::DeviceImp>> device = nullptr;
L0::Device *device = nullptr;
};
} // namespace ult

68
level_zero/core/test/unit_tests/fixtures/module_fixture.h Normal file
View File

@ -0,0 +1,68 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#pragma once
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "level_zero/core/source/module/module.h"
#include "level_zero/core/test/unit_tests/fixtures/device_fixture.h"
#include "level_zero/core/test/unit_tests/mocks/mock_kernel.h"
namespace L0 {
namespace ult {
struct ModuleFixture : public DeviceFixture {
void SetUp() {
DeviceFixture::SetUp();
createModuleFromBinary();
}
void createModuleFromBinary() {
std::string testFile;
retrieveBinaryKernelFilename(testFile, binaryFilename + "_", ".bin");
size_t size = 0;
auto src = loadDataFromFile(
testFile.c_str(),
size);
ASSERT_NE(0u, size);
ASSERT_NE(nullptr, src);
ze_module_desc_t moduleDesc = {ZE_MODULE_DESC_VERSION_CURRENT};
moduleDesc.format = ZE_MODULE_FORMAT_NATIVE;
moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(src.get());
moduleDesc.inputSize = size;
ModuleBuildLog *moduleBuildLog = nullptr;
module.reset(Module::create(device, &moduleDesc, neoDevice, moduleBuildLog));
}
void createKernel() {
ze_kernel_desc_t desc = {ZE_KERNEL_DESC_VERSION_CURRENT};
desc.pKernelName = kernelName.c_str();
kernel = std::make_unique<WhiteBox<::L0::Kernel>>();
kernel->module = module.get();
kernel->initialize(&desc);
}
void TearDown() {
DeviceFixture::TearDown();
}
const std::string binaryFilename = "test_kernel";
const std::string kernelName = "test";
std::unique_ptr<L0::Module> module;
std::unique_ptr<WhiteBox<::L0::Kernel>> kernel;
};
} // namespace ult
} // namespace L0

47
level_zero/core/test/unit_tests/gen_kernel.cmake Normal file
View File

@ -0,0 +1,47 @@
#
# Copyright (C) 2020 Intel Corporation
#
# SPDX-License-Identifier: MIT
#
function(level_zero_gen_kernels target platform_name suffix)
if(NOT DEFINED cloc_cmd_prefix)
if(WIN32)
set(cloc_cmd_prefix ocloc)
else()
if(DEFINED NEO__IGC_LIBRARY_PATH)
set(cloc_cmd_prefix LD_LIBRARY_PATH=${NEO__IGC_LIBRARY_PATH}:$<TARGET_FILE_DIR:ocloc_lib> $<TARGET_FILE:ocloc>)
else()
set(cloc_cmd_prefix LD_LIBRARY_PATH=$<TARGET_FILE_DIR:ocloc_lib> $<TARGET_FILE:ocloc>)
endif()
endif()
endif()
set(outputdir "${TargetDir}/level_zero/${suffix}/test_files/${NEO_ARCH}/")
set(results)
foreach(filepath ${ARGN})
get_filename_component(filename ${filepath} NAME)
get_filename_component(basename ${filepath} NAME_WE)
get_filename_component(workdir ${filepath} DIRECTORY)
set(outputpath_base "${outputdir}${basename}_${suffix}")
set(output_files
${outputpath_base}.bin
${outputpath_base}.gen
)
add_custom_command(
COMMAND echo generate ${cloc_cmd_prefix} -q -file ${filename} -device ${platform_name} -out_dir ${outputdir}
OUTPUT ${output_files}
COMMAND ${cloc_cmd_prefix} -q -file ${filename} -device ${platform_name} -out_dir ${outputdir}
WORKING_DIRECTORY ${workdir}
DEPENDS ${filepath} ocloc
)
list(APPEND results ${output_files})
endforeach()
add_custom_target(${target} DEPENDS ${results} copy_compiler_files)
set_target_properties(${target} PROPERTIES FOLDER ${TARGET_NAME_L0})
endfunction()

11
level_zero/core/test/unit_tests/main.cpp
View File

@ -10,6 +10,7 @@
#include "shared/source/os_interface/hw_info_config.h"
#include "shared/test/unit_test/helpers/default_hw_info.inl"
#include "shared/test/unit_test/helpers/memory_leak_listener.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/helpers/ult_hw_config.inl"
#include "opencl/source/program/kernel_info.h"
@ -192,6 +193,16 @@ int main(int argc, char **argv) {
listeners.Append(customEventListener);
}
binaryNameSuffix.append(NEO::familyName[hwInfoForTests.platform.eRenderCoreFamily]);
binaryNameSuffix.append(hwInfoForTests.capabilityTable.platformType);
std::string testBinaryFiles = getRunPath(argv[0]);
testBinaryFiles.append("/level_zero/");
testBinaryFiles.append(binaryNameSuffix);
testBinaryFiles.append("/");
testBinaryFiles.append(testFiles);
testFiles = testBinaryFiles;
listeners.Append(new NEO::MemoryLeakListener);
NEO::GmmHelper::createGmmContextWrapperFunc =

23
level_zero/core/test/unit_tests/mocks/mock_kernel.h
View File

@ -31,6 +31,29 @@ struct WhiteBox<::L0::KernelImmutableData> : public ::L0::KernelImmutableData {
WhiteBox() : ::L0::KernelImmutableData() {}
};
template <>
struct WhiteBox<::L0::Kernel> : public ::L0::KernelImp {
using BaseClass = ::L0::KernelImp;
using ::L0::KernelImp::createPrintfBuffer;
using ::L0::KernelImp::crossThreadData;
using ::L0::KernelImp::crossThreadDataSize;
using ::L0::KernelImp::groupSize;
using ::L0::KernelImp::kernelImmData;
using ::L0::KernelImp::module;
using ::L0::KernelImp::perThreadDataForWholeThreadGroup;
using ::L0::KernelImp::perThreadDataSize;
using ::L0::KernelImp::printfBuffer;
using ::L0::KernelImp::residencyContainer;
using ::L0::KernelImp::unifiedMemoryControls;
void setBufferSurfaceState(uint32_t argIndex, void *address,
NEO::GraphicsAllocation *alloc) override {}
std::unique_ptr<Kernel> clone() const override { return nullptr; }
WhiteBox() : ::L0::KernelImp(nullptr) {}
};
template <>
struct Mock<::L0::Kernel> : public ::L0::KernelImp {
using BaseClass = ::L0::KernelImp;

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

@ -8,4 +8,5 @@ target_sources(${TARGET_NAME} PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
${CMAKE_CURRENT_SOURCE_DIR}/test_cmdlist.cpp
${CMAKE_CURRENT_SOURCE_DIR}/test_cmdlist_api.cpp
${CMAKE_CURRENT_SOURCE_DIR}/test_cmdlist_append_launch_kernel.cpp
)

48
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist.cpp
View File

@ -32,10 +32,10 @@ TEST(zeCommandListCreateImmediate, redirectsToObject) {
}
TEST_F(CommandListCreate, whenCommandListIsCreatedThenItIsInitialized) {
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device.get()));
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device));
ASSERT_NE(nullptr, commandList);
EXPECT_EQ(device.get(), commandList->device);
EXPECT_EQ(device, commandList->device);
ASSERT_GT(commandList->commandContainer.getCmdBufferAllocations().size(), 0u);
auto numAllocations = 0u;
@ -57,7 +57,7 @@ TEST_F(CommandListCreate, whenCommandListIsCreatedThenItIsInitialized) {
}
TEST_F(CommandListCreate, givenRegularCommandListThenDefaultNumIddPerBlockIsUsed) {
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device.get()));
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device));
ASSERT_NE(nullptr, commandList);
const uint32_t defaultNumIdds = CommandList::defaultNumIddsPerBlock;
@ -71,7 +71,7 @@ TEST_F(CommandListCreate, givenImmediateCommandListThenCustomNumIddPerBlockUsed)
ZE_COMMAND_QUEUE_MODE_DEFAULT,
ZE_COMMAND_QUEUE_PRIORITY_NORMAL,
0};
std::unique_ptr<L0::CommandList> commandList(CommandList::createImmediate(productFamily, device.get(), &desc, false));
std::unique_ptr<L0::CommandList> commandList(CommandList::createImmediate(productFamily, device, &desc, false));
ASSERT_NE(nullptr, commandList);
const uint32_t cmdListImmediateIdds = CommandList::commandListimmediateIddsPerBlock;
@ -85,23 +85,23 @@ TEST_F(CommandListCreate, whenCreatingImmediateCommandListThenItHasImmediateComm
ZE_COMMAND_QUEUE_MODE_DEFAULT,
ZE_COMMAND_QUEUE_PRIORITY_NORMAL,
0};
std::unique_ptr<L0::CommandList> commandList(CommandList::createImmediate(productFamily, device.get(), &desc, false));
std::unique_ptr<L0::CommandList> commandList(CommandList::createImmediate(productFamily, device, &desc, false));
ASSERT_NE(nullptr, commandList);
EXPECT_EQ(device.get(), commandList->device);
EXPECT_EQ(device, commandList->device);
EXPECT_EQ(1u, commandList->cmdListType);
EXPECT_NE(nullptr, commandList->cmdQImmediate);
}
TEST_F(CommandListCreate, givenInvalidProductFamilyThenReturnsNullPointer) {
std::unique_ptr<L0::CommandList> commandList(CommandList::create(IGFX_UNKNOWN, device.get()));
std::unique_ptr<L0::CommandList> commandList(CommandList::create(IGFX_UNKNOWN, device));
EXPECT_EQ(nullptr, commandList);
}
HWTEST_F(CommandListCreate, whenCommandListIsCreatedThenStateBaseAddressCmdIsAddedAndCorrectlyProgrammed) {
using STATE_BASE_ADDRESS = typename FamilyType::STATE_BASE_ADDRESS;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device.get()));
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device));
auto &commandContainer = commandList->commandContainer;
auto gmmHelper = commandContainer.getDevice()->getGmmHelper();
@ -141,37 +141,5 @@ HWTEST_F(CommandListCreate, whenCommandListIsCreatedThenStateBaseAddressCmdIsAdd
EXPECT_EQ(gmmHelper->getMOCS(GMM_RESOURCE_USAGE_OCL_BUFFER), cmdSba->getStatelessDataPortAccessMemoryObjectControlState());
}
HWTEST_F(CommandListCreate, givenNotEnoughSpaceInCommandStreamWhenAppendingFunctionThenBbEndIsAddedAndNewCmdBufferAllocated) {
using MI_BATCH_BUFFER_END = typename FamilyType::MI_BATCH_BUFFER_END;
Mock<Kernel> kernel;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device.get()));
auto &commandContainer = commandList->commandContainer;
const auto stream = commandContainer.getCommandStream();
const auto streamCpu = stream->getCpuBase();
auto available = stream->getAvailableSpace();
stream->getSpace(available - sizeof(MI_BATCH_BUFFER_END) - 16);
auto bbEndPosition = stream->getSpace(0);
ze_group_count_t dispatchFunctionArguments{1, 1, 1};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchFunctionArguments, nullptr, 0, nullptr);
auto usedSpaceAfter = commandContainer.getCommandStream()->getUsed();
ASSERT_GT(usedSpaceAfter, 0u);
const auto streamCpu2 = stream->getCpuBase();
EXPECT_NE(nullptr, streamCpu2);
EXPECT_NE(streamCpu, streamCpu2);
EXPECT_EQ(2u, commandContainer.getCmdBufferAllocations().size());
GenCmdList cmdList;
FamilyType::PARSE::parseCommandBuffer(cmdList, bbEndPosition, 2 * sizeof(MI_BATCH_BUFFER_END));
auto itor = find<MI_BATCH_BUFFER_END *>(cmdList.begin(), cmdList.end());
EXPECT_NE(cmdList.end(), itor);
}
} // namespace ult
} // namespace L0

86
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel.cpp Normal file
View File

@ -0,0 +1,86 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/test/unit_test/cmd_parse/gen_cmd_parse.h"
#include "opencl/source/helpers/hardware_commands_helper.h"
#include "test.h"
#include "level_zero/core/test/unit_tests/fixtures/module_fixture.h"
namespace L0 {
namespace ult {
using CommandListAppendLaunchKernel = Test<ModuleFixture>;
HWTEST_F(CommandListAppendLaunchKernel, givenNotEnoughSpaceInCommandStreamWhenAppendingKernelThenBbEndIsAddedAndNewCmdBufferAllocated) {
using MI_BATCH_BUFFER_END = typename FamilyType::MI_BATCH_BUFFER_END;
createKernel();
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device));
auto &commandContainer = commandList->commandContainer;
const auto stream = commandContainer.getCommandStream();
const auto streamCpu = stream->getCpuBase();
auto available = stream->getAvailableSpace();
stream->getSpace(available - sizeof(MI_BATCH_BUFFER_END) - 16);
auto bbEndPosition = stream->getSpace(0);
ze_group_count_t groupCount{1, 1, 1};
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr);
auto usedSpaceAfter = commandContainer.getCommandStream()->getUsed();
ASSERT_GT(usedSpaceAfter, 0u);
const auto streamCpu2 = stream->getCpuBase();
EXPECT_NE(nullptr, streamCpu2);
EXPECT_NE(streamCpu, streamCpu2);
EXPECT_EQ(2u, commandContainer.getCmdBufferAllocations().size());
GenCmdList cmdList;
FamilyType::PARSE::parseCommandBuffer(cmdList, bbEndPosition, 2 * sizeof(MI_BATCH_BUFFER_END));
auto itor = find<MI_BATCH_BUFFER_END *>(cmdList.begin(), cmdList.end());
EXPECT_NE(cmdList.end(), itor);
}
HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenFunctionWhenBindingTablePrefetchAllowedThenProgramBindingTableEntryCount) {
using MEDIA_INTERFACE_DESCRIPTOR_LOAD = typename FamilyType::MEDIA_INTERFACE_DESCRIPTOR_LOAD;
using INTERFACE_DESCRIPTOR_DATA = typename FamilyType::INTERFACE_DESCRIPTOR_DATA;
createKernel();
ze_group_count_t groupCount{1, 1, 1};
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device));
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr);
auto commandStream = commandList->commandContainer.getCommandStream();
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, commandStream->getCpuBase(), commandStream->getUsed()));
auto itorMIDL = find<MEDIA_INTERFACE_DESCRIPTOR_LOAD *>(cmdList.begin(), cmdList.end());
ASSERT_NE(itorMIDL, cmdList.end());
auto cmd = genCmdCast<MEDIA_INTERFACE_DESCRIPTOR_LOAD *>(*itorMIDL);
ASSERT_NE(cmd, nullptr);
auto dsh = commandList->commandContainer.getIndirectHeap(NEO::HeapType::DYNAMIC_STATE);
auto idd = static_cast<INTERFACE_DESCRIPTOR_DATA *>(ptrOffset(dsh->getCpuBase(), cmd->getInterfaceDescriptorDataStartAddress()));
if (NEO::HardwareCommandsHelper<FamilyType>::doBindingTablePrefetch()) {
uint32_t numArgs = kernel->kernelImmData->getDescriptor().payloadMappings.bindingTable.numEntries;
EXPECT_EQ(numArgs, idd->getBindingTableEntryCount());
} else {
EXPECT_EQ(0u, idd->getBindingTableEntryCount());
}
}
} // namespace ult
} // namespace L0

4
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue.cpp
View File

@ -26,7 +26,7 @@ TEST_F(CommandQueueCreate, whenCreatingCommandQueueThenItIsInitialized) {
auto csr = std::unique_ptr<NEO::CommandStreamReceiver>(neoDevice->createCommandStreamReceiver());
L0::CommandQueue *commandQueue = CommandQueue::create(productFamily,
device.get(),
device,
csr.get(),
&desc);
ASSERT_NE(nullptr, commandQueue);
@ -34,7 +34,7 @@ TEST_F(CommandQueueCreate, whenCreatingCommandQueueThenItIsInitialized) {
L0::CommandQueueImp *commandQueueImp = reinterpret_cast<L0::CommandQueueImp *>(commandQueue);
EXPECT_EQ(csr.get(), commandQueueImp->getCsr());
EXPECT_EQ(device.get(), commandQueueImp->getDevice());
EXPECT_EQ(device, commandQueueImp->getDevice());
EXPECT_EQ(0u, commandQueueImp->getTaskCount());
commandQueue->destroy();

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

@ -11,7 +11,7 @@
#include "test.h"
#include "level_zero/core/source/image/image_hw.h"
#include "level_zero/core/test/unit_tests/mocks/mock_device.h"
#include "level_zero/core/test/unit_tests/fixtures/device_fixture.h"
namespace L0 {
namespace ult {
@ -42,17 +42,7 @@ INSTANTIATE_TEST_CASE_P(
ImageStaticFunctionConvertTypeTest,
testing::ValuesIn(validTypes));
struct ImageStaticFunctionDescriptorFixture {
void SetUp() {
}
void TearDown() {
}
};
using ImageStaticFunctionDescriptorTest = Test<ImageStaticFunctionDescriptorFixture>;
TEST_F(ImageStaticFunctionDescriptorTest, givenZeImageDescWhenConvertDescriptorThenCorrectImageDescriptorReturned) {
TEST(ImageStaticFunctionDescriptorTest, givenZeImageDescWhenConvertDescriptorThenCorrectImageDescriptorReturned) {
ze_image_desc_t zeDesc = {};
zeDesc.arraylevels = 1u;
zeDesc.depth = 1u;
@ -74,19 +64,6 @@ TEST_F(ImageStaticFunctionDescriptorTest, givenZeImageDescWhenConvertDescriptorT
EXPECT_EQ(desc.numSamples, 0u);
}
struct DeviceFixture {
void SetUp() {
neoDevice = NEO::MockDevice::createWithNewExecutionEnvironment<NEO::MockDevice>(NEO::defaultHwInfo.get());
device = std::make_unique<Mock<L0::DeviceImp>>(neoDevice, neoDevice->getExecutionEnvironment());
}
void TearDown() {
}
NEO::MockDevice *neoDevice = nullptr;
std::unique_ptr<Mock<L0::DeviceImp>> device = nullptr;
};
using ImageSupport = IsAtMostProduct<IGFX_TIGERLAKE_LP>;
using ImageCreate = Test<DeviceFixture>;
@ -108,7 +85,7 @@ HWTEST2_F(ImageCreate, givenValidImageDescriptionWhenImageCreateThenImageIsCreat
ZE_IMAGE_FORMAT_SWIZZLE_B,
ZE_IMAGE_FORMAT_SWIZZLE_A};
std::unique_ptr<L0::Image> image(Image::create(productFamily, device.get(), &zeDesc));
std::unique_ptr<L0::Image> image(Image::create(productFamily, device, &zeDesc));
ASSERT_NE(image, nullptr);

37
level_zero/core/test/unit_tests/test_modules/test_kernel.cl Normal file
View File

@ -0,0 +1,37 @@
/*
* Copyright (C) 2020 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
const sampler_t sampler =
CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
kernel void test(const global float *a, const global float *b,
global float *c,
read_only image2d_t input,
write_only image2d_t output) {
const int global_id = get_global_id(0);
const int local_id = get_local_id(0);
local float a_local[16];
float sum = 0.0f;
a_local[local_id] = a[local_id] + b[local_id];
barrier(CLK_LOCAL_MEM_FENCE);
for (int i = 0; i < get_local_size(0); ++i) {
sum += a_local[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
c[global_id] = sum;
int2 coord = {get_global_id(0), get_global_id(1)};
uint4 pixel = read_imageui(input, sampler, coord);
write_imageui(output, coord, pixel);
printf("local_id = %d, global_id = %d \n", local_id, global_id);
}

2
opencl/test/unit_test/api/cl_build_program_tests.inl
View File

@ -8,12 +8,12 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/source/program/kernel_info.h"
#include "opencl/source/program/program.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_clone_kernel_tests.inl
View File

@ -6,9 +6,9 @@
*/
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_compile_program_tests.inl
View File

@ -7,10 +7,10 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_create_kernel_tests.inl
View File

@ -6,10 +6,10 @@
*/
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/source/program/kernel_info.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_program.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_create_kernels_in_program_tests.inl
View File

@ -6,9 +6,9 @@
*/
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_create_program_with_binary_tests.inl
View File

@ -6,9 +6,9 @@
*/
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_get_kernel_arg_info_tests.inl
View File

@ -8,10 +8,10 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"
#include "compiler_options.h"

2
opencl/test/unit_test/api/cl_get_kernel_info_tests.inl
View File

@ -8,10 +8,10 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_get_kernel_work_group_info_tests.inl
View File

@ -8,9 +8,9 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_kernel.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_get_mem_object_info_tests.inl
View File

@ -8,9 +8,9 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_get_program_build_info_tests.inl
View File

@ -12,10 +12,10 @@
#include "shared/source/device_binary_format/elf/elf_encoder.h"
#include "shared/source/device_binary_format/elf/ocl_elf.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_get_program_info_tests.inl
View File

@ -8,10 +8,10 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/api/cl_link_program_tests.inl
View File

@ -8,10 +8,10 @@
#include "shared/source/compiler_interface/compiler_interface.h"
#include "shared/source/device_binary_format/elf/elf_decoder.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"
#include "compiler_options.h"

2
opencl/test/unit_test/api/cl_release_kernel_tests.inl
View File

@ -6,9 +6,9 @@
*/
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/context/context.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "cl_api_tests.h"

2
opencl/test/unit_test/aub_tests/fixtures/run_kernel_fixture.h
View File

@ -9,6 +9,7 @@
#include "shared/source/command_stream/command_stream_receiver.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/program/program.h"
#include "opencl/test/unit_test/aub_tests/command_queue/command_enqueue_fixture.h"
@ -18,7 +19,6 @@
#include "opencl/test/unit_test/command_stream/command_stream_fixture.h"
#include "opencl/test/unit_test/fixtures/run_kernel_fixture.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
namespace NEO {

2
opencl/test/unit_test/built_ins/sip_tests.cpp
View File

@ -7,10 +7,10 @@
#include "shared/source/built_ins/built_ins.h"
#include "shared/source/built_ins/sip.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/mocks/mock_device.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_program.h"
#include "test.h"

2
opencl/test/unit_test/fixtures/built_in_fixture.cpp
View File

@ -9,10 +9,10 @@
#include "shared/source/built_ins/built_ins.h"
#include "shared/source/device/device.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
using namespace NEO;

2
opencl/test/unit_test/fixtures/hello_world_kernel_fixture.h
View File

@ -8,13 +8,13 @@
#pragma once
#include "shared/source/device/device.h"
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/kernel/kernel.h"
#include "opencl/source/platform/platform.h"
#include "opencl/source/program/program.h"
#include "opencl/test/unit_test/fixtures/device_fixture.h"
#include "opencl/test/unit_test/fixtures/program_fixture.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_context.h"
#include "opencl/test/unit_test/mocks/mock_kernel.h"
#include "opencl/test/unit_test/mocks/mock_program.h"

2
opencl/test/unit_test/fixtures/program_fixture.h
View File

@ -7,9 +7,9 @@
#pragma once
#include "shared/source/helpers/file_io.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/program/program.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_program.h"
#include "gtest/gtest.h"

3
opencl/test/unit_test/fixtures/run_kernel_fixture.h
View File

@ -7,8 +7,9 @@
#pragma once
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
namespace NEO {
struct CommandQueueHwFixture;

2
opencl/test/unit_test/gen9/sip_tests_gen9.cpp
View File

@ -7,10 +7,10 @@
#include "shared/source/built_ins/built_ins.h"
#include "shared/source/built_ins/sip.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/mocks/mock_device.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "test.h"
#include "gtest/gtest.h"

2
opencl/test/unit_test/gtpin/gtpin_tests.cpp
View File

@ -13,6 +13,7 @@
#include "shared/source/memory_manager/surface.h"
#include "shared/source/os_interface/os_context.h"
#include "shared/test/unit_test/device_binary_format/patchtokens_tests.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/helpers/variable_backup.h"
#include "shared/test/unit_test/mocks/mock_device.h"
@ -30,7 +31,6 @@
#include "opencl/test/unit_test/fixtures/memory_management_fixture.h"
#include "opencl/test/unit_test/fixtures/platform_fixture.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_buffer.h"
#include "opencl/test/unit_test/mocks/mock_command_queue.h"
#include "opencl/test/unit_test/mocks/mock_context.h"

3
opencl/test/unit_test/helpers/kernel_binary_helper.cpp
View File

@ -7,8 +7,9 @@
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include <string>

4
opencl/test/unit_test/libult/CMakeLists.txt
View File

@ -32,8 +32,6 @@ set(IGDRCL_SRCS_LIB_ULT
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/debug_helpers.cpp
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/execution_environment_helper.cpp
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/execution_environment_helper.h
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/test_files.cpp
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/test_files.h
${NEO_SOURCE_DIR}/opencl/test/unit_test/libult/create_tbx_sockets.cpp
${NEO_SOURCE_DIR}/opencl/test/unit_test/libult/debug_manager.cpp
${NEO_SOURCE_DIR}/opencl/test/unit_test/libult/io_functions.cpp
@ -47,6 +45,8 @@ set(IGDRCL_SRCS_LIB_ULT
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/memory_leak_listener.h
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/memory_management.cpp
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/memory_management.h
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/test_files.cpp
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/test_files.h
)
get_property(NEO_SHARED_TESTS_CMD_PARSE GLOBAL PROPERTY NEO_SHARED_TESTS_CMD_PARSE)

2
opencl/test/unit_test/main.cpp
View File

@ -13,6 +13,7 @@
#include "shared/source/utilities/debug_settings_reader.h"
#include "shared/test/unit_test/helpers/default_hw_info.inl"
#include "shared/test/unit_test/helpers/memory_leak_listener.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/helpers/ult_hw_config.inl"
#include "shared/test/unit_test/tests_configuration.h"
@ -20,7 +21,6 @@
#include "opencl/test/unit_test/custom_event_listener.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/kernel_binary_helper.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_gmm.h"
#include "opencl/test/unit_test/mocks/mock_program.h"
#include "opencl/test/unit_test/mocks/mock_sip.h"

2
opencl/test/unit_test/mocks/mock_compilers.cpp
View File

@ -9,9 +9,9 @@
#include "shared/source/helpers/file_io.h"
#include "shared/source/helpers/hw_info.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/os_interface/os_inc_base.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"
#include "opencl/test/unit_test/mocks/mock_sip.h"

2
opencl/test/unit_test/mocks/mock_sip.cpp
View File

@ -9,10 +9,10 @@
#include "shared/source/helpers/file_io.h"
#include "shared/source/helpers/hw_info.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/memory_manager/os_agnostic_memory_manager.h"
#include "opencl/source/os_interface/os_inc_base.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"
#include "opencl/test/unit_test/mocks/mock_program.h"

2
opencl/test/unit_test/offline_compiler/CMakeLists.txt
View File

@ -60,7 +60,7 @@ set(IGDRCL_SRCS_offline_compiler_tests
${NEO_SHARED_DIRECTORY}/helpers/file_io.cpp
${NEO_SHARED_DIRECTORY}/memory_manager/deferred_deleter.cpp
${NEO_SHARED_DIRECTORY}/memory_manager/deferred_deleter.h
${NEO_SOURCE_DIR}/opencl/test/unit_test/helpers/test_files.cpp
${NEO_SHARED_TEST_DIRECTORY}/unit_test/helpers/test_files.cpp
${IGDRCL_SRCS_cloc}
${IGDRCL_SRCS_offline_compiler_mock}
${IGDRCL_SRCS_tests_compiler_mocks}

3
opencl/test/unit_test/offline_compiler/decoder/encoder_tests.cpp
View File

@ -7,8 +7,7 @@
#include "shared/offline_compiler/source/decoder/binary_decoder.h"
#include "shared/source/helpers/array_count.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "gmock/gmock.h"
#include "mock/mock_encoder.h"

2
opencl/test/unit_test/offline_compiler/environment.h
View File

@ -8,9 +8,9 @@
#pragma once
#include "shared/offline_compiler/source/offline_compiler.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/source/os_interface/os_inc_base.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"
#include "gtest/gtest.h"

2
opencl/test/unit_test/offline_compiler/main.cpp
View File

@ -6,9 +6,9 @@
*/
#include "shared/source/os_interface/os_library.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/custom_event_listener.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "environment.h"
#include "limits.h"

2
opencl/test/unit_test/program/process_elf_binary_tests.cpp
View File

@ -12,9 +12,9 @@
#include "shared/source/helpers/file_io.h"
#include "shared/source/helpers/string.h"
#include "shared/test/unit_test/device_binary_format/patchtokens_tests.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "shared/test/unit_test/mocks/mock_device.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_cl_device.h"
#include "opencl/test/unit_test/mocks/mock_program.h"

2
opencl/test/unit_test/program/program_spec_constants_tests.cpp
View File

@ -10,10 +10,10 @@
#include "shared/source/helpers/file_io.h"
#include "shared/source/helpers/hw_info.h"
#include "shared/test/unit_test/helpers/debug_manager_state_restore.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/fixtures/device_fixture.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_cif.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"
#include "opencl/test/unit_test/mocks/mock_program.h"

2
shared/test/unit_test/compiler_interface/compiler_interface_tests.cpp
View File

@ -11,10 +11,10 @@
#include "shared/source/helpers/hw_cmds.h"
#include "shared/source/helpers/hw_info.h"
#include "shared/test/unit_test/helpers/debug_manager_state_restore.h"
#include "shared/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/fixtures/device_fixture.h"
#include "opencl/test/unit_test/global_environment.h"
#include "opencl/test/unit_test/helpers/test_files.h"
#include "opencl/test/unit_test/mocks/mock_cif.h"
#include "opencl/test/unit_test/mocks/mock_compilers.h"

0
opencl/test/unit_test/helpers/test_files.cpp → shared/test/unit_test/helpers/test_files.cpp
View File

0
opencl/test/unit_test/helpers/test_files.h → shared/test/unit_test/helpers/test_files.h
View File