Add more cmdlist tests

Related-To: NEO-4515

Change-Id: I0ba9f04eb9f3728bab84604d4740650cd9d0e4ae
Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>

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

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<ModuleFixture>;
+using IsSKLOrKBL = IsWithinProducts<IGFX_SKYLAKE, IGFX_KABYLAKE>;
+
+HWTEST2_F(CommandListAppendLaunchKernel, givenKernelWithSLMThenL3IsProgrammedWithSLMValue, IsSKLOrKBL) {
+    using MI_LOAD_REGISTER_IMM = typename FamilyType::MI_LOAD_REGISTER_IMM;
+    createKernel();
+    std::unique_ptr<L0::CommandList> 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<MI_LOAD_REGISTER_IMM *>(*it);
+        if (lri) {
+            if (lri->getRegisterOffset() == NEO::L3CNTLRegisterOffset<FamilyType>::registerOffset) {
+                auto value = lri->getDataDword();
+                auto dataSlm = NEO::PreambleHelper<FamilyType>::getL3Config(commandList->commandContainer.getDevice()->getHardwareInfo(), true);
+                EXPECT_EQ(dataSlm, value);
+                foundL3 = true;
+                break;
+            }
+        }
+    }
+
+    EXPECT_TRUE(foundL3);
+}
+
+} // namespace ult
+} // namespace L0

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

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<L0::CommandList> 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<L0::CommandList> 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<L0::CommandList> 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<IGFX_SKYLAKE>;
+HWTEST2_F(CommandListAppendLaunchKernel, WhenAppendingFunctionThenUsedCmdBufferSizeDoesNotExceedEstimate, SklPlusMatcher) {
+    createKernel();
+    ze_group_count_t groupCount{1, 1, 1};
+
+    auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
+    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<FamilyType>::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<FamilyType>::estimateEncodeDispatchKernelCmdsSize(device->getNEODevice());
+
+    EXPECT_LE(sizeAfter - sizeBefore, estimate);
+    EXPECT_LE(sizeAfter - sizeBefore, estimate);
+}
+
 } // namespace ult
 } // namespace L0