SWDEV-393637 - Standardize the way to implement wait mechanism (#316)

Change-Id: I59027667806878191f9c641f27ce47e6f85ba40c

[ROCm/hip-tests commit: 5f4af7cde4]

projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json

@@ -41,10 +41,6 @@
 	   "Unit_hipStreamCreateWithFlags_DefaultStreamInteraction",
 	   "Unit_hipStreamWaitEvent_UninitializedStream_Negative",
 	   "Unit_hipDeviceSetSharedMemConfig_Negative_Parameters",
-	   "Disabling tests tracked with SWDEV-394083",
-	   "Unit_hipDeviceSynchronize_Positive_Nullstream",
-	   "Disabling tests tracked with SWDEV-393637",
-	   "Unit_hipDeviceSynchronize_Functional",
 	   "Unit_hipMemset3DSync",
 	   "Unit_hipStreamAddCallback_StrmSyncTiming",
 	   "Disabling test tracked SWDEV-394199",

projects/hip-tests/catch/include/hip_test_common.hh

@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #pragma once
 #include "hip_test_context.hh"
+
 #include <catch.hpp>
 #include <atomic>
 #include <chrono>
@@ -30,6 +31,7 @@ THE SOFTWARE.
 #include <iomanip>
 #include <mutex>
 #include <cstdlib>
+#include <thread>
 
 #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
 
@@ -429,6 +431,52 @@ static inline void runKernelForDuration(std::chrono::milliseconds duration,
   hipLaunchKernelGGL(waitKernel_used, dim3(1), dim3(1), 0, stream, ticksPerSecond * millis / 1000);
 }
 
+class BlockingContext {
+  std::atomic_bool blocked{true};
+  hipStream_t stream;
+
+ public:
+  BlockingContext(hipStream_t s) : stream(s), blocked(true) {}
+
+  BlockingContext(const BlockingContext& in) {
+    blocked = in.blocked_val();
+    stream = in.stream_val();
+  }
+
+  BlockingContext(const BlockingContext&& in) {
+    blocked = in.blocked_val();
+    stream = in.stream_val();
+  }
+
+  void reset() { blocked = true; }
+
+  BlockingContext& operator=(const BlockingContext& in) {
+    blocked = in.blocked_val();
+    stream = in.stream_val();
+    return *this;
+  }
+
+  void block_stream() {
+    blocked = true;
+    auto blocking_callback = [](hipStream_t, hipError_t, void* data) {
+      auto blocked = reinterpret_cast<std::atomic_bool*>(data);
+      while (blocked->load()) {
+        // Yield this thread till we are waiting
+        std::this_thread::yield();
+      }
+    };
+    HIP_CHECK(hipStreamAddCallback(stream, blocking_callback, (void*)&blocked, 0));
+  }
+
+  void unblock_stream() {
+    blocked = false;
+  }
+
+  bool is_blocked() const { return hipStreamQuery(stream) == hipErrorNotReady; }
+
+  bool blocked_val() const { return blocked.load(); }
+  hipStream_t stream_val() const { return stream; }
+};
 }  // namespace HipTest
 
 // This must be called in the beginning of image test app's main() to indicate whether image

projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc

@@ -33,15 +33,15 @@ THE SOFTWARE.
 #define NUM_ITERS 1 << 30
 
 static __global__ void Iter(int* Ad, int num) {
-    int tx = threadIdx.x + blockIdx.x * blockDim.x;
-    // Kernel loop designed to execute very slowly.
-    // so we can test timing-related
-    // behavior below
-    if (tx == 0) {
-        for (int i = 0; i < num; i++) {
-            Ad[tx] += 1;
-        }
+  int tx = threadIdx.x + blockIdx.x * blockDim.x;
+  // Kernel loop designed to execute very slowly.
+  // so we can test timing-related
+  // behavior below
+  if (tx == 0) {
+    for (int i = 0; i < num; i++) {
+      Ad[tx] += 1;
     }
+  }
 }
 
 TEST_CASE("Unit_hipDeviceSynchronize_Positive_Empty_Streams") {
@@ -61,42 +61,47 @@ TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") {
   INFO("Current device: " << device);
 
   int *A_h = nullptr, *A_d = nullptr;
+  HipTest::BlockingContext b_context{nullptr};
   HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), _SIZE, hipHostMallocDefault));
   A_h[0] = 1;
   HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), _SIZE));
 
   HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL));
+  b_context.block_stream();
+  REQUIRE(b_context.is_blocked());
   hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30);
   HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL));
 
-  CHECK(1 << 30 != A_h[0] - 1);
+  REQUIRE(1 << 30 != A_h[0] - 1);
+  b_context.unblock_stream();
   HIP_CHECK(hipDeviceSynchronize());
-  CHECK(1 << 30 == A_h[0] - 1);
+  REQUIRE(1 << 30 == A_h[0] - 1);
 }
 
 TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
   int* A[NUM_STREAMS];
   int* Ad[NUM_STREAMS];
   hipStream_t stream[NUM_STREAMS];
+  std::vector<HipTest::BlockingContext> b_context;
+  b_context.reserve(NUM_STREAMS);
+
   for (int i = 0; i < NUM_STREAMS; i++) {
-      HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE,
-                                                   hipHostMallocDefault));
-      A[i][0] = 1;
-      HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
-      HIP_CHECK(hipStreamCreate(&stream[i]));
+    HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE, hipHostMallocDefault));
+    A[i][0] = 1;
+    HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
+    HIP_CHECK(hipStreamCreate(&stream[i]));
+    b_context.emplace_back(HipTest::BlockingContext(stream[i]));
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
-      HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice,
-                                                               stream[i]));
+    HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
-      hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0,
-                                                stream[i], Ad[i], NUM_ITERS);
-      HIP_CHECK(hipGetLastError());
+    b_context[i].block_stream();
+    REQUIRE(b_context[i].is_blocked());
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], NUM_ITERS);
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
-      HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost,
-                                                               stream[i]));
+    HIP_CHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
   }
 
 
@@ -106,7 +111,10 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
   // Conservative implementations which synchronize the hipMemcpyAsync will
   // fail, ie if HIP_LAUNCH_BLOCKING=true.
 
-  CHECK(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
+  REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
+  for (int i = 0; i < NUM_STREAMS; i++) {
+    b_context[i].unblock_stream();
+  }
   HIP_CHECK(hipDeviceSynchronize());
-  CHECK(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
+  REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
 }