diff --git a/level_zero/core/test/black_box_tests/CMakeLists.txt b/level_zero/core/test/black_box_tests/CMakeLists.txt index a256452297..9c2f4b960f 100644 --- a/level_zero/core/test/black_box_tests/CMakeLists.txt +++ b/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() diff --git a/level_zero/core/test/black_box_tests/zello_printf.cpp b/level_zero/core/test/black_box_tests/zello_printf.cpp new file mode 100644 index 0000000000..cee5e7b379 --- /dev/null +++ b/level_zero/core/test/black_box_tests/zello_printf.cpp @@ -0,0 +1,115 @@ +/* + * Copyright (C) 2021 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include + +#include "zello_common.h" +#include "zello_compile.h" + +#include +#include +#include + +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::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; +} diff --git a/level_zero/core/test/unit_tests/sources/module/test_module.cpp b/level_zero/core/test/unit_tests/sources/module/test_module.cpp index 31c5bb4140..c0e1911a59 100644 --- a/level_zero/core/test/unit_tests/sources/module/test_module.cpp +++ b/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; + +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(src.get()); + moduleDesc.inputSize = size; + + auto module = std::unique_ptr(Module::create(device, &moduleDesc, nullptr, ModuleType::User)); + + auto kernel = std::make_unique<::testing::NiceMock>>(); + 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;