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<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]++;
   }
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<std::string, int> 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<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) {
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 <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.";                                       \
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