From 2e9574c656cb06fef9a0fa461ded49bb02d1aa99 Mon Sep 17 00:00:00 2001 From: Krystian Chmielewski Date: Tue, 6 Sep 2022 11:17:04 +0000 Subject: [PATCH] Fix printf for type BYTE and SHORT Generated instructions writing to printf buffer require destination address to be DWORD aligned. Because of that values of type BYTE (1B) and SHORT (2B) need to be written as 4B value. This change adds support for this. When trying to read value of type BYTE or SHORT four bytes are actually read to be aligned with compiler implementation. Signed-off-by: Krystian Chmielewski --- .../test/black_box_tests/zello_printf.cpp | 55 ++++++++----------- .../unit_test/program/printf_helper_tests.cpp | 36 +++++++++--- shared/source/program/print_formatter.h | 4 +- 3 files changed, 52 insertions(+), 43 deletions(-) diff --git a/level_zero/core/test/black_box_tests/zello_printf.cpp b/level_zero/core/test/black_box_tests/zello_printf.cpp index db3c64d82d..9f8512dcc8 100644 --- a/level_zero/core/test/black_box_tests/zello_printf.cpp +++ b/level_zero/core/test/black_box_tests/zello_printf.cpp @@ -10,26 +10,20 @@ #include "zello_common.h" #include "zello_compile.h" -#include -#include #include +#include 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); +__kernel void printf_kernel(char byteValue, short shortValue, int intValue, long longValue){ + printf("byte = %hhd\nshort = %hd\nint = %d\nlong = %ld", byteValue, shortValue, intValue, longValue); } )==="; -void testPrintfKernel(ze_context_handle_t &context, ze_device_handle_t &device) { - ze_module_handle_t module; - ze_kernel_handle_t kernel; +void runPrintfKernel(ze_context_handle_t &context, ze_device_handle_t &device) { ze_command_queue_handle_t cmdQueue; ze_command_list_handle_t cmdList; - ze_group_count_t dispatchTraits; ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC}; - cmdQueueDesc.ordinal = 0; cmdQueueDesc.index = 0; cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS; @@ -44,6 +38,7 @@ void testPrintfKernel(ze_context_handle_t &context, ze_device_handle_t &device) } SUCCESS_OR_TERMINATE((0 == spirV.size())); + ze_module_handle_t module; ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC}; moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; moduleDesc.pInputModule = spirV.data(); @@ -52,41 +47,35 @@ void testPrintfKernel(ze_context_handle_t &context, ze_device_handle_t &device) SUCCESS_OR_TERMINATE(zeModuleCreate(context, device, &moduleDesc, &module, nullptr)); + ze_kernel_handle_t kernel; ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC}; - kernelDesc.pKernelName = "test_printf"; + kernelDesc.pKernelName = "printf_kernel"; SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel)); - uint32_t groupSizeX = 1; - uint32_t groupSizeY = 1; - uint32_t groupSizeZ = 1; - uint32_t globalSizeX = 64; + [[maybe_unused]] int8_t byteValue = std::numeric_limits::max(); + [[maybe_unused]] int16_t shortValue = std::numeric_limits::max(); + [[maybe_unused]] int32_t intValue = std::numeric_limits::max(); + [[maybe_unused]] int64_t longValue = std::numeric_limits::max(); - SUCCESS_OR_TERMINATE(zeKernelSuggestGroupSize(kernel, globalSizeX, 1, 1, &groupSizeX, - &groupSizeY, &groupSizeZ)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 0, sizeof(byteValue), &byteValue)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 1, sizeof(shortValue), &shortValue)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 2, sizeof(intValue), &intValue)); + SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernel, 3, sizeof(longValue), &longValue)); - 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(zeKernelSetGroupSize(kernel, 1U, 1U, 1U)); + ze_group_count_t dispatchTraits; + dispatchTraits.groupCountX = 1u; + dispatchTraits.groupCountY = 1u; + dispatchTraits.groupCountZ = 1u; 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)); } @@ -102,7 +91,7 @@ int main(int argc, char *argv[]) { SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties)); printDeviceProperties(deviceProperties); - testPrintfKernel(context, device); + runPrintfKernel(context, device); SUCCESS_OR_TERMINATE(zeContextDestroy(context)); diff --git a/opencl/test/unit_test/program/printf_helper_tests.cpp b/opencl/test/unit_test/program/printf_helper_tests.cpp index 6d4091100d..15c22e7d85 100644 --- a/opencl/test/unit_test/program/printf_helper_tests.cpp +++ b/opencl/test/unit_test/program/printf_helper_tests.cpp @@ -33,8 +33,6 @@ class PrintFormatterTest : public testing::Test { uint8_t buffer; MockGraphicsAllocation *data; - MockKernel *kernel; - std::unique_ptr program; std::unique_ptr kernelInfo; ClDevice *device; @@ -50,9 +48,6 @@ class PrintFormatterTest : public testing::Test { data = new MockGraphicsAllocation(underlyingBuffer, maxPrintfOutputLength); kernelInfo = std::make_unique(); - device = new MockClDevice{MockDevice::createWithNewExecutionEnvironment(nullptr)}; - program = std::make_unique(toClDeviceVector(*device)); - kernel = new MockKernel(program.get(), *kernelInfo, *device); printFormatter = std::unique_ptr(new PrintFormatter(static_cast(data->getUnderlyingBuffer()), printfBufferSize, is32bit, &kernelInfo->kernelDescriptor.kernelMetadata.printfStringsMap)); @@ -64,8 +59,6 @@ class PrintFormatterTest : public testing::Test { void TearDown() override { delete data; - delete kernel; - delete device; } enum class PRINTF_DATA_TYPE : int { @@ -86,6 +79,7 @@ class PrintFormatterTest : public testing::Test { VECTOR_DOUBLE }; + PRINTF_DATA_TYPE getPrintfDataType(char value) { return PRINTF_DATA_TYPE::BYTE; }; PRINTF_DATA_TYPE getPrintfDataType(int8_t value) { return PRINTF_DATA_TYPE::BYTE; }; PRINTF_DATA_TYPE getPrintfDataType(uint8_t value) { return PRINTF_DATA_TYPE::BYTE; }; PRINTF_DATA_TYPE getPrintfDataType(int16_t value) { return PRINTF_DATA_TYPE::SHORT; }; @@ -100,8 +94,14 @@ class PrintFormatterTest : public testing::Test { template void injectValue(T value) { - storeData(getPrintfDataType(value)); - storeData(value); + auto dataType = getPrintfDataType(value); + storeData(dataType); + if (dataType == PRINTF_DATA_TYPE::BYTE || + dataType == PRINTF_DATA_TYPE::SHORT) { + storeData(static_cast(value)); + } else { + storeData(value); + } } void injectStringValue(int value) { @@ -906,6 +906,24 @@ TEST_F(PrintFormatterTest, GivenNoStringMapAndBufferWithFormatStringAnd2StringsT EXPECT_STREQ(expectedOutput, output); } +TEST_F(PrintFormatterTest, GivenTypeSmallerThan4BThenItIsReadAs4BValue) { + printFormatter.reset(new PrintFormatter(static_cast(data->getUnderlyingBuffer()), printfBufferSize, true)); + const char *formatString = "%c %hd %d"; + storeData(formatString); + + char byteValue = 'a'; + injectValue(byteValue); + short shortValue = 123; + injectValue(shortValue); + int intValue = 456; + injectValue(intValue); + + const char *expectedOutput = "a 123 456"; + char output[maxPrintfOutputLength]; + printFormatter->printKernelOutput([&output](char *str) { strncpy_s(output, maxPrintfOutputLength, str, maxPrintfOutputLength - 1); }); + EXPECT_STREQ(expectedOutput, output); +} + TEST(printToSTDOUTTest, GivenStringWhenPrintingToStdoutThenOutputOccurs) { testing::internal::CaptureStdout(); printToSTDOUT("test"); diff --git a/shared/source/program/print_formatter.h b/shared/source/program/print_formatter.h index f2da586765..64947fa972 100644 --- a/shared/source/program/print_formatter.h +++ b/shared/source/program/print_formatter.h @@ -82,8 +82,10 @@ class PrintFormatter { template size_t typedPrintToken(char *output, size_t size, const char *formatString) { - T value = {0}; + T value{0}; read(&value); + constexpr auto offsetToBeDwordAligned = static_cast(std::max(int64_t(sizeof(int) - sizeof(T)), int64_t(0))); + currentOffset += offsetToBeDwordAligned; return simpleSprintf(output, size, formatString, value); }