Add printf L0 blackbox test

- add printf allocation test


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

level_zero/core/test/black_box_tests/CMakeLists.txt

@@ -20,6 +20,7 @@ if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug")
       zello_world_global_work_offset
       zello_scratch
       zello_fence
+      zello_printf
   )
 
   include_directories(common)
@@ -55,6 +56,7 @@ if("${CMAKE_BUILD_TYPE}" STREQUAL "Debug")
   target_link_libraries(zello_world_jitc_ocloc PUBLIC ocloc_lib)
   target_link_libraries(zello_scratch PUBLIC ocloc_lib)
   target_link_libraries(zello_fence PUBLIC ocloc_lib)
+  target_link_libraries(zello_printf PUBLIC ocloc_lib)
   if(UNIX)
     target_link_libraries(zello_world_global_work_offset PUBLIC ocloc_lib)
   endif()

level_zero/core/test/black_box_tests/zello_printf.cpp

@@ -0,0 +1,115 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include <level_zero/ze_api.h>
+
+#include "zello_common.h"
+#include "zello_compile.h"
+
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+
+extern bool verbose;
+bool verbose = false;
+
+const char *source = R"===(
+__kernel void test_printf(__global char *dst, __global char *src){
+    uint gid = get_global_id(0);
+    printf("global_id = %d\n", gid);
+}
+)===";
+
+void testPrintfKernel(ze_context_handle_t context, ze_device_handle_t &device) {
+    ze_module_handle_t module;
+    ze_kernel_handle_t kernel;
+    ze_command_queue_handle_t cmdQueue;
+    ze_command_list_handle_t cmdList;
+    ze_group_count_t dispatchTraits;
+
+    ze_command_queue_desc_t cmdQueueDesc = {};
+
+    cmdQueueDesc.ordinal = 0;
+    cmdQueueDesc.index = 0;
+    cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
+
+    SUCCESS_OR_TERMINATE(zeCommandQueueCreate(context, device, &cmdQueueDesc, &cmdQueue));
+    SUCCESS_OR_TERMINATE(createCommandList(context, device, cmdList));
+
+    std::string buildLog;
+    auto spirV = compileToSpirV(source, "", buildLog);
+    if (buildLog.size() > 0) {
+        std::cout << "Build log " << buildLog;
+    }
+    SUCCESS_OR_TERMINATE((0 == spirV.size()));
+
+    ze_module_desc_t moduleDesc = {};
+    moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
+    moduleDesc.pInputModule = spirV.data();
+    moduleDesc.inputSize = spirV.size();
+    moduleDesc.pBuildFlags = "";
+
+    SUCCESS_OR_TERMINATE(zeModuleCreate(context, device, &moduleDesc, &module, nullptr));
+
+    ze_kernel_desc_t kernelDesc = {};
+    kernelDesc.pKernelName = "test_printf";
+    SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel));
+
+    uint32_t groupSizeX = 1;
+    uint32_t groupSizeY = 1;
+    uint32_t groupSizeZ = 1;
+    uint32_t globalSizeX = 64;
+
+    SUCCESS_OR_TERMINATE(zeKernelSuggestGroupSize(kernel, globalSizeX, 1, 1, &groupSizeX,
+                                                  &groupSizeY, &groupSizeZ));
+
+    SUCCESS_OR_TERMINATE(zeKernelSetGroupSize(kernel, groupSizeX, groupSizeY, groupSizeZ));
+
+    dispatchTraits.groupCountX = globalSizeX / groupSizeX;
+    dispatchTraits.groupCountY = 1;
+    dispatchTraits.groupCountZ = 1;
+
+    if (verbose) {
+        std::cout << "Number of groups : (" << dispatchTraits.groupCountX << ", "
+                  << dispatchTraits.groupCountY << ", " << dispatchTraits.groupCountZ << ")"
+                  << std::endl;
+    }
+
+    SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 0, sizeof(size_t), nullptr));
+    SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 1, sizeof(size_t), nullptr));
+
+    SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatchTraits, 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()));
+
+    SUCCESS_OR_TERMINATE(zeKernelDestroy(kernel));
+    SUCCESS_OR_TERMINATE(zeModuleDestroy(module));
+
+    SUCCESS_OR_TERMINATE(zeCommandListDestroy(cmdList));
+    SUCCESS_OR_TERMINATE(zeCommandQueueDestroy(cmdQueue));
+}
+
+int main(int argc, char *argv[]) {
+    verbose = isVerbose(argc, argv);
+
+    ze_context_handle_t context = nullptr;
+    auto devices = zelloInitContextAndGetDevices(context);
+    auto device = devices[0];
+
+    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";
+
+    testPrintfKernel(context, device);
+    // always pass - no printf capturing
+    std::cout << "\nZello Printf Always PASS " << std::endl;
+
+    return 0;
+}

level_zero/core/test/unit_tests/sources/module/test_module.cpp

@@ -896,6 +896,40 @@ HWTEST_F(ModuleTranslationUnitTest, WhenBuildOptionsAreNullThenReuseExistingOpti
     EXPECT_NE(pMockCompilerInterface->inputInternalOptions.find("cl-intel-greater-than-4GB-buffer-required"), std::string::npos);
 }
 
+using PrintfModuleTest = Test<DeviceFixture>;
+
+HWTEST_F(PrintfModuleTest, GivenModuleWithPrintfWhenKernelIsCreatedThenPrintfAllocationIsPlacedInResidencyContainer) {
+    std::string testFile;
+    retrieveBinaryKernelFilenameNoRevision(testFile, "test_kernel_", ".gen");
+
+    size_t size = 0;
+    auto src = loadDataFromFile(testFile.c_str(), size);
+
+    ASSERT_NE(0u, size);
+    ASSERT_NE(nullptr, src);
+
+    ze_module_desc_t moduleDesc = {};
+    moduleDesc.format = ZE_MODULE_FORMAT_NATIVE;
+    moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(src.get());
+    moduleDesc.inputSize = size;
+
+    auto module = std::unique_ptr<L0::Module>(Module::create(device, &moduleDesc, nullptr, ModuleType::User));
+
+    auto kernel = std::make_unique<::testing::NiceMock<Mock<Kernel>>>();
+    ASSERT_NE(nullptr, kernel);
+
+    kernel->module = module.get();
+    ze_kernel_desc_t kernelDesc = {};
+    kernelDesc.pKernelName = "test";
+    kernel->initialize(&kernelDesc);
+
+    auto &container = kernel->residencyContainer;
+    auto printfPos = std::find(container.begin(), container.end(), kernel->getPrintfBufferAllocation());
+    EXPECT_NE(container.end(), printfPos);
+    bool correctPos = printfPos >= container.begin() + kernel->getImmutableData()->getDescriptor().payloadMappings.explicitArgs.size();
+    EXPECT_TRUE(correctPos);
+}
+
 TEST(BuildOptions, givenNoSrcOptionNameInSrcNamesWhenMovingBuildOptionsThenFalseIsReturned) {
     std::string srcNames = NEO::CompilerOptions::concatenate(NEO::CompilerOptions::fastRelaxedMath, NEO::CompilerOptions::finiteMathOnly);
     std::string dstNames;