From 223dddae6dc9ffb961270917333ada02e7275a95 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Sun, 7 Mar 2021 21:04:23 -0500 Subject: [PATCH] SWDEV-232428 - Fix HIP printf tests on Windows MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit On Windows there's something fundamentally broken about redirecting IO into a file and then restoring that said IO to it's original state. Even though no syscalls would fail, the output would sometimes either go into CLI or straight up nowhere. Simply using pipes instead of a temporary file magically resolves the above issue ¯\_(ツ)_/¯ Unfortunately the max pipe size on Linux is 1Mb, which is not enough to store all the data printed by the kernel. This leads to a softhang in vprintf(). Stick to using a temporary file on Linux, but switch to pipes on Windows. Slightly refactor the CaptureStream struct to accomadate this difference. Change-Id: Id8e68f150df47815a4f652ee2bcd6cfb7c3e3bac --- tests/src/printf/hipPrintfAltForms.cpp | 9 +- tests/src/printf/hipPrintfBasic.cpp | 68 +++++++++---- tests/src/printf/hipPrintfFlags.cpp | 9 +- tests/src/printf/hipPrintfManyDevices.cpp | 13 ++- tests/src/printf/hipPrintfManyWaves.cpp | 59 +++++++---- tests/src/printf/hipPrintfSpecifiers.cpp | 9 +- tests/src/printf/hipPrintfStar.cpp | 9 +- tests/src/printf/hipPrintfWidthPrecision.cpp | 9 +- tests/src/printf/printf_common.h | 102 +++++++++++++++---- tests/src/test_common.h | 1 + 10 files changed, 211 insertions(+), 77 deletions(-) diff --git a/tests/src/printf/hipPrintfAltForms.cpp b/tests/src/printf/hipPrintfAltForms.cpp index f2cd0e6233..c6df9b0d36 100644 --- a/tests/src/printf/hipPrintfAltForms.cpp +++ b/tests/src/printf/hipPrintfAltForms.cpp @@ -65,11 +65,14 @@ int main(int argc, char **argv) { 0x00000042 )here"); - CaptureStream captured(stdout); + CaptureStream capture(stdout); + + capture.Begin(); hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); - std::string device_output = gulp(CapturedData); + capture.End(); + + std::string device_output = capture.getData(); HIPASSERT(device_output == reference); passed(); diff --git a/tests/src/printf/hipPrintfBasic.cpp b/tests/src/printf/hipPrintfBasic.cpp index 828d299e61..27e7ec8a69 100644 --- a/tests/src/printf/hipPrintfBasic.cpp +++ b/tests/src/printf/hipPrintfBasic.cpp @@ -41,24 +41,29 @@ __global__ void kernel_uniform0(int *retval) { static void test_uniform0(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_uniform0, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { HIPASSERT(retval[ii] == strlen("Hello World\n")); } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -73,24 +78,29 @@ __global__ void kernel_uniform1(int *retval) { static void test_uniform1(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_uniform1, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { HIPASSERT(retval[ii] == strlen("Six times Eight is 42") + 1); } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -105,17 +115,18 @@ __global__ void kernel_divergent0(int *retval) { static void test_divergent0(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_divergent0, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != 10; ++ii) { HIPASSERT(retval[ii] == 13); @@ -125,8 +136,12 @@ static void test_divergent0(int *retval, uint num_blocks, HIPASSERT(retval[ii] == 14); } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::vector threadIds; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { auto pos = line.find(':'); HIPASSERT(line.substr(0, pos) == "Thread ID"); threadIds.push_back(std::stoul(line.substr(pos + 2))); @@ -148,17 +163,18 @@ __global__ void kernel_divergent1(int *retval) { static void test_divergent1(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_divergent1, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { if (ii % 2) { @@ -168,8 +184,12 @@ static void test_divergent1(int *retval, uint num_blocks, } } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -191,25 +211,30 @@ __global__ void kernel_series(int *retval) { } static void test_series(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_series, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { HIPASSERT(retval[ii] == strlen(msg_long1) + strlen(msg_short) + strlen(msg_long2) + 3); } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -231,20 +256,25 @@ __global__ void kernel_divergent_loop() { } static void test_divergent_loop(uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; + capture.Begin(); hipLaunchKernelGGL(kernel_divergent_loop, dim3(num_blocks), dim3(threads_per_block), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); + + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; std::map count; while (true) { int i; - CapturedData >> i; - if (CapturedData.fail()) + dataStream >> i; + if (dataStream.fail()) break; count[i]++; } diff --git a/tests/src/printf/hipPrintfFlags.cpp b/tests/src/printf/hipPrintfFlags.cpp index 5eeb1a34c6..33a7f1ae4d 100644 --- a/tests/src/printf/hipPrintfFlags.cpp +++ b/tests/src/printf/hipPrintfFlags.cpp @@ -57,11 +57,14 @@ xyzzy 00000042 )here"); - CaptureStream captured(stdout); + CaptureStream capture(stdout); + + capture.Begin(); hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); - std::string device_output = gulp(CapturedData); + capture.End(); + + std::string device_output = capture.getData(); HIPASSERT(device_output == reference); passed(); diff --git a/tests/src/printf/hipPrintfManyDevices.cpp b/tests/src/printf/hipPrintfManyDevices.cpp index da3edcf89b..cec1fff694 100644 --- a/tests/src/printf/hipPrintfManyDevices.cpp +++ b/tests/src/printf/hipPrintfManyDevices.cpp @@ -49,20 +49,25 @@ int main() { uint threads_per_block = 250; uint threads_per_device = num_blocks * threads_per_block; + CaptureStream capture(stdout); + int num_devices = 0; hipGetDeviceCount(&num_devices); - - CaptureStream captured(stdout); + capture.Begin(); 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(); + capture.End(); + + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } diff --git a/tests/src/printf/hipPrintfManyWaves.cpp b/tests/src/printf/hipPrintfManyWaves.cpp index ca0b336a6e..1a42739e40 100644 --- a/tests/src/printf/hipPrintfManyWaves.cpp +++ b/tests/src/printf/hipPrintfManyWaves.cpp @@ -58,17 +58,18 @@ __global__ void kernel_mixed0(int *retval) { } static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_mixed0, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { switch (ii % 3) { @@ -84,8 +85,12 @@ static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) { } } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -115,17 +120,18 @@ __global__ void kernel_mixed1(int *retval) { } static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_mixed1, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { switch (ii % 3) { @@ -141,8 +147,12 @@ static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) { } } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -165,25 +175,30 @@ __global__ void kernel_mixed2(int *retval) { } static void test_mixed2(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_mixed2, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { HIPASSERT(retval[ii] == strlen(msg_short) + strlen(msg_long1) + strlen(msg_long2) + 1); } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -216,17 +231,18 @@ __global__ void kernel_mixed3(int *retval) { } static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; for (uint i = 0; i != num_threads; ++i) { retval[i] = 0x23232323; } + capture.Begin(); hipLaunchKernelGGL(kernel_mixed3, dim3(num_blocks), dim3(threads_per_block), 0, 0, retval); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); for (uint ii = 0; ii != num_threads; ++ii) { if (ii % 3 == 0) { @@ -237,8 +253,12 @@ static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) { } } + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; + std::map linecount; - for (std::string line; std::getline(CapturedData, line);) { + for (std::string line; std::getline(dataStream, line);) { linecount[line]++; } @@ -257,19 +277,24 @@ __global__ void kernel_numbers() { } static void test_numbers(uint num_blocks, uint threads_per_block) { - CaptureStream captured(stdout); + CaptureStream capture(stdout); uint num_threads = num_blocks * threads_per_block; + capture.Begin(); hipLaunchKernelGGL(kernel_numbers, dim3(num_blocks), dim3(threads_per_block), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); + capture.End(); + + std::string data = capture.getData(); + std::stringstream dataStream; + dataStream << data; std::vector points; while (true) { uint i; - CapturedData >> i; - if (CapturedData.fail()) + dataStream >> i; + if (dataStream.fail()) break; points.push_back(i); } @@ -278,8 +303,6 @@ static void test_numbers(uint num_blocks, uint threads_per_block) { 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) { diff --git a/tests/src/printf/hipPrintfSpecifiers.cpp b/tests/src/printf/hipPrintfSpecifiers.cpp index 2e427ea537..7a8cb6d7d8 100644 --- a/tests/src/printf/hipPrintfSpecifiers.cpp +++ b/tests/src/printf/hipPrintfSpecifiers.cpp @@ -102,11 +102,14 @@ x )here"); #endif - CaptureStream captured(stdout); + CaptureStream capture(stdout); + + capture.Begin(); hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); - std::string device_output = gulp(CapturedData); + capture.End(); + + std::string device_output = capture.getData(); HIPASSERT(device_output == reference); passed(); diff --git a/tests/src/printf/hipPrintfStar.cpp b/tests/src/printf/hipPrintfStar.cpp index 38f31f406b..dc868fb713 100644 --- a/tests/src/printf/hipPrintfStar.cpp +++ b/tests/src/printf/hipPrintfStar.cpp @@ -43,11 +43,14 @@ int main(int argc, char **argv) { 123.45600000 hello * world )here"); - CaptureStream captured(stdout); + CaptureStream capture(stdout); + + capture.Begin(); hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); - std::string device_output = gulp(CapturedData); + capture.End(); + + std::string device_output = capture.getData(); HIPASSERT(device_output == reference); passed(); diff --git a/tests/src/printf/hipPrintfWidthPrecision.cpp b/tests/src/printf/hipPrintfWidthPrecision.cpp index 1ac7b0a0c3..132e249921 100644 --- a/tests/src/printf/hipPrintfWidthPrecision.cpp +++ b/tests/src/printf/hipPrintfWidthPrecision.cpp @@ -63,11 +63,14 @@ int main(int argc, char **argv) { hello )here"); - CaptureStream captured(stdout); + CaptureStream capture(stdout); + + capture.Begin(); hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); hipStreamSynchronize(0); - auto CapturedData = captured.getCapturedData(); - std::string device_output = gulp(CapturedData); + capture.End(); + + std::string device_output = capture.getData(); HIPASSERT(device_output == reference); passed(); diff --git a/tests/src/printf/printf_common.h b/tests/src/printf/printf_common.h index 60b0aff46c..4f75b13d62 100644 --- a/tests/src/printf/printf_common.h +++ b/tests/src/printf/printf_common.h @@ -2,14 +2,84 @@ #define COMMON_H #include -#include #include #include #include #include #include -#include +#include +#include +#include +#if defined(_WIN32) +#include +#else +#include +#include +#endif + +#if defined(_WIN32) +class CaptureStream { +private: + FILE* stream; + int fdPipe[2]; + int fd; + + static constexpr size_t bufferSize = 25 * 1024 * 1024; + +public: + CaptureStream(FILE *original) { + stream = original; + + if (pipe(fdPipe, bufferSize, O_TEXT) != 0) { + fprintf(stderr, "pipe(3) failed with error %d\n", errno); + assert(false); + } + + if ((fd = dup(fileno(stream))) == -1) { + fprintf(stderr, "dup(1) failed with error %d\n", errno); + assert(false); + } + } + + ~CaptureStream() { + close(fd); + close(fdPipe[1]); + close(fdPipe[0]); + } + + void Begin() { + fflush(stream); + + if (dup2(fdPipe[1], fileno(stream)) == -1) { + fprintf(stderr, "dup2(2) failed with error %d\n", errno); + assert(false); + } + + setvbuf(stream, NULL, _IONBF, 0); + } + + void End() { + if (dup2(fd, fileno(stream)) == -1) { + fprintf(stderr, "dup2(2) failed with error %d\n", errno); + assert(false); + } + } + + std::string getData() { + std::string data; + data.resize(bufferSize); + + int numRead = read(fdPipe[0], const_cast(data.c_str()), bufferSize); + data[numRead] = '\0'; + + data.resize(strlen(data.c_str())); + data.shrink_to_fit(); + + return data; + } +}; +#else struct CaptureStream { int saved_fd; int orig_fd; @@ -25,7 +95,9 @@ struct CaptureStream { error(0, errno, "Error"); assert(false); } + } + void Begin() { fflush(nullptr); if (dup2(temp_fd, orig_fd) == -1) { error(0, errno, "Error"); @@ -37,9 +109,7 @@ struct CaptureStream { } } - void restoreStream() { - if (saved_fd == -1) - return; + void End() { fflush(nullptr); if (dup2(saved_fd, orig_fd) == -1) { error(0, errno, "Error"); @@ -49,33 +119,23 @@ struct CaptureStream { error(0, errno, "Error"); assert(false); } - saved_fd = -1; } - std::ifstream getCapturedData() { - restoreStream(); - std::ifstream temp(tempname); - return temp; + std::string getData() { + std::ifstream tmpFileStream(tempname); + std::stringstream strStream; + strStream << tmpFileStream.rdbuf(); + return strStream.str(); } ~CaptureStream() { - restoreStream(); if (remove(tempname) != 0) { 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; -} +#endif #define DECLARE_DATA() \ const char *msg_short = "Carpe diem."; \ diff --git a/tests/src/test_common.h b/tests/src/test_common.h index 080d5e144a..e7612b079a 100755 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -137,6 +137,7 @@ inline int hip_skip_retcode() { #define dup(x) _dup(x) #define dup2(x,y) _dup2(x,y) #define close(x) _close(x) +#define pipe(x,y,z) _pipe(x,y,z) #else #define aligned_free(x) free(x) #endif