Added testing for hipStreamSynchronize and hipStreamQuery (#2572)

[ROCm/hip commit: b7f5db36d2]

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

@@ -14,6 +14,8 @@ set(TEST_SRC
     hipStreamValue.cc
     hipStreamWithCUMask.cc
     hipStreamACb_MultiThread.cc
+    hipStreamSynchronize.cc
+    hipStreamQuery.cc
     hipStreamWaitEvent.cc
     )
 else()
@@ -32,6 +34,8 @@ set(TEST_SRC
                                    # Fixing would break ABI, to be re-enabled when the fix is made.
     streamCommon.cc
     hipStreamValue.cc
+    hipStreamSynchronize.cc
+    hipStreamQuery.cc
     )
 
     # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)

projects/hip/tests/catch/unit/stream/hipStreamQuery.cc

@@ -0,0 +1,125 @@
+/*
+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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include "streamCommon.hh"
+
+/**
+ * @brief Check that querying a stream with no work returns hipSuccess
+ *
+ **/
+TEST_CASE("Unit_hipStreamQuery_WithNoWork") {
+  hipStream_t stream{nullptr};
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamQuery(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+/**
+ * @brief Check that querying a stream with finished work returns hipSuccess
+ *
+ **/
+TEST_CASE("Unit_hipStreamQuery_WithFinishedWork") {
+  hipStream_t stream{nullptr};
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  hip::stream::empty_kernel<<<dim3(1), dim3(1), 0, stream>>>();
+  HIP_CHECK(hipStreamSynchronize(stream));
+
+  HIP_CHECK(hipStreamQuery(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+#if !HT_NVIDIA
+/**
+ * @brief Check that submitting work to a destroyed stream sets its status as
+ * hipErrorContextIsDestroyed
+ *
+ * Test removed for Nvidia devices because it returns unexpected error
+ */
+TEST_CASE("Unit_hipStreamQuery_WithDestroyedStream") {
+  hipStream_t stream{nullptr};
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed);
+}
+
+/**
+ * @brief Check that submitting work to an uninitialized stream sets its status as
+ * hipErrorContextIsDestroyed
+ *
+ * Test removed for Nvidia devices because it returns unexpected error
+ */
+TEST_CASE("Unit_hipStreamQuery_WithUninitializedStream") {
+  hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
+  HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed);
+}
+#endif
+
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+
+/**
+ * @brief Check that submitting work to a stream sets the status of the nullStream to
+ * hipErrorNotReady
+ *
+ */
+TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") {
+  {
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+
+    HIP_CHECK(hipStreamQuery(hip::nullStream));
+    HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream);
+    HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+    HIP_CHECK(hipDeviceSynchronize());
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+/**
+ * @brief Check that submitting work to the nullStream properly sets its status as
+ * hipErrorNotReady.
+ *
+ */
+TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") {
+  HIP_CHECK(hipStreamQuery(hip::nullStream));
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize(hip::nullStream));
+}
+
+/**
+ * @brief Check that querying a stream with pending work returns hipErrorNotReady
+ *
+ **/
+TEST_CASE("Unit_hipStreamQuery_WithPendingWork") {
+  hipStream_t waitingStream{nullptr};
+  HIP_CHECK(hipStreamCreate(&waitingStream));
+
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), waitingStream);
+
+  HIP_CHECK_ERROR(hipStreamQuery(waitingStream), hipErrorNotReady);
+  HIP_CHECK(hipStreamSynchronize(waitingStream));
+  HIP_CHECK(hipStreamQuery(waitingStream));
+
+  HIP_CHECK(hipStreamDestroy(waitingStream));
+}
+#endif

projects/hip/tests/catch/unit/stream/hipStreamSynchronize.cc

@@ -0,0 +1,156 @@
+/*
+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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include "streamCommon.hh"
+
+namespace hipStreamSynchronizeTest {
+
+/**
+ * @brief Check that hipStreamSynchronize handles empty streams properly.
+ *
+ */
+TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") {
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamSynchronize(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+
+/**
+ * @brief Check that all work executing in a stream is finished after a call to
+ * hipStreamSynchronize.
+ *
+ */
+TEST_CASE("Unit_hipStreamSynchronize_FinishWork") {
+  const hipStream_t explicitStream = reinterpret_cast<hipStream_t>(-1);
+  hipStream_t stream = GENERATE_COPY(explicitStream, hip::nullStream, hip::streamPerThread);
+
+  const bool isExplicitStream = stream == explicitStream;
+  if (isExplicitStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+  }
+
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream);
+  HIP_CHECK(hipStreamSynchronize(stream));
+  HIP_CHECK(hipStreamQuery(stream));
+
+  if (isExplicitStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+/**
+ * @brief Check that synchronizing the nullStream implicitly synchronizes all executing streams.
+ */
+TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") {
+  int totalStreams = 10;
+
+  std::vector<hipStream_t> streams{};
+
+  for (int i = 0; i < totalStreams; ++i) {
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    streams.push_back(stream);
+  }
+
+  for (int i = 0; i < totalStreams; ++i) {
+    HipTest::runKernelForDuration(std::chrono::milliseconds(1000), streams[i]);
+  }
+
+  for (int i = 0; i < totalStreams; ++i) {
+    HIP_CHECK_ERROR(hipStreamQuery(streams[i]), hipErrorNotReady);
+  }
+
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize(hip::nullStream));
+  HIP_CHECK(hipStreamQuery(hip::nullStream));
+
+  for (int i = 0; i < totalStreams; ++i) {
+    HIP_CHECK(hipStreamQuery(streams[i]));
+  }
+
+  for (int i = 0; i < totalStreams; ++i) {
+    HIP_CHECK(hipStreamDestroy(streams[i]));
+  }
+}
+
+/**
+ * @brief Check that synchronizing one stream does implicitly synchronize other streams.
+ *        Check that submiting work to the nullStream does not affect synchronization of other
+ * streams. Check that querying the nullStream does not affect synchronization of other streams.
+ */
+TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-22");
+#else
+
+  hipStream_t stream1;
+  hipStream_t stream2;
+
+  HIP_CHECK(hipStreamCreate(&stream1));
+  HIP_CHECK(hipStreamCreate(&stream2));
+
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream1);
+  HipTest::runKernelForDuration(std::chrono::milliseconds(2000), stream2);
+
+  SECTION("Do not use NullStream") {}
+  SECTION("Submit Kernel to NullStream") {
+    hip::stream::empty_kernel<<<1, 1, 0, hip::nullStream> > >();
+  }
+  SECTION("Query NullStream") {
+    HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+  }
+
+  HIP_CHECK_ERROR(hipStreamQuery(stream1), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
+
+
+  HIP_CHECK(hipStreamSynchronize(stream1));
+  HIP_CHECK(hipStreamQuery(stream1));
+  HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize(stream2));
+  HIP_CHECK(hipStreamQuery(stream2));
+
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+#endif
+}
+
+/**
+ * @brief Check that synchronizing the nullStream also synchronizes the hipStreamPerThread
+ * special stream.
+ *
+ */
+TEST_CASE("Unit_hipStreamSynchronize_NullStreamAndStreamPerThread") {
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::streamPerThread);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipErrorNotReady);
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream);
+  HIP_CHECK(hipStreamSynchronize(hip::nullStream))
+  HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipSuccess);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipSuccess);
+}
+#endif
+}  // namespace hipStreamSynchronizeTest

projects/hip/tests/catch/unit/stream/streamCommon.cc

@@ -66,44 +66,8 @@ bool checkStreamFlags_(hipStream_t stream, bool checkFlags = false, unsigned fla
 
 inline namespace stream {
 
-__device__ int defaultSemaphore = 0;
-
-__global__ void signaling_kernel(int* semaphore) {
-  size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
-  if (tid == 0) {
-    if (semaphore == nullptr) {
-      atomicAdd(&defaultSemaphore, 1);
-    } else {
-      atomicAdd(semaphore, 1);
-    }
-  }
-}
-
-__global__ void waiting_kernel(int* semaphore) {
-  size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
-  if (tid == 0) {
-    if (semaphore == nullptr) {
-      while (atomicCAS(&defaultSemaphore, 1, 2) == 0) {
-      }
-    } else {
-      while (atomicCAS(semaphore, 1, 2) == 0) {
-      }
-    }
-  }
-}
-
-std::thread startSignalingThread(int* semaphore) {
-  std::thread signalingThread([semaphore]() {
-    hipStream_t signalingStream;
-    HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking));
-
-    signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore);
-    HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream));
-    HIP_CHECK_THREAD(hipStreamDestroy(signalingStream));
-  });
-
-  return signalingThread;
-}
+/* Empty kernel to ensure work finishes on the stream quickly */
+__global__ void empty_kernel() {}
 
 bool checkStream(hipStream_t stream) {
   {  // Check default flags

projects/hip/tests/catch/unit/stream/streamCommon.hh

@@ -24,33 +24,12 @@ THE SOFTWARE.
 namespace hip {
 inline namespace stream {
 
+/* Empty kernel to ensure work finishes on the stream quickly */
+__global__ void empty_kernel();
+
 const hipStream_t nullStream = nullptr;
 const hipStream_t streamPerThread = hipStreamPerThread;
 
-/**
- * @brief Kernel that signals a semaphore to change value from 0 to 1.
- *
- * @param semaphore the semaphore that needs to be signaled.
- */
-__global__ void signaling_kernel(int* semaphore = nullptr);
-
-/**
- * @brief Kernel that busy waits until the specified semaphore goes from 0 to 1.
- *
- * @param semaphore the semaphore to wait for.
- */
-__global__ void waiting_kernel(int* semaphore = nullptr);
-
-/**
- * @brief Creates a thread that runs a signaling_kernel on a non-blocking stream.
- * hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream.
- * You must call HIP_CHECK_THREAD_FINALIZE after joining this thread.
- *
- * @param semaphore memory location to signal
- * @return std::thread thread that has to be joined after the testing is done.
- */
-std::thread startSignalingThread(int* semaphore = nullptr);
-
 // Checks stream for valid values of flags and priority
 bool checkStream(hipStream_t stream);