SWDEV-227201: Introduce tests for printf on hostcall

Tests that check POSIX specifiers with a single thread:
 - hipPrintfSpecifiers.cpp     : all conversion specifiers
 - hipPrintfFlags.cpp          : common flags that modify conversions
 - hipPrintfAltForms.cpp       : alternate forms ('#')
 - hipPrintfStar.cpp           : additional arguments ('*')
 - hipPrintfWidthPrecision.cpp : floating point details

Tests that check functionality on top of hostcall
 - hipPrintfBasic.cpp       : divergent calls, series of calls, return value, etc
 - hipPrintfManyWaves.cpp   : many waves printing together
 - hipPrintfManyDevices.cpp : many waves on many devices

Change-Id: I35e069f4c542f896999239996dc89eda0faad7b8

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();
+}

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 <vector>
+
+// 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<std::string, int> 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<std::string, int> 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<uint> 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<std::string, int> 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<std::string, int> 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<int *>(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();
+}

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();
+}

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<std::string, int> 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();
+}

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 <vector>
+
+// 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<std::string, int> 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<std::string, int> 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<std::string, int> 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<std::string, int> 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<uint> 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<int *>(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();
+}

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();
+}

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();
+}

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();
+}

tests/src/printf/printf_common.h

@@ -0,0 +1,94 @@
+#ifndef COMMON_H
+#define COMMON_H
+
+#include <errno.h>
+#include <error.h>
+#include <fstream>
+#include <iostream>
+#include <map>
+#include <stdlib.h>
+#include <string>
+#include <unistd.h>
+
+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