diff --git a/projects/hip/tests/catch/unit/stream/CMakeLists.txt b/projects/hip/tests/catch/unit/stream/CMakeLists.txt index 9b829207cf..8518ff07cb 100644 --- a/projects/hip/tests/catch/unit/stream/CMakeLists.txt +++ b/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) diff --git a/projects/hip/tests/catch/unit/stream/hipStreamQuery.cc b/projects/hip/tests/catch/unit/stream/hipStreamQuery.cc new file mode 100644 index 0000000000..ea2e9d0476 --- /dev/null +++ b/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 +#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<<>>(); + 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(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 diff --git a/projects/hip/tests/catch/unit/stream/hipStreamSynchronize.cc b/projects/hip/tests/catch/unit/stream/hipStreamSynchronize.cc new file mode 100644 index 0000000000..0555620b1f --- /dev/null +++ b/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 +#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(-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 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 \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/stream/streamCommon.cc b/projects/hip/tests/catch/unit/stream/streamCommon.cc index 14ac4000eb..6a2904bbee 100644 --- a/projects/hip/tests/catch/unit/stream/streamCommon.cc +++ b/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 diff --git a/projects/hip/tests/catch/unit/stream/streamCommon.hh b/projects/hip/tests/catch/unit/stream/streamCommon.hh index 1d5a1ea958..2017b9298a 100644 --- a/projects/hip/tests/catch/unit/stream/streamCommon.hh +++ b/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);