From 05951d3a9ece969879b52de24aee3f94fab4704b Mon Sep 17 00:00:00 2001 From: Mateusz Hoppe Date: Wed, 8 Apr 2020 15:37:38 +0200 Subject: [PATCH] Add more cmdlist tests Related-To: NEO-4515 Change-Id: I0ba9f04eb9f3728bab84604d4740650cd9d0e4ae Signed-off-by: Mateusz Hoppe --- .../core/test/unit_tests/gen9/CMakeLists.txt | 1 + ...test_cmdlist_append_launch_kernel_gen9.cpp | 57 ++++++++++++ .../sources/cmdlist/test_cmdlist_api.cpp | 16 ++++ .../test_cmdlist_append_launch_kernel.cpp | 88 ++++++++++++++++++- 4 files changed, 161 insertions(+), 1 deletion(-) create mode 100644 level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp diff --git a/level_zero/core/test/unit_tests/gen9/CMakeLists.txt b/level_zero/core/test/unit_tests/gen9/CMakeLists.txt index dfcb4af4fa..c29991a6b1 100644 --- a/level_zero/core/test/unit_tests/gen9/CMakeLists.txt +++ b/level_zero/core/test/unit_tests/gen9/CMakeLists.txt @@ -8,5 +8,6 @@ if(TESTS_GEN9) target_sources(${TARGET_NAME} PRIVATE ${COMPUTE_RUNTIME_ULT_GEN9} ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt + ${CMAKE_CURRENT_SOURCE_DIR}/test_cmdlist_append_launch_kernel_gen9.cpp ) endif() diff --git a/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp b/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp new file mode 100644 index 0000000000..499b44fd3e --- /dev/null +++ b/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp @@ -0,0 +1,57 @@ +/* + * Copyright (C) 2020 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include "shared/source/gen9/reg_configs.h" +#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; +using IsSKLOrKBL = IsWithinProducts; + +HWTEST2_F(CommandListAppendLaunchKernel, givenKernelWithSLMThenL3IsProgrammedWithSLMValue, IsSKLOrKBL) { + using MI_LOAD_REGISTER_IMM = typename FamilyType::MI_LOAD_REGISTER_IMM; + createKernel(); + std::unique_ptr commandList(CommandList::create(productFamily, device)); + ze_group_count_t groupCount{1, 1, 1}; + + EXPECT_LE(0u, kernel->kernelImmData->getDescriptor().kernelAttributes.slmInlineSize); + + auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr); + EXPECT_EQ(ZE_RESULT_SUCCESS, result); + + auto usedSpaceAfter = commandList->commandContainer.getCommandStream()->getUsed(); + + GenCmdList cmdList; + ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer( + cmdList, ptrOffset(commandList->commandContainer.getCommandStream()->getCpuBase(), 0), usedSpaceAfter)); + + bool foundL3 = false; + for (auto it = cmdList.begin(); it != cmdList.end(); it++) { + auto lri = genCmdCast(*it); + if (lri) { + if (lri->getRegisterOffset() == NEO::L3CNTLRegisterOffset::registerOffset) { + auto value = lri->getDataDword(); + auto dataSlm = NEO::PreambleHelper::getL3Config(commandList->commandContainer.getDevice()->getHardwareInfo(), true); + EXPECT_EQ(dataSlm, value); + foundL3 = true; + break; + } + } + } + + EXPECT_TRUE(foundL3); +} + +} // namespace ult +} // namespace L0 \ No newline at end of file diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_api.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_api.cpp index 058cd81e12..13d211f544 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_api.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_api.cpp @@ -9,6 +9,7 @@ #include "level_zero/core/test/unit_tests/fixtures/device_fixture.h" #include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h" +#include "level_zero/core/test/unit_tests/mocks/mock_kernel.h" namespace L0 { namespace ult { @@ -61,5 +62,20 @@ TEST(zeCommandListAppendMemoryFill, whenCalledThenRedirectedToObject) { ASSERT_EQ(ZE_RESULT_SUCCESS, res); } +TEST(zeCommandListAppendLaunchKernel, whenCalledThenRedirectedToObject) { + Mock commandList; + Mock<::L0::Kernel> kernel; + ze_group_count_t dispatchFunctionArguments; + + EXPECT_CALL(commandList, appendLaunchKernel(kernel.toHandle(), &dispatchFunctionArguments, + nullptr, 0, nullptr)) + .Times(1); + + auto result = + zeCommandListAppendLaunchKernel(commandList.toHandle(), kernel.toHandle(), + &dispatchFunctionArguments, nullptr, 0, nullptr); + EXPECT_EQ(ZE_RESULT_SUCCESS, result); +} + } // namespace ult } // namespace L0 \ No newline at end of file diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel.cpp index 4eaaac366e..b24118fa79 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel.cpp @@ -5,13 +5,15 @@ * */ +#include "shared/source/command_container/command_encoder.h" +#include "shared/source/helpers/preamble.h" #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" - +#include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h" namespace L0 { namespace ult { @@ -82,5 +84,89 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenFunctionWhenBind } } +HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfUsedWhenAppendedToCommandListThenKernelIsStored) { + createKernel(); + std::unique_ptr commandList(CommandList::create(productFamily, device)); + ze_group_count_t groupCount{1, 1, 1}; + + EXPECT_TRUE(kernel->kernelImmData->getDescriptor().kernelAttributes.flags.usesPrintf); + + auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr); + EXPECT_EQ(ZE_RESULT_SUCCESS, result); + + EXPECT_EQ(1u, commandList->getPrintfFunctionContainer().size()); + EXPECT_EQ(kernel.get(), commandList->getPrintfFunctionContainer()[0]); +} + +HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfUsedWhenAppendedToCommandListMultipleTimesThenKernelIsStoredOnce) { + createKernel(); + std::unique_ptr commandList(CommandList::create(productFamily, device)); + ze_group_count_t groupCount{1, 1, 1}; + + EXPECT_TRUE(kernel->kernelImmData->getDescriptor().kernelAttributes.flags.usesPrintf); + + auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr); + EXPECT_EQ(ZE_RESULT_SUCCESS, result); + + EXPECT_EQ(1u, commandList->getPrintfFunctionContainer().size()); + EXPECT_EQ(kernel.get(), commandList->getPrintfFunctionContainer()[0]); + + result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr); + EXPECT_EQ(ZE_RESULT_SUCCESS, result); + EXPECT_EQ(1u, commandList->getPrintfFunctionContainer().size()); +} + +HWTEST_F(CommandListAppendLaunchKernel, WhenAppendingMultipleTimesThenSshIsNotDepletedButReallocated) { + createKernel(); + std::unique_ptr commandList(CommandList::create(productFamily, device)); + ze_group_count_t groupCount{1, 1, 1}; + + auto kernelSshSize = kernel->getSurfaceStateHeapDataSize(); + auto ssh = commandList->commandContainer.getIndirectHeap(NEO::HeapType::SURFACE_STATE); + auto sshHeapSize = ssh->getMaxAvailableSpace(); + auto initialAllocation = ssh->getGraphicsAllocation(); + EXPECT_NE(nullptr, initialAllocation); + + for (size_t i = 0; i < sshHeapSize / kernelSshSize + 1; i++) { + auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr); + ASSERT_EQ(ZE_RESULT_SUCCESS, result); + } + + auto reallocatedAllocation = ssh->getGraphicsAllocation(); + EXPECT_NE(nullptr, reallocatedAllocation); + EXPECT_NE(initialAllocation, reallocatedAllocation); +} + +using SklPlusMatcher = IsAtLeastProduct; +HWTEST2_F(CommandListAppendLaunchKernel, WhenAppendingFunctionThenUsedCmdBufferSizeDoesNotExceedEstimate, SklPlusMatcher) { + createKernel(); + ze_group_count_t groupCount{1, 1, 1}; + + auto commandList = std::make_unique>>(); + bool ret = commandList->initialize(device); + ASSERT_TRUE(ret); + + auto sizeBefore = commandList->commandContainer.getCommandStream()->getUsed(); + + auto result = commandList->appendLaunchKernelWithParams(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false, false); + ASSERT_EQ(ZE_RESULT_SUCCESS, result); + + auto sizeAfter = commandList->commandContainer.getCommandStream()->getUsed(); + auto estimate = NEO::EncodeDispatchKernel::estimateEncodeDispatchKernelCmdsSize(device->getNEODevice()); + + EXPECT_LE(sizeAfter - sizeBefore, estimate); + + sizeBefore = commandList->commandContainer.getCommandStream()->getUsed(); + + result = commandList->appendLaunchKernelWithParams(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, true, false); + ASSERT_EQ(ZE_RESULT_SUCCESS, result); + + sizeAfter = commandList->commandContainer.getCommandStream()->getUsed(); + estimate = NEO::EncodeDispatchKernel::estimateEncodeDispatchKernelCmdsSize(device->getNEODevice()); + + EXPECT_LE(sizeAfter - sizeBefore, estimate); + EXPECT_LE(sizeAfter - sizeBefore, estimate); +} + } // namespace ult } // namespace L0 \ No newline at end of file