diff --git a/tests/src/printf/hipPrintfAltForms.cpp b/tests/src/printf/hipPrintfAltForms.cpp new file mode 100644 index 0000000000..dc0b325bba --- /dev/null +++ b/tests/src/printf/hipPrintfAltForms.cpp @@ -0,0 +1,76 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__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); +} + +int main(int argc, char **argv) { + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfBasic.cpp b/tests/src/printf/hipPrintfBasic.cpp new file mode 100644 index 0000000000..e51373c251 --- /dev/null +++ b/tests/src/printf/hipPrintfBasic.cpp @@ -0,0 +1,238 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" +#include + +// Global string constants don't work inside device functions, so we +// use a macro to repeat the declaration in host and device contexts. +DECLARE_DATA(); + +__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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == strlen("Hello World\n")); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == strlen("Six times Eight is 42") + 1); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != 10; ++ii) { + HIPASSERT(retval[ii] == 13); + } + + for (uint ii = 10; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == 14); + } + + std::vector threadIds; + for (std::string line; std::getline(CapturedData, line);) { + auto pos = line.find(':'); + HIPASSERT(line.substr(0, pos) == "Thread ID"); + threadIds.push_back(std::stoul(line.substr(pos + 2))); + } + + std::sort(threadIds.begin(), threadIds.end()); + HIPASSERT(threadIds.size() == num_threads); + HIPASSERT(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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 2) { + HIPASSERT(retval[ii] == strlen("Hello World\n")); + } else { + HIPASSERT(retval[ii] == -1); + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(linecount["Hello World"] == num_threads / 2); +} + +__global__ void kernel_series(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + + result += printf("%s\n", msg_long1); + result += printf("%s\n", msg_short); + result += printf("%s\n", msg_long2); + + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(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]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(linecount[msg_short] == num_threads); +} + +int main() { + uint num_blocks = 1; + uint threads_per_block = 64; + uint num_threads = num_blocks * threads_per_block; + + void *retval_void; + HIPCHECK(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); + + passed(); +} diff --git a/tests/src/printf/hipPrintfFlags.cpp b/tests/src/printf/hipPrintfFlags.cpp new file mode 100644 index 0000000000..5093c5e1c9 --- /dev/null +++ b/tests/src/printf/hipPrintfFlags.cpp @@ -0,0 +1,68 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + printf("%08d\n", 42); + printf("%08i\n", -42); + printf("%08u\n", 42); + printf("%08g\n", 123.456); + printf("%0+8d\n", 42); + printf("%+d\n", -42); + printf("%+08d\n", 42); + printf("%-8s\n", "xyzzy"); + printf("% i\n", -42); + printf("%-16.8d\n", 42); + printf("%16.8d\n", 42); +} + +int main(int argc, char **argv) { + std::string reference(R"here(00000042 +-0000042 +00000042 +0123.456 ++0000042 +-42 ++0000042 +xyzzy +-42 +00000042 + 00000042 +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfManyDevices.cpp b/tests/src/printf/hipPrintfManyDevices.cpp new file mode 100644 index 0000000000..23751e26ea --- /dev/null +++ b/tests/src/printf/hipPrintfManyDevices.cpp @@ -0,0 +1,77 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +DECLARE_DATA(); + +__global__ void print_things() { + DECLARE_DATA(); + + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const char *msg[] = {msg_short, msg_long1, msg_long2}; + + printf("%s\n", msg[tid % 3]); + if (tid % 3 == 0) + printf("%s\n", msg_short); + printf("%s\n", msg[(tid + 1) % 3]); + printf("%s\n", msg[(tid + 2) % 3]); +} + +int main() { + uint num_blocks = 14; + uint threads_per_block = 250; + uint threads_per_device = num_blocks * threads_per_block; + + int num_devices = 0; + hipGetDeviceCount(&num_devices); + + CaptureStream captured(stdout); + for (int i = 0; i != num_devices; ++i) { + hipSetDevice(i); + hipLaunchKernelGGL(print_things, dim3(num_blocks), dim3(threads_per_block), + 0, 0); + 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; + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(linecount[msg_short] == + num_threads + ((threads_per_device + 2) / 3) * num_devices); + + passed(); +} diff --git a/tests/src/printf/hipPrintfManyWaves.cpp b/tests/src/printf/hipPrintfManyWaves.cpp new file mode 100644 index 0000000000..c15e695424 --- /dev/null +++ b/tests/src/printf/hipPrintfManyWaves.cpp @@ -0,0 +1,301 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" +#include + +// Global string constants don't work inside device functions, so we +// use a macro to repeat the declaration in host and device contexts. +DECLARE_DATA(); + +__global__ void kernel_mixed0(int *retval) { + DECLARE_DATA(); + + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + ulong result = 0; + + // Three strings passed as divergent values to the same hostcall. + const char *msg; + switch (tid % 3) { + case 0: + msg = msg_short; + break; + case 1: + msg = msg_long1; + break; + case 2: + msg = msg_long2; + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + HIPASSERT(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + HIPASSERT(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + HIPASSERT(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_short] == (num_threads + 2) / 3); + HIPASSERT(linecount[msg_long1] == (num_threads + 1) / 3); + HIPASSERT(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed1(int *retval) { + DECLARE_DATA(); + + 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); + break; + case 1: + retval[tid] = printf("%s\n", msg_long1); + break; + case 2: + retval[tid] = printf("%s\n", msg_long2); + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + HIPASSERT(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + HIPASSERT(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + HIPASSERT(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_short] == (num_threads + 2) / 3); + HIPASSERT(linecount[msg_long1] == (num_threads + 1) / 3); + HIPASSERT(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed2(int *retval) { + DECLARE_DATA(); + + 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, msg_long1, msg_long2}; + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(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); + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[str1] == (num_threads + 2) / 3); + HIPASSERT(linecount[str2] == (num_threads + 1) / 3); + HIPASSERT(linecount[str3] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed3(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + + result += printf("%s\n", msg_long1); + if (tid % 3 == 0) { + result += printf("%s\n", msg_short); + } + result += printf("%s\n", msg_long2); + + 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); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 3 == 0) { + HIPASSERT(retval[ii] == + strlen(msg_long1) + strlen(msg_short) + strlen(msg_long2) + 3); + } else { + HIPASSERT(retval[ii] == strlen(msg_long1) + strlen(msg_long2) + 2); + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(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); + 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()); + HIPASSERT(points.size() == 21 * num_threads); + HIPASSERT(points.back() == 21 * num_threads - 1); + + passed(); +} + +int main(int argc, char **argv) { + uint num_blocks = 150; + uint threads_per_block = 250; + uint num_threads = num_blocks * threads_per_block; + + void *retval_void; + HIPCHECK(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); + + passed(); +} diff --git a/tests/src/printf/hipPrintfSpecifiers.cpp b/tests/src/printf/hipPrintfSpecifiers.cpp new file mode 100644 index 0000000000..03507658b9 --- /dev/null +++ b/tests/src/printf/hipPrintfSpecifiers.cpp @@ -0,0 +1,90 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + const char *N = nullptr; + const char *s = "hello world"; + + printf("xyzzy\n"); + printf("%%\n"); + printf("hello %% world\n"); + printf("%%s\n"); + // Two special tests to make sure that the compiler pass correctly + // skips over a '%%' without affecting the logic for locating + // string arguments. + printf("%%s%p\n", (void *)0xf01dab1eca55e77e); + printf("%%c%s\n", "xyzzy"); + printf("%c%c%c\n", 's', 'e', 'p'); + printf("%d\n", -42); + printf("%u\n", 42); + 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("%c\n", 'x'); + printf("%s\n", N); + printf("%p\n", N); + printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 5, s, (void *)0xf01dab1eca55e77e); +} + +int main(int argc, char **argv) { + std::string reference(R"here(xyzzy +% +hello % world +%s +%s0xf01dab1eca55e77e +%cxyzzy +sep +-42 +42 +123.456000 +-123.456000 +-1.234560e+02 +1.234560E+02 +123.456 +-123.456 +x + +(nil) +3.14159000 hello 0xf01dab1eca55e77e +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfStar.cpp b/tests/src/printf/hipPrintfStar.cpp new file mode 100644 index 0000000000..5e97d6eae5 --- /dev/null +++ b/tests/src/printf/hipPrintfStar.cpp @@ -0,0 +1,54 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + 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"); +} + +int main(int argc, char **argv) { + std::string reference(R"here( 42 +00000042 +00000042 + 123.45600000 hello * world +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfWidthPrecision.cpp b/tests/src/printf/hipPrintfWidthPrecision.cpp new file mode 100644 index 0000000000..db003377a9 --- /dev/null +++ b/tests/src/printf/hipPrintfWidthPrecision.cpp @@ -0,0 +1,74 @@ +/* +Copyright (c) 2020 - present 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + 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"); +} + +int main(int argc, char **argv) { + 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, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/printf_common.h b/tests/src/printf/printf_common.h new file mode 100644 index 0000000000..a2df88db9f --- /dev/null +++ b/tests/src/printf/printf_common.h @@ -0,0 +1,94 @@ +#ifndef COMMON_H +#define COMMON_H + +#include +#include +#include +#include +#include +#include +#include +#include + +struct CaptureStream { + int saved_fd; + int orig_fd; + int temp_fd; + + char tempname[13] = "mytestXXXXXX"; + + CaptureStream(FILE *original) { + orig_fd = fileno(original); + saved_fd = dup(orig_fd); + + temp_fd = mkstemp(tempname); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + + fflush(nullptr); + dup2(temp_fd, orig_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + close(temp_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + } + + void restoreStream() { + if (saved_fd == -1) + return; + fflush(nullptr); + dup2(saved_fd, orig_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + close(saved_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + saved_fd = -1; + } + + std::ifstream getCapturedData() { + restoreStream(); + std::ifstream temp(tempname); + return temp; + } + + ~CaptureStream() { + restoreStream(); + remove(tempname); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + } +}; + +static 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; +} + +#define DECLARE_DATA() \ + const char *msg_short = "Carpe diem."; \ + const char *msg_long1 = "Lorem ipsum dolor sit amet, consectetur nullam. " \ + "In mollis imperdiet nibh nec ullamcorper."; \ + const char *msg_long2 = "Curabitur nec metus sit amet augue vehicula " \ + "ultrices ut id leo. Lorem ipsum dolor sit amet, " \ + "consectetur adipiscing elit amet."; + +#endif