From aea5bb334caeac28cc141e015e7ff5c75e34a196 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 1 Apr 2022 04:32:45 +0100 Subject: [PATCH] Add and restructure tests for hipStreamCreate (#2560) * Add and restructure tests for hipStreamCreate * Add streamCreate to CMakeLists file [ROCm/hip-tests commit: 600157719de4904068de9e8e69533d223c0ed3f2] --- .../catch/unit/stream/CMakeLists.txt | 2 + .../catch/unit/stream/hipStreamCreate.cc | 44 ++---- .../stream/hipStreamCreateWithPriority.cc | 74 +++++++--- .../catch/unit/stream/hipStreamGetPriority.cc | 56 ------- .../catch/unit/stream/streamCommon.cc | 137 ++++++++++++++++++ .../catch/unit/stream/streamCommon.hh | 61 ++++++++ 6 files changed, 268 insertions(+), 106 deletions(-) create mode 100644 projects/hip-tests/catch/unit/stream/streamCommon.cc create mode 100644 projects/hip-tests/catch/unit/stream/streamCommon.hh diff --git a/projects/hip-tests/catch/unit/stream/CMakeLists.txt b/projects/hip-tests/catch/unit/stream/CMakeLists.txt index fc528f2b72..a9d3c35d10 100644 --- a/projects/hip-tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/stream/CMakeLists.txt @@ -9,6 +9,7 @@ set(TEST_SRC hipStreamCreateWithPriority.cc hipStreamGetCUMask.cc hipAPIStreamDisable.cc + streamCommon.cc ) #skipped in windows - duplicate HipTest::vector_square sym (compiler issue) @@ -29,6 +30,7 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipAPIStreamDisable.cc + streamCommon.cc ) endif() diff --git a/projects/hip-tests/catch/unit/stream/hipStreamCreate.cc b/projects/hip-tests/catch/unit/stream/hipStreamCreate.cc index c0caf29b3d..21fff252bc 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamCreate.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamCreate.cc @@ -16,36 +16,20 @@ 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 -TEST_CASE("Unit_hipStreamCreate_default") { - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); -} -TEST_CASE("Unit_hipStreamCreateWithFlags_Negative") { - hipStream_t stream; - auto status = hipStreamCreateWithFlags(&stream, 0xFF); - REQUIRE(status == hipErrorInvalidValue); - status = hipStreamCreateWithFlags(nullptr, hipStreamDefault); - REQUIRE(status == hipErrorInvalidValue); -} -TEST_CASE("Unit_hipStreamCreateWithFlags") { - hipStream_t stream; - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - HIP_CHECK(hipStreamDestroy(stream)); -} -TEST_CASE("Unit_hipStreamCreateWithPriority") { - int priority_low = 0; - int priority_high = 0; - HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); - hipStream_t stream; - SECTION("Setting high prirority") { - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); - } - SECTION("Setting low priority") { - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); - } +#include "streamCommon.hh" + +TEST_CASE("Unit_hipStreamCreate_default") { + int id = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(id)); + + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); // Check if stream has a valid ptr + REQUIRE(hip::checkStream(stream)); // check its flags and priority HIP_CHECK(hipStreamDestroy(stream)); } + +TEST_CASE("Unit_hipStreamCreate_Negative") { + REQUIRE(hipErrorInvalidValue == hipStreamCreate(nullptr)); +} diff --git a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc index 2a75684b50..fadc2966a8 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamCreateWithPriority.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 @@ -43,7 +43,7 @@ kernel tasks on these streams from multiple threads. Validate all the results. 8) Validate stream priorities with event after classifying them as low, medium, high. */ -#include +#include "streamCommon.hh" #include #include #include @@ -500,8 +500,6 @@ bool validateStreamPrioritiesWithEvents() { } // namespace hipStreamCreateWithPriorityTest - - /** Tests following scenarios. 1)Create streams with default flag for all available priority levels and @@ -565,14 +563,12 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") { flag = 0xffffffff. */ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") { - hipStream_t stream; - int priority_low; - int priority_high; - hipError_t ret; + hipStream_t stream{nullptr}; + int priority_low{0}; + int priority_high{0}; // Test is to get the Stream Priority Range - HIP_CHECK( - hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); // Check if priorities are indeed supported if (priority_low == priority_high) { WARN("Stream priority range not supported. Skipping test."); @@ -580,26 +576,64 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") { } SECTION("stream = nullptr test") { - ret = hipStreamCreateWithPriority(nullptr, hipStreamDefault, - priority_low); - - REQUIRE(ret != hipSuccess); + REQUIRE(hipErrorInvalidValue == + hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low)); } SECTION("flag value invalid test") { - ret = hipStreamCreateWithPriority(&stream, 0xffffffff, - priority_low); - - REQUIRE(ret != hipSuccess); + REQUIRE(hipErrorInvalidValue == hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low)); } } +TEST_CASE("Unit_hipStreamCreateWithPriority") { + int id = GENERATE(range(0, HipTest::getDeviceCount())); + + HIP_CHECK(hipSetDevice(id)); + + int priority_low = 0, priority_high = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + hipStream_t stream{nullptr}; + + SECTION("Setting high priority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + REQUIRE(stream != nullptr); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high)); + } + + SECTION("Setting low priority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + REQUIRE(stream != nullptr); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low)); + } + + SECTION("Setting lowest possible priority") { + HIP_CHECK( + hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits::max())); + REQUIRE(stream != nullptr); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low)); + } + + SECTION("Setting highest possible priority") { + HIP_CHECK( + hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits::min())); + REQUIRE(stream != nullptr); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high)); + } + + SECTION("Setting flags to hipStreamNonBlocking") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high)); + REQUIRE(stream != nullptr); + REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, hipStreamNonBlocking)); + } + + HIP_CHECK(hipStreamDestroy(stream)); +} + /** * Validate stream priorities with event after classifying them as low, medium and high. */ TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") { bool TestPassed = true; - TestPassed = hipStreamCreateWithPriorityTest:: - validateStreamPrioritiesWithEvents(); + TestPassed = hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents(); REQUIRE(TestPassed); } diff --git a/projects/hip-tests/catch/unit/stream/hipStreamGetPriority.cc b/projects/hip-tests/catch/unit/stream/hipStreamGetPriority.cc index 29f182ecdf..72839db20d 100644 --- a/projects/hip-tests/catch/unit/stream/hipStreamGetPriority.cc +++ b/projects/hip-tests/catch/unit/stream/hipStreamGetPriority.cc @@ -35,44 +35,6 @@ TEST_CASE("Unit_hipStreamGetPriority_Negative") { REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue); } -/** - * Create stream and check default priority of stream is within range. - */ -TEST_CASE("Unit_hipStreamGetPriority_default") { - 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; - HIP_CHECK(hipStreamCreate(&stream)); - int priority = 0; - HIP_CHECK(hipStreamGetPriority(stream, &priority)); - // valid priority - // Lower the value higher the priority, higher the value lower the priority - REQUIRE(priority_low >= priority); - REQUIRE(priority >= priority_high); - HIP_CHECK(hipStreamDestroy(stream)); -} - -/** - * Create stream with high priority and check priority is set as expected. - */ -TEST_CASE("Unit_hipStreamGetPriority_high") { - 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; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, - priority_high)); - int priority = 0; - HIP_CHECK(hipStreamGetPriority(stream, &priority)); - REQUIRE(priority == priority_high); - HIP_CHECK(hipStreamDestroy(stream)); -} - /** * Create stream with higher priority for the priority range returned. */ @@ -91,24 +53,6 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") { HIP_CHECK(hipStreamDestroy(stream)); } -/** - * Create stream with low priority and check priority is set as expected. - */ -TEST_CASE("Unit_hipStreamGetPriority_low") { - 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; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, - priority_low)); - int priority = 0; - HIP_CHECK(hipStreamGetPriority(stream, &priority)); - REQUIRE(priority_low == priority); - HIP_CHECK(hipStreamDestroy(stream)); -} - /** * Create stream with lower priority for the priority range returned. */ diff --git a/projects/hip-tests/catch/unit/stream/streamCommon.cc b/projects/hip-tests/catch/unit/stream/streamCommon.cc new file mode 100644 index 0000000000..265142e23b --- /dev/null +++ b/projects/hip-tests/catch/unit/stream/streamCommon.cc @@ -0,0 +1,137 @@ +/* +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 "streamCommon.hh" + +namespace hip { + +inline namespace internal { + +bool checkStreamPriority_(hipStream_t stream, bool checkPriority = false, int priority_ = 0) { + int priority{0}; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + if (checkPriority) { + if (priority_ != priority) { + UNSCOPED_INFO("Priority Mismatch, Expected Priority: " << priority_ + << " Actual Priority: " << priority); + return false; + } + } else { + int priority_low{0}, priority_high{0}; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + if (priority_low < priority || priority_high > priority) { + UNSCOPED_INFO("Priority Mismatch, Expected Priority Range: " + << priority_low << " - " << priority_high << " Actual Priority: " << priority); + return false; + } + } + return true; +} + +bool checkStreamFlags_(hipStream_t stream, bool checkFlags = false, unsigned flags_ = 0) { + unsigned flags{0}; + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + if (checkFlags) { + if (flags_ != flags) { + UNSCOPED_INFO("Flags Mismatch, Expected Flag: " << flags_ << " Actual Flag: " << flags); + return false; + } + } else { + if (flags != hipStreamDefault && flags != hipStreamNonBlocking) { + UNSCOPED_INFO("Flags Mismatch, Expected Flag: " << hipStreamDefault << " or " + << hipStreamNonBlocking + << " Actual Flag: " << flags); + return false; + } + } + return true; +} +} // namespace internal + +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(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); + + signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore); + HIP_CHECK(hipStreamSynchronize(signalingStream)); + HIP_CHECK(hipStreamDestroy(signalingStream)); + }); + + return signalingThread; +} + +bool checkStream(hipStream_t stream) { + { // Check default flags + auto res = checkStreamFlags_(stream, true, hipStreamDefault); + if (!res) return false; + } + + { // Check default Priority + auto res = checkStreamPriority_(stream); + if (!res) return false; + } + + return true; +} + +bool checkStreamPriorityAndFlags(hipStream_t stream, int priority, unsigned int flags) { + { // Check flags + auto res = checkStreamFlags_(stream, true, flags); + if (!res) return false; + } + + { // Check priority + auto res = checkStreamPriority_(stream, true, priority); + if (!res) return false; + } + + return true; +} + +} // namespace stream +} // namespace hip diff --git a/projects/hip-tests/catch/unit/stream/streamCommon.hh b/projects/hip-tests/catch/unit/stream/streamCommon.hh new file mode 100644 index 0000000000..db73f1f668 --- /dev/null +++ b/projects/hip-tests/catch/unit/stream/streamCommon.hh @@ -0,0 +1,61 @@ +/* +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. +*/ + +#pragma once + +#include + +namespace hip { +inline namespace stream { + +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. + * + * @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); + +// Checks stream for valid flags and a particular value of priority +bool checkStreamPriorityAndFlags(hipStream_t stream, int priority, + unsigned int flags = hipStreamDefault); + +} // namespace stream +} // namespace hip