From fb0006cafe893875ad56d86634e151bec013c8e0 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 3 Aug 2022 12:17:34 +0100 Subject: [PATCH] Add tests for hipStreamWaitEvent (#2559) [ROCm/hip-tests commit: 867f1276b791c7d516dd191b7b5765a27e7e3f3c] --- projects/hip-tests/catch/CMakeLists.txt | 5 + .../catch/unit/stream/CMakeLists.txt | 3 +- .../catch/unit/stream/hipStreamWaitEvent.cc | 121 ++++++++++++++++++ 3 files changed, 128 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index a8a488002d..f9e11b9c82 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -201,6 +201,11 @@ set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n") set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n") file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo}) +# Enable device lambda on nvidia platforms +if(HIP_COMPILER MATCHES "nvcc") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --extended-lambda") +endif() + # Disable CXX extensions (gnu++11 etc) set(CMAKE_CXX_EXTENSIONS OFF) diff --git a/projects/hip-tests/catch/unit/stream/CMakeLists.txt b/projects/hip-tests/catch/unit/stream/CMakeLists.txt index 229df157c1..6cdbac9a4e 100644 --- a/projects/hip-tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/stream/CMakeLists.txt @@ -13,6 +13,7 @@ set(TEST_SRC hipStreamValue.cc hipStreamWithCUMask.cc hipStreamACb_MultiThread.cc + hipStreamWaitEvent.cc ) else() set(TEST_SRC @@ -31,7 +32,7 @@ set(TEST_SRC hipStreamValue.cc ) - set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) + # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc b/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc new file mode 100644 index 0000000000..cea39fcd54 --- /dev/null +++ b/projects/hip-tests/catch/unit/stream/hipStreamWaitEvent.cc @@ -0,0 +1,121 @@ +/* +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 + +TEST_CASE("Unit_hipStreamWaitEvent_Negative") { + enum class StreamTestType { NullStream = 0, StreamPerThread, CreatedStream }; + + auto streamType = GENERATE(StreamTestType::NullStream, StreamTestType::StreamPerThread, + StreamTestType::CreatedStream); + + hipStream_t stream{nullptr}; + hipEvent_t event{nullptr}; + + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + + HIP_CHECK(hipEventCreate(&event)); + + REQUIRE((stream != nullptr) != (streamType == StreamTestType::NullStream)); + REQUIRE(event != nullptr); + + SECTION("Invalid Event") { + INFO("Running against Invalid Event"); + HIP_CHECK_ERROR(hipStreamWaitEvent(stream, nullptr, 0), hipErrorInvalidResourceHandle); + } + + SECTION("Invalid Flags") { + INFO("Running against Invalid Flags"); + constexpr unsigned flag = ~0u; + REQUIRE(flag != 0); + HIP_CHECK_ERROR(hipStreamWaitEvent(stream, event, flag), hipErrorInvalidValue); + } + + HIP_CHECK(hipEventDestroy(event)); + + if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +// Since we can not use atomic*_system on every gpu, we will use wait based on clock rate. +// This wont be accurate since clock rate of a GPU varies depending on many variables including +// thermals, load, utilization +__global__ void waitKernel(int clockRate, int seconds) { + auto start = clock(); + auto ms = seconds * 1000; + long long waitTill = clockRate * (long long)ms; + while (1) { + auto end = clock(); + if ((end - start) > waitTill) { + return; + } + } +} + +TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { + hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr}; + hipEvent_t waitEvent{nullptr}; + + HIP_CHECK(hipStreamCreate(&blockedStreamA)); + HIP_CHECK(hipStreamCreate(&streamBlockedOnStreamA)); + HIP_CHECK(hipStreamCreate(&unblockingStream)); + HIP_CHECK(hipEventCreate(&waitEvent)); + + REQUIRE(blockedStreamA != nullptr); + REQUIRE(streamBlockedOnStreamA != nullptr); + REQUIRE(waitEvent != nullptr); + + hipDeviceProp_t prop{}; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + auto clockRate = prop.clockRate; + + waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate, + 3); // wait for 5 seconds + HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA)); + + // Make sure stream is waiting for data to be set + HIP_CHECK_ERROR(hipEventQuery(waitEvent), hipErrorNotReady); + + HIP_CHECK(hipStreamWaitEvent(streamBlockedOnStreamA, waitEvent, 0)); + + waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds + + // Make sure stream is waiting for event on blockedStreamA + HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady); + + HIP_CHECK(hipStreamSynchronize(unblockingStream)); + + HIP_CHECK(hipStreamSynchronize(blockedStreamA)); + HIP_CHECK(hipStreamSynchronize(streamBlockedOnStreamA)); + + // Check that both streams have finished + HIP_CHECK(hipStreamQuery(blockedStreamA)); + HIP_CHECK(hipStreamQuery(streamBlockedOnStreamA)); + + HIP_CHECK(hipStreamDestroy(blockedStreamA)); + HIP_CHECK(hipStreamDestroy(streamBlockedOnStreamA)); + HIP_CHECK(hipStreamDestroy(unblockingStream)); + HIP_CHECK(hipEventDestroy(waitEvent)); +}