Add tests for hipStreamWaitEvent (#2559)

[ROCm/hip-tests commit: 867f1276b7]

projects/hip-tests/catch/CMakeLists.txt

@@ -201,6 +201,11 @@ set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n")
 set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n")
 file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo})
 
+# Enable device lambda on nvidia platforms
+if(HIP_COMPILER MATCHES "nvcc")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --extended-lambda")
+endif()
+
 # Disable CXX extensions (gnu++11 etc)
 set(CMAKE_CXX_EXTENSIONS OFF)
 

projects/hip-tests/catch/unit/stream/CMakeLists.txt

@@ -13,6 +13,7 @@ set(TEST_SRC
     hipStreamValue.cc
     hipStreamWithCUMask.cc
     hipStreamACb_MultiThread.cc
+    hipStreamWaitEvent.cc
     )
 else()
 set(TEST_SRC
@@ -31,7 +32,7 @@ set(TEST_SRC
     hipStreamValue.cc
     )
 
-    set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
+    # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
 endif()
 
 hip_add_exe_to_target(NAME StreamTest

projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc

@@ -0,0 +1,121 @@
+/*
+Copyright (c) 2022 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.
+*/
+
+
+#include <hip_test_common.hh>
+
+TEST_CASE("Unit_hipStreamWaitEvent_Negative") {
+  enum class StreamTestType { NullStream = 0, StreamPerThread, CreatedStream };
+
+  auto streamType = GENERATE(StreamTestType::NullStream, StreamTestType::StreamPerThread,
+                             StreamTestType::CreatedStream);
+
+  hipStream_t stream{nullptr};
+  hipEvent_t event{nullptr};
+
+  if (streamType == StreamTestType::StreamPerThread) {
+    stream = hipStreamPerThread;
+  } else if (streamType == StreamTestType::CreatedStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+  }
+
+  HIP_CHECK(hipEventCreate(&event));
+
+  REQUIRE((stream != nullptr) != (streamType == StreamTestType::NullStream));
+  REQUIRE(event != nullptr);
+
+  SECTION("Invalid Event") {
+    INFO("Running against Invalid Event");
+    HIP_CHECK_ERROR(hipStreamWaitEvent(stream, nullptr, 0), hipErrorInvalidResourceHandle);
+  }
+
+  SECTION("Invalid Flags") {
+    INFO("Running against Invalid Flags");
+    constexpr unsigned flag = ~0u;
+    REQUIRE(flag != 0);
+    HIP_CHECK_ERROR(hipStreamWaitEvent(stream, event, flag), hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipEventDestroy(event));
+
+  if (streamType == StreamTestType::CreatedStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate.
+// This wont be accurate since clock rate of a GPU varies depending on many variables including
+// thermals, load, utilization
+__global__ void waitKernel(int clockRate, int seconds) {
+  auto start = clock();
+  auto ms = seconds * 1000;
+  long long waitTill = clockRate * (long long)ms;
+  while (1) {
+    auto end = clock();
+    if ((end - start) > waitTill) {
+      return;
+    }
+  }
+}
+
+TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") {
+  hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr};
+  hipEvent_t waitEvent{nullptr};
+
+  HIP_CHECK(hipStreamCreate(&blockedStreamA));
+  HIP_CHECK(hipStreamCreate(&streamBlockedOnStreamA));
+  HIP_CHECK(hipStreamCreate(&unblockingStream));
+  HIP_CHECK(hipEventCreate(&waitEvent));
+
+  REQUIRE(blockedStreamA != nullptr);
+  REQUIRE(streamBlockedOnStreamA != nullptr);
+  REQUIRE(waitEvent != nullptr);
+
+  hipDeviceProp_t prop{};
+  HIP_CHECK(hipGetDeviceProperties(&prop, 0));
+  auto clockRate = prop.clockRate;
+
+  waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate,
+                                          3);  // wait for 5 seconds
+  HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA));
+
+  // Make sure stream is waiting for data to be set
+  HIP_CHECK_ERROR(hipEventQuery(waitEvent), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamWaitEvent(streamBlockedOnStreamA, waitEvent, 0));
+
+  waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2);  // Wait for 2 seconds
+
+  // Make sure stream is waiting for event on blockedStreamA
+  HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize(unblockingStream));
+
+  HIP_CHECK(hipStreamSynchronize(blockedStreamA));
+  HIP_CHECK(hipStreamSynchronize(streamBlockedOnStreamA));
+
+  // Check that both streams have finished
+  HIP_CHECK(hipStreamQuery(blockedStreamA));
+  HIP_CHECK(hipStreamQuery(streamBlockedOnStreamA));
+
+  HIP_CHECK(hipStreamDestroy(blockedStreamA));
+  HIP_CHECK(hipStreamDestroy(streamBlockedOnStreamA));
+  HIP_CHECK(hipStreamDestroy(unblockingStream));
+  HIP_CHECK(hipEventDestroy(waitEvent));
+}