diff --git a/projects/hip/tests/catch/unit/stream/hipStreamCreateWithFlags.cc b/projects/hip/tests/catch/unit/stream/hipStreamCreateWithFlags.cc index 95606f418a..ef38755028 100644 --- a/projects/hip/tests/catch/unit/stream/hipStreamCreateWithFlags.cc +++ b/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 #include +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(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