EXSWCPHIPT-140 - Added hipStreamCreateWithFlags tests (#2549)

[ROCm/hip commit: 46717c6ffe]

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

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -17,23 +17,69 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-/**
-Testcase Scenarios :
-1) Validates functionality of hipStreamCreateWithFlags when stream = nullptr.
-2) Validates functionality of hipStreamCreateWithFlags when flag = 0xffffffff.
-*/
-
+#include <chrono>
 #include <hip_test_common.hh>
 
+namespace hipStreamCreateWithFlagsTests {
 
-TEST_CASE("Unit_hipStreamCreateWithFlags_ArgValidation") {
-  // stream = nullptr test
-  SECTION("stream is nullptr") {
-    REQUIRE(hipStreamCreateWithFlags(nullptr, hipStreamDefault) != hipSuccess);
-  }
-  // flag value invalid test
-  SECTION("flag value invalid") {
-    hipStream_t stream;
-    REQUIRE(hipStreamCreateWithFlags(&stream, 0xffffffff) != hipSuccess);
-  }
+TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_NullStream") {
+  HIP_CHECK_ERROR(hipStreamCreateWithFlags(nullptr, hipStreamDefault), hipErrorInvalidValue);
 }
+
+TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_InvalidFlag") {
+  hipStream_t stream{};
+  unsigned int flag = 0xFF;
+  REQUIRE(flag != hipStreamDefault);
+  REQUIRE(flag != hipStreamNonBlocking);
+  HIP_CHECK_ERROR(hipStreamCreateWithFlags(&stream, flag), hipErrorInvalidValue);
+}
+
+// create a stream and check the properties are correctly set
+TEST_CASE("Unit_hipStreamCreateWithFlags_Default") {
+  const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest));
+
+  unsigned int flag{};
+  HIP_CHECK(hipStreamGetFlags(stream, &flag));
+  REQUIRE(flag == flagUnderTest);
+
+  int priority{};
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  // zero is considered default priority
+  REQUIRE(priority == 0);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+// a stream will default to blocking the null stream, but will not block the null stream when
+// created with hipStreamNonBlocking
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+TEST_CASE("Unit_hipStreamCreateWithFlags_DefaultStreamInteraction") {
+  const hipStream_t defaultStream = GENERATE(static_cast<hipStream_t>(nullptr), hipStreamPerThread);
+  const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  const hipError_t expectedError = (flagUnderTest == hipStreamDefault) && (defaultStream == nullptr)
+      ? hipErrorNotReady
+      : hipSuccess;
+  CAPTURE(defaultStream, flagUnderTest, expectedError, hipGetErrorString(expectedError));
+
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest));
+
+  constexpr auto delay = std::chrono::milliseconds(500);
+
+  SECTION("default stream waiting for created stream") {
+    HipTest::runKernelForDuration(delay, stream);
+    REQUIRE(hipStreamQuery(defaultStream) == expectedError);
+  }
+  SECTION("created stream waiting for default stream") {
+    HipTest::runKernelForDuration(delay, defaultStream);
+    REQUIRE(hipStreamQuery(stream) == expectedError);
+  }
+
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif
+
+}  // namespace hipStreamCreateWithFlagsTests