From 6c8eeb09d751f3d499ef361fb8759451550c07f7 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Mon, 21 Jun 2021 13:19:44 -0400 Subject: [PATCH] SWDEV-289401 - Catch2 tests for hip event APIs Change-Id: I7c28f842282e07c21656fb92ddbb1a9ad32d752c [ROCm/hip commit: 5816d2075216184ba59a648d1a67aea6cb62226d] --- .../tests/catch/hipTestMain/CMakeLists.txt | 1 + .../tests/catch/include/hip_test_common.hh | 29 +++ .../catch/multiproc/hipMallocConcurrency.cc | 19 +- projects/hip/tests/catch/unit/CMakeLists.txt | 1 + .../hip/tests/catch/unit/event/CMakeLists.txt | 14 ++ .../tests/catch/unit/event/Unit_hipEvent.cc | 187 ++++++++++++++++++ .../unit/event/Unit_hipEventElapsedTime.cc | 100 ++++++++++ .../catch/unit/event/Unit_hipEventIpc.cc | 104 ++++++++++ .../catch/unit/event/Unit_hipEventRecord.cc | 91 +++++++++ .../unit/event/Unit_hipEvent_Negative.cc | 48 +++++ 10 files changed, 576 insertions(+), 18 deletions(-) create mode 100644 projects/hip/tests/catch/unit/event/CMakeLists.txt create mode 100644 projects/hip/tests/catch/unit/event/Unit_hipEvent.cc create mode 100644 projects/hip/tests/catch/unit/event/Unit_hipEventElapsedTime.cc create mode 100644 projects/hip/tests/catch/unit/event/Unit_hipEventIpc.cc create mode 100644 projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc create mode 100644 projects/hip/tests/catch/unit/event/Unit_hipEvent_Negative.cc diff --git a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt index c19f48aa85..c80bbb585e 100644 --- a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt +++ b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt @@ -12,6 +12,7 @@ endif() target_link_libraries(UnitTests PRIVATE UnitDeviceTests MemoryTest StreamTest + EventTest OccupancyTest DeviceTest RTC diff --git a/projects/hip/tests/catch/include/hip_test_common.hh b/projects/hip/tests/catch/include/hip_test_common.hh index d5041be9b9..59a0971f0b 100644 --- a/projects/hip/tests/catch/include/hip_test_common.hh +++ b/projects/hip/tests/catch/include/hip_test_common.hh @@ -49,6 +49,11 @@ THE SOFTWARE. #define HIP_ASSERT(x) \ { REQUIRE((x)); } +#ifdef __cplusplus + #include + #include + #include +#endif // Utility Functions namespace HipTest { @@ -57,4 +62,28 @@ static inline int getDeviceCount() { HIP_CHECK(hipGetDeviceCount(&dev)); return dev; } + +// Returns the current system time in microseconds +static inline long long get_time() { + return std::chrono::high_resolution_clock::now().time_since_epoch() + /std::chrono::microseconds(1); +} + +static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) { + return ((double)(stopTimeUs - startTimeUs)) / ((double)(1000)); +} + +static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) { + int device; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device)); + + unsigned blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } + + return blocks; +} } diff --git a/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc b/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc index 23d28fd40e..f2a01235e8 100644 --- a/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc +++ b/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc @@ -10,29 +10,12 @@ #include #include - #include size_t N = 4 * 1024 * 1024; unsigned blocksPerCU = 6; // to hide latency unsigned threadsPerBlock = 256; - -unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device)); - - unsigned blocks = props.multiProcessorCount * blocksPerCU; - if (blocks * threadsPerBlock > N) { - blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - } - - return blocks; -} - - /** * Validates data consitency on supplied gpu */ @@ -48,7 +31,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { printf("tgs allocating..\n"); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - unsigned blocks = setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); diff --git a/projects/hip/tests/catch/unit/CMakeLists.txt b/projects/hip/tests/catch/unit/CMakeLists.txt index be8a4f2d20..f0ae8b6b93 100644 --- a/projects/hip/tests/catch/unit/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/CMakeLists.txt @@ -1,6 +1,7 @@ add_subdirectory(memory) add_subdirectory(deviceLib) add_subdirectory(stream) +add_subdirectory(event) add_subdirectory(occupancy) add_subdirectory(device) add_subdirectory(rtc) diff --git a/projects/hip/tests/catch/unit/event/CMakeLists.txt b/projects/hip/tests/catch/unit/event/CMakeLists.txt new file mode 100644 index 0000000000..a2fdc1cdd7 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/CMakeLists.txt @@ -0,0 +1,14 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + Unit_hipEvent_Negative.cc + Unit_hipEvent.cc + Unit_hipEventElapsedTime.cc + Unit_hipEventRecord.cc + Unit_hipEventIpc.cc +) + +# Create shared lib of all tests +add_library(EventTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) + +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests EventTest) diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEvent.cc b/projects/hip/tests/catch/unit/event/Unit_hipEvent.cc new file mode 100644 index 0000000000..32cd3c9ca8 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/Unit_hipEvent.cc @@ -0,0 +1,187 @@ +/* +Copyright (c) 2021 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. +*/ +// Tests hipEventRecord and hipEventElapsedTime with different scenarios +// and confirms if these are working as expected +#include +#include + +#include + +int tests = -1; +enum SyncMode { + syncNone, + syncStream, + syncStopEvent, +}; + +const char* syncModeString(int syncMode) { + switch (syncMode) { + case syncNone: + return "syncNone"; + case syncStream: + return "syncStream"; + case syncStopEvent: + return "syncStopEvent"; + default: + return "unknown"; + }; +}; + +void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, hipStream_t stream, + int waitStart, SyncMode syncMode) { + if (!(testMask & tests)) { + return; + } + std::cout << "\n test 0x " << testMask << ": stream= " << stream << " waitStart= " << waitStart + << " syncMode= " << syncModeString(syncMode) << std::endl; + + size_t sizeBytes = numElements * sizeof(int); + + int count = 100; + int init0 = 0; + HIP_CHECK(hipMemset(C_d, init0, sizeBytes)); + for (int i = 0; i < numElements; i++) { + C_h[i] = -1; // initialize + } + + hipEvent_t neverCreated = 0, neverRecorded, timingDisabled; + HIP_CHECK(hipEventCreate(&neverRecorded)); + HIP_CHECK(hipEventCreateWithFlags(&timingDisabled, hipEventDisableTiming)); + + hipEvent_t start, stop; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&stop)); + unsigned blocksPerCU = 6; + unsigned threadsPerBlock = 256; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); + + HIP_CHECK(hipEventRecord(timingDisabled, stream)); + // sandwhich a kernel: + HIP_CHECK(hipEventRecord(start, stream)); + hipLaunchKernelGGL(HipTest::addCountReverse, dim3(blocks), dim3(threadsPerBlock), 0, stream, + static_cast(C_d), C_h, numElements, count); + HIP_CHECK(hipEventRecord(stop, stream)); + + if (waitStart) { + HIP_CHECK(hipEventSynchronize(start)); + } + + hipError_t expectedStopError = hipSuccess; + + // How to wait for the events to finish: + switch (syncMode) { + case syncNone: + expectedStopError = hipErrorNotReady; + break; + case syncStream: + HIP_CHECK(hipStreamSynchronize(stream)); // wait for recording to finish... + break; + case syncStopEvent: + HIP_CHECK(hipEventSynchronize(stop)); + break; + default: + assert(0); + }; + + float t; + + hipError_t e = hipEventElapsedTime(&t, start, start); + if ((e != hipSuccess) && (e != hipErrorNotReady || syncMode != syncNone)) { + printf("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e)); + REQUIRE(false); + } + + if (e == hipSuccess) assert(t == 0.0f); + + // stop usually ready unless we skipped the synchronization (syncNone) + HIP_ASSERT(hipEventElapsedTime(&t, stop, stop) == expectedStopError); + if (e == hipSuccess) assert(t == 0.0f); + + e = hipEventElapsedTime(&t, start, stop); + HIP_ASSERT(e == expectedStopError); + if (expectedStopError == hipSuccess) assert(t > 0.0f); + printf("time=%6.2f error=%s\n", t, hipGetErrorName(e)); + + e = hipEventElapsedTime(&t, stop, start); + HIP_ASSERT(e == expectedStopError); + if (expectedStopError == hipSuccess) assert(t < 0.0f); + printf("negtime=%6.2f error=%s\n", t, hipGetErrorName(e)); + + { + // Check some error conditions for incomplete events: + HIP_ASSERT(hipEventElapsedTime(&t, timingDisabled, stop) == hipErrorInvalidHandle); + HIP_ASSERT(hipEventElapsedTime(&t, start, timingDisabled) == hipErrorInvalidHandle); + + HIP_ASSERT(hipEventElapsedTime(&t, neverCreated, stop) == hipErrorInvalidHandle); + HIP_ASSERT(hipEventElapsedTime(&t, start, neverCreated) == hipErrorInvalidHandle); + + HIP_ASSERT(hipEventElapsedTime(&t, neverRecorded, stop) == hipErrorInvalidHandle); + HIP_ASSERT(hipEventElapsedTime(&t, start, neverRecorded) == hipErrorInvalidHandle); + } + + HIP_CHECK(hipEventDestroy(neverRecorded)); + HIP_CHECK(hipEventDestroy(timingDisabled)); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + + // Clear out everything: + HIP_CHECK(hipDeviceSynchronize()); + + printf("test: OK \n"); +} + +void runTests(int64_t numElements) { + size_t sizeBytes = numElements * sizeof(int); + + printf("test: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes, + sizeBytes / 1024.0 / 1024.0); + + + int *C_h, *C_d; + HIP_CHECK(hipMalloc(&C_d, sizeBytes)); + HIP_CHECK(hipHostMalloc(&C_h, sizeBytes)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithFlags(&stream, 0x0)); + + for (int waitStart = 1; waitStart >= 0; waitStart--) { + unsigned W = waitStart ? 0x1000 : 0; + test(W | 0x01, C_d, C_h, numElements, 0, 0, syncNone); + test(W | 0x02, C_d, C_h, numElements, stream, 0, syncNone); + test(W | 0x04, C_d, C_h, numElements, 0, waitStart, syncStream); + test(W | 0x08, C_d, C_h, numElements, stream, waitStart, syncStream); + test(W | 0x10, C_d, C_h, numElements, 0, waitStart, syncStopEvent); + test(W | 0x20, C_d, C_h, numElements, stream, waitStart, syncStopEvent); + } + + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipHostFree(C_h)); +} + + +TEST_CASE("Unit_hipEvent") { + runTests(80000000); +} diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEventElapsedTime.cc b/projects/hip/tests/catch/unit/event/Unit_hipEventElapsedTime.cc new file mode 100644 index 0000000000..e2fd695fa9 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/Unit_hipEventElapsedTime.cc @@ -0,0 +1,100 @@ +/* +Copyright (c) 2021 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 +#include + +TEST_CASE("Unit_hipEventElapsedTime_NullCheck") { + hipEvent_t start = nullptr, end = nullptr; + float tms = 1.0f; + HIP_ASSERT(hipEventElapsedTime(nullptr, start, end) == hipErrorInvalidValue); +#ifndef __HIP_PLATFORM_NVIDIA__ + // On NVCC platform API throws seg fault hence skipping + HIP_ASSERT(hipEventElapsedTime(&tms, nullptr, end) == hipErrorInvalidHandle); + HIP_ASSERT(hipEventElapsedTime(&tms, start, nullptr) == hipErrorInvalidHandle); +#endif +} + +TEST_CASE("Unit_hipEventElapsedTime_DisableTiming") { + float timeElapsed = 1.0f; + hipEvent_t start, stop; + HIP_CHECK(hipEventCreateWithFlags(&start, hipEventDisableTiming)); + HIP_CHECK(hipEventCreateWithFlags(&stop, hipEventDisableTiming)); + HIP_ASSERT(hipEventElapsedTime(&timeElapsed, start, stop) == hipErrorInvalidHandle); + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); +} + +TEST_CASE("Unit_hipEventElapsedTime_DifferentDevices") { + int devCount = 0; + HIP_CHECK(hipGetDeviceCount(&devCount)); + if (devCount > 1) { + // create event on dev=0 + HIP_CHECK(hipSetDevice(0)); + hipEvent_t start; + hipEvent_t start1; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&start1)); + + HIP_CHECK(hipEventRecord(start, nullptr)); + HIP_CHECK(hipEventSynchronize(start)); + + // create event on dev=1 + HIP_CHECK(hipSetDevice(1)); + hipEvent_t stop; + HIP_CHECK(hipEventCreate(&stop)); + + // start1 on device 0 but null stream on device 1 + HIP_ASSERT(hipEventRecord(start1, nullptr) == hipErrorInvalidHandle); + + HIP_CHECK(hipEventRecord(stop, nullptr)); + HIP_CHECK(hipEventSynchronize(stop)); + + float tElapsed = 1.0f; + // start on device 0 but stop on device 1 + HIP_ASSERT(hipEventElapsedTime(&tElapsed,start,stop) == hipErrorInvalidHandle); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(start1)); + HIP_CHECK(hipEventDestroy(stop)); + } +} + +TEST_CASE("Unit_hipEventElapsedTime") { + hipEvent_t start; + HIP_CHECK(hipEventCreate(&start)); + + hipEvent_t stop; + HIP_CHECK(hipEventCreate(&stop)); + + HIP_CHECK(hipEventRecord(start, nullptr)); + HIP_CHECK(hipEventSynchronize(start)); + + HIP_CHECK(hipEventRecord(stop, nullptr)); + HIP_CHECK(hipEventSynchronize(stop)); + + float tElapsed = 1.0f; + HIP_CHECK(hipEventElapsedTime(&tElapsed, start, stop)); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); +} diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEventIpc.cc b/projects/hip/tests/catch/unit/event/Unit_hipEventIpc.cc new file mode 100644 index 0000000000..1d59089bf4 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/Unit_hipEventIpc.cc @@ -0,0 +1,104 @@ +/* +Copyright (c) 2021 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. +*/ + +// Test hipEventIpc behavior. + +#include +#include + +#include + + +TEST_CASE("Unit_hipEventIpc") { + size_t N = 4 * 1024 * 1024; + unsigned threadsPerBlock = 256; + int iterations = 1; + + unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + if (blocks > 1024) blocks = 1024; + if (blocks == 0) blocks = 1; + + printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N, + ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations); + printf("iterations=%d\n", iterations); + + 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 start, stop; + + // NULL stream check: + HIP_CHECK(hipEventCreateWithFlags(&start, hipEventDisableTiming|hipEventInterprocess)); + HIP_CHECK(hipEventCreateWithFlags(&stop, hipEventDisableTiming|hipEventInterprocess)); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + + for (int i = 0; i < iterations; i++) { + //--- START TIMED REGION + long long hostStart = HipTest::get_time(); + // Record the start event + HIP_CHECK(hipEventRecord(start, NULL)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A_d), static_cast(B_d), C_d, N); + + + HIP_CHECK(hipEventRecord(stop, NULL)); + HIP_CHECK(hipEventSynchronize(stop)); + HIP_CHECK(hipEventQuery(stop)); + long long hostStop = HipTest::get_time(); + //--- STOP TIMED REGION + + + float eventMs = 1.0f; + // should fail due to hipEventDisableTiming + REQUIRE(hipSuccess != hipEventElapsedTime(&eventMs, start, stop)); + float hostMs = HipTest::elapsed_time(hostStart, hostStop); + + printf("host_time (chrono) =%6.3fms\n", hostMs); + printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs); + printf("\n"); + + } + + hipIpcEventHandle_t ipc_handle; + HIP_CHECK(hipIpcGetEventHandle(&ipc_handle, start)); + + hipEvent_t ipc_event; + hipError_t err = hipIpcOpenEventHandle(&ipc_event, ipc_handle); + + // hipIpcOpenEventHandle() should be called in a different process, hence it should fail here + REQUIRE(err == hipErrorInvalidContext); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + +} diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc b/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc new file mode 100644 index 0000000000..c3dc1a6086 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -0,0 +1,91 @@ +/* +Copyright (c) 2021 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. +*/ +// Test hipEventRecord serialization behavior. +// Through manual inspection of the reported timestamps, can determine if recording a NULL event +// forces synchronization : set +#include +#include + +#include + +TEST_CASE("Unit_hipEventRecord") { + size_t N = 4 * 1024 * 1024; + unsigned threadsPerBlock = 256; + int iterations = 1; + + unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + if (blocks > 1024) blocks = 1024; + if (blocks == 0) blocks = 1; + + printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N, + ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations); + printf("iterations=%d\n", iterations); + + 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 start, stop; + + // NULL stream check: + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&stop)); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + for (int i = 0; i < iterations; i++) { + //--- START TIMED REGION + long long hostStart = HipTest::get_time(); + // Record the start event + HIP_CHECK(hipEventRecord(start, NULL)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A_d), static_cast(B_d), C_d, N); + + HIP_CHECK(hipEventRecord(stop, NULL)); + HIP_CHECK(hipEventSynchronize(stop)); + long long hostStop = HipTest::get_time(); + //--- STOP TIMED REGION + + float eventMs = 1.0f; + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + float hostMs = HipTest::elapsed_time(hostStart, hostStop); + + printf("host_time (chrono) =%6.3fms\n", hostMs); + printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs); + printf("\n"); + + // Make sure timer is timing something... + REQUIRE(eventMs > 0.0f); + } + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); + + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + +} diff --git a/projects/hip/tests/catch/unit/event/Unit_hipEvent_Negative.cc b/projects/hip/tests/catch/unit/event/Unit_hipEvent_Negative.cc new file mode 100644 index 0000000000..1d8b2b29a4 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/Unit_hipEvent_Negative.cc @@ -0,0 +1,48 @@ +/* +Copyright (c) 2021 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_hipEventCreate_NullCheck") { + hipEvent_t start_event; + auto res = hipEventCreate(nullptr); + REQUIRE(res != hipSuccess); + res = hipEventCreateWithFlags(nullptr, 0); + REQUIRE(res != hipSuccess); + res = hipEventCreateWithFlags(&start_event, 10); + REQUIRE(res != hipSuccess); +} + +TEST_CASE("Unit_hipEventSynchronize_NullCheck") { + auto res = hipEventSynchronize(nullptr); + REQUIRE(res != hipSuccess); +} + +TEST_CASE("Unit_hipEventQuery_NullCheck") { + auto res = hipEventQuery(nullptr); + REQUIRE(res != hipSuccess); +} + +TEST_CASE("Unit_hipEventDestroy_NullCheck") { + auto res = hipEventDestroy(nullptr); + REQUIRE(res != hipSuccess); +}