diff --git a/projects/hip-tests/catch/unit/printf/CMakeLists.txt b/projects/hip-tests/catch/unit/printf/CMakeLists.txt index a127092a3c..a1466664f0 100644 --- a/projects/hip-tests/catch/unit/printf/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/printf/CMakeLists.txt @@ -5,8 +5,7 @@ set(TEST_SRC printfSpecifiers.cc printfFlagsNonHost.cc printfSpecifiersNonHost.cc - printfHost.cc -) + printfHost.cc) if(HIP_PLATFORM MATCHES "nvidia") set(LINKER_LIBS nvrtc) @@ -16,7 +15,13 @@ endif() if(UNIX) set(AMD_TEST_SRC - printfNonHost.cc) + printfNonHost.cc + hipPrintfManyDevices.cc + hipPrintfStar.cc + hipPrintfManyWaves.cc + hipPrintfWidthPrecision.cc + hipPrintfBasic.cc + hipPrintfAltForms.cc) endif() if(HIP_PLATFORM MATCHES "amd") diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfAltForms.cc b/projects/hip-tests/catch/unit/printf/hipPrintfAltForms.cc new file mode 100644 index 0000000000..82ef2377b0 --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfAltForms.cc @@ -0,0 +1,93 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include "printf_common.h" // NOLINT + +__global__ void test_kernel() { + printf("%#o\n", 042); + printf("%#x\n", 0x42); + printf("%#X\n", 0x42); + printf("%#08x\n", 0x42); + printf("%#f\n", -123.456); + printf("%#F\n", 123.456); + printf("%#e\n", 123.456); + printf("%#E\n", -123.456); + printf("%#g\n", -123.456); + printf("%#G\n", 123.456); + printf("%#a\n", 123.456); + printf("%#A\n", -123.456); + printf("%#.8x\n", 0x42); + printf("%#16.8x\n", 0x42); + printf("%-#16.8x\n", 0x42); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify alternate forms of printf API. +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfAltForms.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_PrintfAltFormsTsts") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + std::string reference(R"here(042 +0x42 +0X42 +0x000042 +-123.456000 +123.456000 +1.234560e+02 +-1.234560E+02 +-123.456 +123.456 +0x1.edd2f1a9fbe77p+6 +-0X1.EDD2F1A9FBE77P+6 +0x00000042 + 0x00000042 +0x00000042 +)here"); + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + std::string device_output = captured.gulp(CapturedData); + REQUIRE(device_output == reference); +} +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfBasic.cc b/projects/hip-tests/catch/unit/printf/hipPrintfBasic.cc new file mode 100644 index 0000000000..ae46254abe --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfBasic.cc @@ -0,0 +1,225 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include "printf_common.h" // NOLINT + +const char *msg_short = "Carpe diem."; +const char *msg_long1 = "Lorem ipsum dolor sit amet, consectetur nullam. In mollis imperdiet nibh nec ullamcorper."; // NOLINT +const char *msg_long2 = "Curabitur nec metus sit amet augue vehicula ultrices ut id leo. Lorem ipsum dolor sit amet, consectetur adipiscing elit amet."; // NOLINT + +__global__ void kernel_uniform0(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Hello World\n"); +} + +static void test_uniform0(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_uniform0, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + REQUIRE(retval[ii] == strlen("Hello World\n")); + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 1); + REQUIRE(linecount["Hello World"] == num_threads); +} + +__global__ void kernel_uniform1(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Six times Eight is %d\n", 42); +} + +static void test_uniform1(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_uniform1, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + REQUIRE(retval[ii] == strlen("Six times Eight is 42") + 1); + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 1); + REQUIRE(linecount["Six times Eight is 42"] == num_threads); +} + +__global__ void kernel_divergent0(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Thread ID: %d\n", tid); +} + +static void test_divergent0(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_divergent0, dim3(num_blocks), + dim3(threads_per_block), 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != 10; ++ii) { + REQUIRE(retval[ii] == 13); + } + for (uint ii = 10; ii != num_threads; ++ii) { + REQUIRE(retval[ii] == 14); + } + std::vector threadIds; + for (std::string line; std::getline(CapturedData, line);) { + auto pos = line.find(':'); + REQUIRE(line.substr(0, pos) == "Thread ID"); + threadIds.push_back(std::stoul(line.substr(pos + 2))); + } + std::sort(threadIds.begin(), threadIds.end()); + REQUIRE(threadIds.size() == num_threads); + REQUIRE(threadIds.back() == num_threads - 1); +} + +__global__ void kernel_divergent1(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + if (tid % 2) { + retval[tid] = printf("Hello World\n"); + } else { + retval[tid] = -1; + } +} + +static void test_divergent1(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_divergent1, dim3(num_blocks), + dim3(threads_per_block), 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 2) { + REQUIRE(retval[ii] == strlen("Hello World\n")); + } else { + REQUIRE(retval[ii] == -1); + } + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 1); + REQUIRE(linecount["Hello World"] == num_threads / 2); +} + +__global__ void kernel_series(int *retval) { + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + result += printf("%s\n", msg_long1_dev); + result += printf("%s\n", msg_short_dev); + result += printf("%s\n", msg_long2_dev); + retval[tid] = result; +} + +static void test_series(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_series, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + REQUIRE(retval[ii] == + strlen(msg_long1) + strlen(msg_short) + + strlen(msg_long2) + 3); + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[msg_long1] == num_threads); + REQUIRE(linecount[msg_long2] == num_threads); + REQUIRE(linecount[msg_short] == num_threads); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify basic functionality of printf API. +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfBasic.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_PrintfBasicTsts") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + uint num_blocks = 1; + uint threads_per_block = 64; + uint num_threads = num_blocks * threads_per_block; + void *retval_void; + HIP_CHECK(hipHostMalloc(&retval_void, 4 * num_threads)); + auto retval = reinterpret_cast(retval_void); + test_uniform0(retval, num_blocks, threads_per_block); + test_uniform1(retval, num_blocks, threads_per_block); + test_divergent0(retval, num_blocks, threads_per_block); + test_divergent1(retval, num_blocks, threads_per_block); + test_series(retval, num_blocks, threads_per_block); +} +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfManyDevices.cc b/projects/hip-tests/catch/unit/printf/hipPrintfManyDevices.cc new file mode 100644 index 0000000000..d1dd2ecb2b --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfManyDevices.cc @@ -0,0 +1,87 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include "printf_common.h" // NOLINT + +__global__ void print_things() { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const char *msg[] = {msg_short_dev, msg_long1_dev, msg_long2_dev}; + printf("%s\n", msg[tid % 3]); + if (tid % 3 == 0) + printf("%s\n", msg_short_dev); + printf("%s\n", msg[(tid + 1) % 3]); + printf("%s\n", msg[(tid + 2) % 3]); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify printf API functionality on many devices +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfManyDevices.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_ManyDevicesTest") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + uint num_blocks = 14; + uint threads_per_block = 250; + uint threads_per_device = num_blocks * threads_per_block; + int num_devices = 0; + HIP_CHECK(hipGetDeviceCount(&num_devices)); + CaptureStream captured(stdout); + for (int i = 0; i != num_devices; ++i) { + HIP_CHECK(hipSetDevice(i)); + hipLaunchKernelGGL(print_things, dim3(num_blocks), dim3(threads_per_block), + 0, 0); + HIP_CHECK(hipDeviceSynchronize()); + } + auto CapturedData = captured.getCapturedData(); + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + uint num_threads = threads_per_device * num_devices; + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[msg_long1] == num_threads); + REQUIRE(linecount[msg_long2] == num_threads); + REQUIRE(linecount[msg_short] == + num_threads + ((threads_per_device + 2) / 3) * num_devices); +} +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfManyWaves.cc b/projects/hip-tests/catch/unit/printf/hipPrintfManyWaves.cc new file mode 100644 index 0000000000..c8821d9453 --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfManyWaves.cc @@ -0,0 +1,279 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include +#include "printf_common.h" // NOLINT + +__global__ void kernel_mixed0(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + // Three strings passed as divergent values to the same hostcall. + const char *msg; + switch (tid % 3) { + case 0: + msg = msg_short_dev; + break; + case 1: + msg = msg_long1_dev; + break; + case 2: + msg = msg_long2_dev; + break; + } + retval[tid] = printf("%s\n", msg); +} + +static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_mixed0, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + REQUIRE(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + REQUIRE(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + REQUIRE(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[msg_short] == (num_threads + 2) / 3); + REQUIRE(linecount[msg_long1] == (num_threads + 1) / 3); + REQUIRE(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed1(int *retval) { + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + // Three strings passed to divergent hostcalls. + switch (tid % 3) { + case 0: + retval[tid] = printf("%s\n", msg_short_dev); + break; + case 1: + retval[tid] = printf("%s\n", msg_long1_dev); + break; + case 2: + retval[tid] = printf("%s\n", msg_long2_dev); + break; + } +} + +static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_mixed1, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + REQUIRE(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + REQUIRE(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + REQUIRE(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[msg_short] == (num_threads + 2) / 3); + REQUIRE(linecount[msg_long1] == (num_threads + 1) / 3); + REQUIRE(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed2(int *retval) { + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + // Three different strings. All workitems print all three, but + // in different orders. + const char *msg[] = {msg_short_dev, msg_long1_dev, msg_long2_dev}; + retval[tid] = + printf("%s%s%s\n", msg[tid % 3], msg[(tid + 1) % 3], msg[(tid + 2) % 3]); +} + +static void test_mixed2(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_mixed2, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + REQUIRE(retval[ii] == + strlen(msg_short) + strlen(msg_long1) + + strlen(msg_long2) + 1); + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + std::string str1 = + std::string(msg_short) + std::string(msg_long1) + + std::string(msg_long2); + std::string str2 = + std::string(msg_long1) + std::string(msg_long2) + + std::string(msg_short); + std::string str3 = + std::string(msg_long2) + std::string(msg_short) + + std::string(msg_long1); + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[str1] == (num_threads + 2) / 3); + REQUIRE(linecount[str2] == (num_threads + 1) / 3); + REQUIRE(linecount[str3] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed3(int *retval) { + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + result += printf("%s\n", msg_long1_dev); + if (tid % 3 == 0) { + result += printf("%s\n", msg_short_dev); + } + result += printf("%s\n", msg_long2_dev); + retval[tid] = result; +} + +static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + hipLaunchKernelGGL(kernel_mixed3, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 3 == 0) { + REQUIRE(retval[ii] == + strlen(msg_long1) + strlen(msg_short) + + strlen(msg_long2) + 3); + } else { + REQUIRE(retval[ii] == strlen(msg_long1) + + strlen(msg_long2) + 2); + } + } + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + REQUIRE(linecount.size() == 3); + REQUIRE(linecount[msg_long1] == num_threads); + REQUIRE(linecount[msg_long2] == num_threads); + REQUIRE(linecount[msg_short] == (num_threads + 2) / 3); +} + +__global__ void kernel_numbers() { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + for (uint i = 0; i != 7; ++i) { + uint base = tid * 21 + i * 3; + printf("%d %d %d\n", base, base + 1, base + 2); + } +} + +static void test_numbers(uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + hipLaunchKernelGGL(kernel_numbers, dim3(num_blocks), dim3(threads_per_block), + 0, 0); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + std::vector points; + while (true) { + uint i; + CapturedData >> i; + if (CapturedData.fail()) + break; + points.push_back(i); + } + std::sort(points.begin(), points.end()); + points.erase(std::unique(points.begin(), points.end()), points.end()); + REQUIRE(points.size() == 21 * num_threads); + REQUIRE(points.back() == 21 * num_threads - 1); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify printf API functionality with different strings +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfManyWaves.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_PrintfManyWaves") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + uint num_blocks = 150; + uint threads_per_block = 250; + uint num_threads = num_blocks * threads_per_block; + void *retval_void; + HIP_CHECK(hipHostMalloc(&retval_void, 4 * num_threads)); + auto retval = reinterpret_cast(retval_void); + test_mixed0(retval, num_blocks, threads_per_block); + test_mixed1(retval, num_blocks, threads_per_block); + test_mixed2(retval, num_blocks, threads_per_block); + test_mixed3(retval, num_blocks, threads_per_block); + test_numbers(num_blocks, threads_per_block); +} +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfStar.cc b/projects/hip-tests/catch/unit/printf/hipPrintfStar.cc new file mode 100644 index 0000000000..ea8299511e --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfStar.cc @@ -0,0 +1,72 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include "printf_common.h" // NOLINT + +__global__ void test_kernel_star() { + printf("%*d\n", 16, 42); + printf("%.*d\n", 8, 42); + printf("%*.*d\n", -16, 8, 42); + printf("%*.*f %s * %.*s\n", 16, 8, 123.456, "hello", 5, "worldxyz"); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify the additional arguments (*) in the printf API +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfStar.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_PrintfStar") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + std::string reference(R"here( 42 +00000042 +00000042 + 123.45600000 hello * world +)here"); + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel_star, dim3(1), dim3(1), 0, 0); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + std::string device_output = captured.gulp(CapturedData); + REQUIRE(device_output == reference); +} + +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/hipPrintfWidthPrecision.cc b/projects/hip-tests/catch/unit/printf/hipPrintfWidthPrecision.cc new file mode 100644 index 0000000000..c9a6a40856 --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/hipPrintfWidthPrecision.cc @@ -0,0 +1,92 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include "printf_common.h" // NOLINT + +__global__ void test_kernel_width() { + printf("%16d\n", 42); + printf("%.8d\n", 42); + printf("%16.5d\n", -42); + printf("%.8x\n", 0x42); + printf("%.8o\n", 042); + printf("%16.8e\n", 12345.67891); + printf("%16.8f\n", -12345.67891); + printf("%16.8g\n", 12345.67891); + printf("%8.4e\n", -12345.67891); + printf("%8.4f\n", 12345.67891); + printf("%8.4g\n", 12345.67891); + printf("%4.2f\n", 12345.67891); + printf("%.1f\n", 12345.67891); + printf("%.5s\n", "helloxyz"); +} +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ +/** +* Test Description +* ------------------------ +* - Test case to verify the floating point details via printf API +* Test source +* ------------------------ +* - catch/unit/printf/hipPrintfWidthPrecision.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.2 +*/ +TEST_CASE("Unit_Printf_PrintfWidthPrecision") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, + hipDeviceAttributeHostNativeAtomicSupported, + 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + std::string reference(R"here( 42 +00000042 + -00042 +00000042 +00000042 + 1.23456789e+04 + -12345.67891000 + 12345.679 +-1.2346e+04 +12345.6789 +1.235e+04 +12345.68 +12345.7 +hello +)here"); + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel_width, dim3(1), dim3(1), 0, 0); + HIP_CHECK(hipStreamSynchronize(0)); + auto CapturedData = captured.getCapturedData(); + std::string device_output = captured.gulp(CapturedData); + REQUIRE(device_output == reference); +} + +/** +* End doxygen group PrintfTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/printf/printf_common.h b/projects/hip-tests/catch/unit/printf/printf_common.h new file mode 100644 index 0000000000..d904d95df5 --- /dev/null +++ b/projects/hip-tests/catch/unit/printf/printf_common.h @@ -0,0 +1,106 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once +#ifdef __linux__ +#include +#include +#include +#include +#include +#include +#include +#include + +struct CaptureStream { + int saved_fd; + int orig_fd; + int temp_fd; + + char tempname[13] = "mytestXXXXXX"; + + explicit CaptureStream(FILE *original) { + orig_fd = fileno(original); + saved_fd = dup(orig_fd); + + if ((temp_fd = mkstemp(tempname)) == -1) { + error(0, errno, "Error"); + assert(false); + } + + fflush(nullptr); + if (dup2(temp_fd, orig_fd) == -1) { + error(0, errno, "Error"); + assert(false); + } + if (close(temp_fd) != 0) { + error(0, errno, "Error"); + assert(false); + } + } + + void restoreStream() { + if (saved_fd == -1) + return; + fflush(nullptr); + if (dup2(saved_fd, orig_fd) == -1) { + error(0, errno, "Error"); + assert(false); + } + if (close(saved_fd) != 0) { + error(0, errno, "Error"); + assert(false); + } + saved_fd = -1; + } + + const char *getTempFilename() { + return (const char*)tempname; + } + + std::ifstream getCapturedData() { + restoreStream(); + std::ifstream temp(tempname); + return temp; + } + + ~CaptureStream() { + restoreStream(); + if (remove(tempname) != 0) { + error(0, errno, "Error"); + assert(false); + } + } + std::string gulp(std::ifstream &input) { + std::string retval; + input.seekg(0, std::ios_base::end); + retval.resize(input.tellg()); + input.seekg(0, std::ios_base::beg); + input.read(&retval[0], retval.size()); + input.close(); + return retval; + } +}; +extern const char *msg_short; +extern const char *msg_long1; +extern const char *msg_long2; + +__device__ const char *msg_short_dev = "Carpe diem."; +__device__ const char *msg_long1_dev = "Lorem ipsum dolor sit amet, consectetur nullam. In mollis imperdiet nibh nec ullamcorper."; // NOLINT +__device__ const char *msg_long2_dev = "Curabitur nec metus sit amet augue vehicula ultrices ut id leo. Lorem ipsum dolor sit amet, consectetur adipiscing elit amet."; // NOLINT +#endif