From 1d464f391ea53f07087f1f62c1e53bfc0ba32350 Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <sameer.sahasrabuddhe@amd.com>
Date: Thu, 9 Apr 2020 10:20:11 +0530
Subject: [PATCH] printf test: loop with divergent exit condition

Change-Id: I1071e4a240a280332bde669701c72228b9dea2df
---
 tests/src/printf/hipPrintfBasic.cpp | 37 +++++++++++++++++++++++++++++
 tests/src/test_common.h             |  4 ++++
 2 files changed, 41 insertions(+)

diff --git a/tests/src/printf/hipPrintfBasic.cpp b/tests/src/printf/hipPrintfBasic.cpp
index e51373c251..7d0c57cdc3 100644
--- a/tests/src/printf/hipPrintfBasic.cpp
+++ b/tests/src/printf/hipPrintfBasic.cpp
@@ -219,6 +219,42 @@ static void test_series(int *retval, uint num_blocks, uint threads_per_block) {
   HIPASSERT(linecount[msg_short] == num_threads);
 }
 
+__global__ void kernel_divergent_loop() {
+  DECLARE_DATA();
+
+  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int result = 0;
+
+  for (int i = 0; i <= tid; ++i) {
+    printf("%d\n", i);
+  }
+}
+
+static void test_divergent_loop(uint num_blocks, uint threads_per_block) {
+  CaptureStream captured(stdout);
+
+  uint num_threads = num_blocks * threads_per_block;
+
+  hipLaunchKernelGGL(kernel_divergent_loop, dim3(num_blocks), dim3(threads_per_block),
+                     0, 0);
+  hipStreamSynchronize(0);
+  auto CapturedData = captured.getCapturedData();
+
+  std::map<int, int> count;
+  while (true) {
+    int i;
+    CapturedData >> i;
+    if (CapturedData.fail())
+      break;
+    count[i]++;
+  }
+
+  HIPASSERT(count.size() == num_threads);
+  for (int i = 0; i != num_threads; ++i) {
+    HIPASSERT(count[i] == num_threads - i);
+  }
+}
+
 int main() {
   uint num_blocks = 1;
   uint threads_per_block = 64;
@@ -233,6 +269,7 @@ int main() {
   test_divergent0(retval, num_blocks, threads_per_block);
   test_divergent1(retval, num_blocks, threads_per_block);
   test_series(retval, num_blocks, threads_per_block);
+  test_divergent_loop(num_blocks, threads_per_block);
 
   passed();
 }
diff --git a/tests/src/test_common.h b/tests/src/test_common.h
index 7d8c39e74c..ce65dfaf33 100644
--- a/tests/src/test_common.h
+++ b/tests/src/test_common.h
@@ -55,11 +55,15 @@ THE SOFTWARE.
     printf("%sPASSED!%s\n", KGRN, KNRM);                                                           \
     exit(0);
 
+// The real "assert" would have written to stderr. But it is
+// sufficient to just fflush here without getting pedantic. This also
+// ensures that we don't lose any earlier writes to stdout.
 #define failed(...)                                                                                \
     printf("%serror: ", KRED);                                                                     \
     printf(__VA_ARGS__);                                                                           \
     printf("\n");                                                                                  \
     printf("error: TEST FAILED\n%s", KNRM);                                                        \
+    fflush(nullptr);                                                                               \
     abort();
 
 #define warn(...)                                                                                  \