From 9e373a398fba7ba87efad3dc532767dc2cc4802c Mon Sep 17 00:00:00 2001 From: nives-vukovic <110852104+nives-vukovic@users.noreply.github.com> Date: Mon, 26 Sep 2022 07:32:00 +0200 Subject: [PATCH] EXSWHTEC-16 - Implement hipEventSynchronize tests (#2904) - Add hipEventSynchronize.cc file with positive unit tests: Simple recorded event synchronization and without event recording - Add newly created file into CMakeLists.txt [ROCm/hip-tests commit: bc83dae816aad96130760772243019de749a03b3] --- .../hip-tests/catch/unit/event/CMakeLists.txt | 1 + .../catch/unit/event/hipEventSynchronize.cc | 130 ++++++++++++++++++ 2 files changed, 131 insertions(+) create mode 100644 projects/hip-tests/catch/unit/event/hipEventSynchronize.cc diff --git a/projects/hip-tests/catch/unit/event/CMakeLists.txt b/projects/hip-tests/catch/unit/event/CMakeLists.txt index ba5d37e746..b9f57c1451 100644 --- a/projects/hip-tests/catch/unit/event/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/event/CMakeLists.txt @@ -7,6 +7,7 @@ set(TEST_SRC Unit_hipEventIpc.cc hipEventDestroy.cc hipEventCreateWithFlags.cc + hipEventSynchronize.cc ) # The test used wait mechanism and doesnt play well with all arch of nvidia diff --git a/projects/hip-tests/catch/unit/event/hipEventSynchronize.cc b/projects/hip-tests/catch/unit/event/hipEventSynchronize.cc new file mode 100644 index 0000000000..e3c07bb395 --- /dev/null +++ b/projects/hip-tests/catch/unit/event/hipEventSynchronize.cc @@ -0,0 +1,130 @@ +/* +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. +*/ +/* +Testcase Scenarios : +Unit_hipEventSynchronize_Default_Positive- Test synchronization of an event that is completed after a simple kernel launch (on null/created stream) +Unit_hipEventSynchronize_NoEventRecord_Positive - Test synchronization of an event that has not been recorded +*/ + +#include + +#include +#include + +void testSynchronize(hipStream_t stream) { + + constexpr size_t N = 1024; + + constexpr int blocks = 1024; + + constexpr size_t Nbytes = N * sizeof(float); + + float *A_h, *B_h, *C_h; + float *A_d, *B_d, *C_d; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + hipEvent_t end_event; + HIP_CHECK(hipEventCreate(&end_event)); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + HipTest::launchKernel(HipTest::vectorADD, blocks, 1, 0, stream, + static_cast(A_d), static_cast(B_d), + C_d, N); + + if ( stream != nullptr ) + { + HIP_CHECK(hipStreamSynchronize(stream)); + } + + // Record the end_event + HIP_CHECK(hipEventRecord(end_event, nullptr)); + // Wait for the end_event to complete + HIP_CHECK(hipEventSynchronize(end_event)); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipEventDestroy(end_event)); + + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +TEST_CASE("Unit_hipEventSynchronize_Default_Positive") { + hipStream_t stream{nullptr}; + + SECTION("Kernel launched in null stream") { + testSynchronize(stream); + } + + SECTION ("Kernel launched in created stream") { + HIP_CHECK(hipStreamCreate(&stream)); + testSynchronize(stream); + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +TEST_CASE("Unit_hipEventSynchronize_NoEventRecord_Positive") { + constexpr size_t N = 1024; + + constexpr int blocks = 1024; + + constexpr size_t Nbytes = N * sizeof(float); + + float *A_h, *B_h, *C_h; + float *A_d, *B_d, *C_d; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + hipEvent_t dummy_event; + HIP_CHECK(hipEventCreate(&dummy_event)); + + hipEvent_t end_event; + HIP_CHECK(hipEventCreate(&end_event)); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + HipTest::launchKernel(HipTest::vectorADD, blocks, 1, 0, 0, + static_cast(A_d), static_cast(B_d), + C_d, N); + + // Record the end_event + HIP_CHECK(hipEventRecord(end_event, NULL)); + + // When hipEventSynchronized is called on event that has not been recorded, + // the function returns immediately + HIP_CHECK(hipEventSynchronize(dummy_event)); + + // End event has not been completed + HIP_CHECK_ERROR(hipEventQuery(end_event), hipErrorNotReady); + // Wait for end_event to complete + HIP_CHECK(hipEventSynchronize(end_event)); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipEventDestroy(dummy_event)); + HIP_CHECK(hipEventDestroy(end_event)); + + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +}