SWDEV-232428 - Fix HIP printf tests on Windows

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


[ROCm/clr commit: 55377dfee2]

projects/clr/hipamd/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();

projects/clr/hipamd/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<std::string, int> 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<std::string, int> 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<uint> 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<std::string, int> 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<std::string, int> 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<int, int> count;
   while (true) {
     int i;
-    CapturedData >> i;
-    if (CapturedData.fail())
+    dataStream >> i;
+    if (dataStream.fail())
       break;
     count[i]++;
   }

projects/clr/hipamd/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();

projects/clr/hipamd/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<std::string, int> linecount;
-  for (std::string line; std::getline(CapturedData, line);) {
+  for (std::string line; std::getline(dataStream, line);) {
     linecount[line]++;
   }
 

projects/clr/hipamd/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<std::string, int> 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<std::string, int> 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<std::string, int> 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<std::string, int> 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<uint> 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) {

projects/clr/hipamd/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();

projects/clr/hipamd/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();

projects/clr/hipamd/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();

projects/clr/hipamd/tests/src/printf/printf_common.h

@@ -2,14 +2,84 @@
 #define COMMON_H
 
 #include <errno.h>
-#include <error.h>
 #include <fstream>
 #include <iostream>
 #include <map>
 #include <stdlib.h>
 #include <string>
-#include <unistd.h>
+#include <fcntl.h>
+#include <sys/stat.h>
+#include <chrono>
 
+#if defined(_WIN32)
+#include <io.h>
+#else
+#include <error.h>
+#include <unistd.h>
+#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<char*>(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.";                                       \

projects/clr/hipamd/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