SWDEV-466663-[catch2][dtest] _spt api names missing from code coverage-Phase-1

Change-Id: Ifd374fcc242767ab7a9e999dba79f2f3c2ef7c4a
2024-07-26 20:45:14 +05:30
commit 4e659f42de
@@ -1174,6 +1174,13 @@
        "Unit_hipHostAlloc_Negative_NonCoherent",
        "Unit_hipHostAlloc_Negative_Coherent",
        "Unit_hipHostAlloc_Negative_NumaUser",
+        "=== Following tests disabled due to SWDEV-486363",
+        "Unit_hipStreamQuery_spt_WithFinishedWork",
+        "Unit_hipStreamQuery_spt_NegativeCases",
+        "Unit_hipStreamQuery_spt_WithPendingWork",
+        "Unit_hipStreamSynchronize_spt_FinishWork",
+        "Unit_hipStreamSynchronize_spt_SynchronizeStreamAndQueryNullStream",
+        "====================================================",
    #endif
        "End of json"
    ]
@@ -24,7 +24,11 @@ set(TEST_SRC

 if(HIP_PLATFORM MATCHES "amd")
  set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc
-               hipStreamACb_MultiThread.cc hipStreamWaitEvent.cc)
+               hipStreamACb_MultiThread.cc hipStreamWaitEvent.cc
+	       hipStreamGetFlags_spt.cc
+               hipStreamGetPriority_spt.cc
+               hipStreamQuery_spt.cc
+               hipStreamSynchronize_spt.cc)
 else()
  set(TEST_SRC
      ${TEST_SRC}
@@ -0,0 +1,111 @@
+/*
+Copyright (c) 2024 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>
+#include <hip_test_defgroups.hh>
+/**
+ * @addtogroup hipStreamGetFlags_spt hipStreamGetFlags_spt
+ * @{
+ * @ingroup StreamTest
+ * `hipStreamGetFlags_spt(hipStream_t stream, unsigned int* flags)` -
+ * Return flags associated with a stream on stream per thread
+ */
+
+ /**
+ * Test Description
+ * ------------------------
+ *  - Basic test to verify that hipStreamGetFlags_spt returns the same
+ * flags that were used to create the stream.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetFlags_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamGetFlags_spt_Basic") {
+  unsigned int expectedFlag = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  unsigned int returnedFlags;
+  hipStream_t stream;
+
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, expectedFlag));
+  HIP_CHECK(hipStreamGetFlags_spt(stream, &returnedFlags));
+  REQUIRE((returnedFlags & expectedFlag) == expectedFlag);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+ /**
+ * Test Description
+ * ------------------------
+ *  - Tests to verify Negative scenarios for the hipStreamGetFlags_spt API
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetFlags_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+
+TEST_CASE("Unit_hipStreamGetFlags_spt_Negative") {
+  hipStream_t validStream;
+  unsigned int flags;
+
+  HIP_CHECK(hipStreamCreate(&validStream));
+
+  SECTION("Nullptr Stream && Valid Flags") {
+    #if HT_AMD
+      HIP_CHECK_ERROR(hipStreamGetFlags_spt(nullptr, &flags), hipSuccess);
+    #elif HT_NVIDIA
+      HIP_CHECK(hipStreamGetFlags_spt(nullptr, &flags));
+    #endif
+  }
+
+  SECTION("Valid Stream && Nullptr Flags") {
+    HIP_CHECK_ERROR(hipStreamGetFlags_spt(validStream, nullptr),
+                                          hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipStreamDestroy(validStream));
+}
+ /**
+ * Test Description
+ * ------------------------
+ *  - Test flag value when streams created with CUMask
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetFlags_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+
+#if HT_AMD
+TEST_CASE("Unit_hipStreamGetFlags_spt_StreamsCreatedWithCUMask") {
+  hipStream_t stream;
+  unsigned int flags;
+  const uint32_t cuMask = 0xffffffff;
+  HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
+  HIP_CHECK(hipStreamGetFlags_spt(stream, &flags));
+  REQUIRE(flags == hipStreamDefault);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif
+
+/**
+* End doxygen group StreamTest.
+* @}
+*/
@@ -0,0 +1,156 @@
+/*
+Copyright (c) 2024 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>
+#include <hip_test_defgroups.hh>
+
+/**
+ * @addtogroup hipStreamGetPriority_spt hipStreamGetPriority_spt
+ * @{
+ * @ingroup StreamTest
+ * `hipStreamGetPriority_spt(hipStream_t stream, int* priority)` -
+ * Query the priority of a stream
+ */
+
+ /**
+ * Test Description
+ * ------------------------
+ * - Test to create a stream and check priority.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetPriority_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamGetPriority_spt_BasicTst") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream{};
+  int priority = 0;
+  SECTION("Null Stream") {
+    HIP_CHECK(hipStreamGetPriority_spt(nullptr, &priority));
+    // valid priority
+    REQUIRE(priority_low >= priority);
+    REQUIRE(priority >= priority_high);
+  }
+  SECTION("Created Stream") {
+    SECTION("Default Priority") {
+      HIP_CHECK(hipStreamCreate(&stream));
+      HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+      // valid priority
+      // Lower the value higher the priority,
+      // higher the value lower the priority
+      REQUIRE(priority_low >= priority);
+      REQUIRE(priority >= priority_high);
+    }
+    SECTION("High Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
+                                            priority_high));
+      HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+      REQUIRE(priority == priority_high);
+    }
+    SECTION("Higher Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
+                                            priority_high - 1));
+      HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+      REQUIRE(priority == priority_high);
+    }
+    SECTION("Low Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
+                                            priority_low));
+      HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+      REQUIRE(priority_low == priority);
+    }
+    SECTION("Lower Priority") {
+      HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
+                                            priority_low + 1));
+      HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+      REQUIRE(priority_low == priority);
+    }
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+ /**
+ * Test Description
+ * ------------------------
+ * - Test check stream priority with negative cases.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetPriority_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamGetPriority_spt_NegativeCases") {
+  // stream as nullptr and priority as nullptr
+  auto res = hipStreamGetPriority_spt(nullptr, nullptr);
+  REQUIRE(res == hipErrorInvalidValue);
+  // priority as nullptr
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+  res = hipStreamGetPriority_spt(stream, nullptr);
+  REQUIRE(res == hipErrorInvalidValue);
+  HIP_CHECK(hipStreamDestroy(stream));
+  // stream as nullptr
+  int priority = -1;
+  HIP_CHECK(hipStreamGetPriority_spt(nullptr, &priority));
+}
+ /**
+ * Test Description
+ * ------------------------
+ * - Create stream with CUMask and check priority is returned as expected.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamGetPriority_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+
+#if HT_AMD
+TEST_CASE("Unit_hipStreamGetPriority_spt_StreamsWithCUMask") {
+  hipStream_t stream{};
+  int priority = 0;
+  int priority_normal = 0;
+  int priority_low = 0;
+  int priority_high = 0;
+  // Test is to get the Stream Priority Range
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  priority_normal = (priority_low + priority_high) / 2;
+  // Check if priorities are indeed supported
+  REQUIRE_FALSE(priority_low == priority_high);
+  // Creating a stream with hipExtStreamCreateWithCUMask and checking
+  // priority.
+  const uint32_t cuMask = 0xffffffff;
+  HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
+  HIP_CHECK(hipStreamGetPriority_spt(stream, &priority));
+  REQUIRE_FALSE(priority_normal != priority);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif
+
+/**
+* End doxygen group StreamTest.
+* @}
+*/
@@ -0,0 +1,151 @@
+/*
+Copyright (c) 2024 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>
+#include <utils.hh>
+#include "streamCommon.hh"
+/**
+ * @addtogroup hipStreamQuery_spt hipStreamQuery_spt
+ * @{
+ * @ingroup StreamTest
+ * `hipStreamQuery_spt(hipStream_t stream)` -
+ * Returns status of a stream
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ * - Test to query a stream with no work returns hipSuccess or not.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamQuery_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamQuery_spt_WithNoWork") {
+  hipStream_t stream{nullptr};
+
+  SECTION("Null Stream") {
+    HIP_CHECK(hipStreamQuery_spt(stream));
+  }
+
+  SECTION("Created Stream") {
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamQuery_spt(stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+/**
+ * Test Description
+ * ------------------------
+ * - Test to query a stream with finished work returns hipSuccess or not.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamQuery_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamQuery_spt_WithFinishedWork") {
+  hipStream_t stream{nullptr};
+
+  SECTION("Null Stream") {
+    hip::stream::empty_kernel<<<dim3(1), dim3(1), 0, stream>>>();
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HIP_CHECK(hipStreamQuery_spt(stream));
+  }
+
+  SECTION("Created Stream") {
+    HIP_CHECK(hipStreamCreate(&stream));
+    hip::stream::empty_kernel<<<dim3(1), dim3(1), 0, stream>>>();
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HIP_CHECK(hipStreamQuery_spt(stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+/**
+ * Test Description
+ * ------------------------
+ * - Negative cases to query a stream.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamQuery_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+#if HT_AMD
+TEST_CASE("Unit_hipStreamQuery_spt_NegativeCases") {
+  SECTION("Query Destroyed Stream") {
+    hipStream_t stream{nullptr};
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK_ERROR(hipStreamQuery_spt(stream), hipErrorContextIsDestroyed);
+  }
+  SECTION("Query Uninitialized Stream") {
+    hipStream_t streamUnInit{reinterpret_cast<hipStream_t>(0xFFFF)};
+    HIP_CHECK_ERROR(hipStreamQuery_spt(streamUnInit),
+                                       hipErrorContextIsDestroyed);
+  }
+  SECTION("Submit Work On Stream And Query Null Stream") {
+    hipStream_t ValidStream;
+    HIP_CHECK(hipStreamCreate(&ValidStream));
+
+    HIP_CHECK(hipStreamQuery_spt(hip::nullStream));
+    LaunchDelayKernel(std::chrono::milliseconds(500), ValidStream);
+    HIP_CHECK_ERROR(hipStreamQuery_spt(hip::nullStream), hipSuccess);
+    HIP_CHECK(hipDeviceSynchronize());
+    HIP_CHECK(hipStreamDestroy(ValidStream));
+  }
+  SECTION("Submit work to the nullStream") {
+    HIP_CHECK(hipStreamQuery_spt(hip::nullStream));
+    LaunchDelayKernel(std::chrono::milliseconds(500), hip::nullStream);
+    HIP_CHECK_ERROR(hipStreamQuery_spt(hip::nullStream), hipSuccess);
+    HIP_CHECK(hipStreamSynchronize(hip::nullStream));
+  }
+}
+/**
+ * Test Description
+ * ------------------------
+ * - Test to querying a stream with pending work returns hipErrorNotReady or not.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamQuery_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamQuery_spt_WithPendingWork") {
+  hipStream_t waitingStream{nullptr};
+  HIP_CHECK(hipStreamCreate(&waitingStream));
+  LaunchDelayKernel(std::chrono::milliseconds(500), waitingStream);
+  HIP_CHECK_ERROR(hipStreamQuery_spt(waitingStream), hipErrorNotReady);
+  HIP_CHECK(hipStreamSynchronize(waitingStream));
+  HIP_CHECK(hipStreamQuery_spt(waitingStream));
+  HIP_CHECK(hipStreamDestroy(waitingStream));
+}
+#endif
+
+/**
+* End doxygen group StreamTest.
+* @}
+*/
+
@@ -0,0 +1,139 @@
+/*
+Copyright (c) 2024 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>
+#include <hip_test_defgroups.hh>
+#include <utils.hh>
+#include "streamCommon.hh"
+/**
+ * @addtogroup hipStreamSynchronize_spt hipStreamSynchronize_spt
+ * @{
+ * @ingroup StreamTest
+ * `hipStreamSynchronize_spt(hipStream_t stream)` -
+ * Query the priority of a stream
+ */
+
+ /**
+ * Test Description
+ * ------------------------
+ * - Test to verify that hipStreamSynchronize_spt handles empty streams properly.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamSynchronize_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamSynchronize_spt_EmptyStream") {
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamSynchronize_spt(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+ /**
+ * Test Description
+ * ------------------------
+ * - Check that synchronization of uninitialized stream sets its status to
+ * hipErrorContextIsDestroyed
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamSynchronize_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+
+#if HT_AMD
+TEST_CASE("Unit_hipStreamSynchronize_spt_UninitializedStream") {
+  hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
+  HIP_CHECK_ERROR(hipStreamSynchronize_spt(stream), hipErrorContextIsDestroyed);
+}
+ /**
+ * Test Description
+ * ------------------------
+ * - Check that all work executing in a stream is finished after a call to
+ * hipStreamSynchronize_spt.
+ * Test source
+ * ------------------------
+ *  - /unit/stream/hipStreamSynchronize_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamSynchronize_spt_FinishWork") {
+  hipStream_t explicitStream  = reinterpret_cast<hipStream_t>(-1);
+  hipStream_t stream = GENERATE_COPY(explicitStream, hip::nullStream,
+                                     hip::streamPerThread);
+  if (explicitStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+  }
+
+  LaunchDelayKernel(std::chrono::milliseconds(500), stream);
+  HIP_CHECK(hipStreamSynchronize_spt(stream));
+  HIP_CHECK(hipStreamQuery(stream));
+
+  if (explicitStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+ /**
+ * Test Description
+ * ------------------------
+ * - 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 source
+ * ------------------------
+ *  - /unit/stream/hipStreamSynchronize_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+
+TEST_CASE("Unit_hipStreamSynchronize_spt_SynchronizeStreamAndQueryNullStream") {
+  hipStream_t stream1;
+  hipStream_t stream2;
+
+  HIP_CHECK(hipStreamCreate(&stream1));
+  HIP_CHECK(hipStreamCreate(&stream2));
+
+  LaunchDelayKernel(std::chrono::milliseconds(500), stream1);
+  LaunchDelayKernel(std::chrono::milliseconds(2000), stream2);
+  hip::stream::empty_kernel<<<1, 1, 0, hip::nullStream>>>();
+
+  HIP_CHECK_ERROR(hipStreamQuery(stream1), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize_spt(stream1));
+  HIP_CHECK(hipStreamQuery(stream1));
+  HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady);
+
+  HIP_CHECK(hipStreamSynchronize_spt(stream2));
+  HIP_CHECK(hipStreamQuery(stream2));
+
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+}
+#endif
+/**
+* End doxygen group StreamTest.
+* @}
+*/