From fb938a607f66a1aef2289b7e42e4e7020fe03fa0 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 19 Oct 2022 19:12:01 +0530 Subject: [PATCH] SWDEV-286739 - Support hipDeviceAttributeWallClockRate (#2984) Part 4: Add unit test of wall_clock64() Change-Id: I52cbba6d67d21cde6da19c5ab533159f426a9bf7 --- tests/catch/unit/CMakeLists.txt | 3 + tests/catch/unit/clock/CMakeLists.txt | 29 +++++ tests/catch/unit/clock/hipClockCheck.cc | 116 ++++++++++++++++++ .../module/hipExtLaunchKernelGGL.cpp | 10 +- 4 files changed, 153 insertions(+), 5 deletions(-) create mode 100644 tests/catch/unit/clock/CMakeLists.txt create mode 100644 tests/catch/unit/clock/hipClockCheck.cc diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 9874ce29d2..0ce9be5ed9 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/tests/catch/unit/CMakeLists.txt @@ -31,3 +31,6 @@ add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) +if(HIP_PLATFORM STREQUAL "amd") +add_subdirectory(clock) +endif() diff --git a/tests/catch/unit/clock/CMakeLists.txt b/tests/catch/unit/clock/CMakeLists.txt new file mode 100644 index 0000000000..c181f1a76b --- /dev/null +++ b/tests/catch/unit/clock/CMakeLists.txt @@ -0,0 +1,29 @@ +# 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. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + hipClockCheck.cc +) + +hip_add_exe_to_target(NAME ClockCheckTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/tests/catch/unit/clock/hipClockCheck.cc b/tests/catch/unit/clock/hipClockCheck.cc new file mode 100644 index 0000000000..4d817a3764 --- /dev/null +++ b/tests/catch/unit/clock/hipClockCheck.cc @@ -0,0 +1,116 @@ +/* +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 +#include +#include + +#define ONESECOND 1000 // in ms +#define HALFSECOND 500 // in ms + +enum CLOCK_MODE { + CLOCK_MODE_CLOCK64, + CLOCK_MODE_WALL_CLOCK64 +}; + +__global__ void kernel_c(int clockRate, uint64_t wait_t) { + uint64_t start = clock64() / clockRate, cur = 0; // in ms + do { cur = clock64() / clockRate-start;} while (cur < wait_t); +} + +__global__ void kernel_w(int clockRate, uint64_t wait_t) { + uint64_t start = wall_clock64() / clockRate, cur = 0; // in ms + do { cur = wall_clock64() / clockRate-start;} while (cur < wait_t); +} + +bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2, + float expectedTime1, float expectedTime2) { + bool testStatus = false; + float ratio = m == CLOCK_MODE_CLOCK64 ? 0.5 : 0.01; + + if (fabs(time1 - expectedTime1) < ratio * expectedTime1 + && fabs(time2 - expectedTime2) < ratio * expectedTime2) { + WARN("Succeeded: Expected Vs Actual: Kernel1 - " << expectedTime1 << " Vs " << time1 + << ", Kernel2 - " << expectedTime2 << " Vs " << time2); + testStatus = true; + } else { + FAIL_CHECK("Failed: Expected Vs Actual: Kernel1 -" << expectedTime1 << " Vs " << time1 + << ", Kernel2 - " << expectedTime2 << " Vs " << time2); + testStatus = false; + } + return testStatus; +} + +/* + * Launching kernel1 and kernel2 and then we try to + * get the event elapsed time of each kernel using the start and + * end events.The event elapsed time should return us the kernel + * execution time for that particular kernel +*/ +bool kernelTimeExecution(CLOCK_MODE m, int clockRate, + uint64_t expectedTime1, uint64_t expectedTime2) { + hipStream_t stream; + hipEvent_t start_event1, end_event1, start_event2, end_event2; + float time1 = 0, time2 = 0; + HIPCHECK(hipEventCreate(&start_event1)); + HIPCHECK(hipEventCreate(&end_event1)); + HIPCHECK(hipEventCreate(&start_event2)); + HIPCHECK(hipEventCreate(&end_event2)); + HIPCHECK(hipStreamCreate(&stream)); + hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, + dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, clockRate, expectedTime1); + hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, + dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, clockRate, expectedTime2); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipEventElapsedTime(&time1, start_event1, end_event1)); + HIPCHECK(hipEventElapsedTime(&time2, start_event2, end_event2)); + + HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipEventDestroy(start_event1)); + HIPCHECK(hipEventDestroy(end_event1)); + HIPCHECK(hipEventDestroy(start_event2)); + HIPCHECK(hipEventDestroy(end_event2)); + + return verifyTimeExecution(m, time1, time2, expectedTime1, expectedTime2); +} + +TEST_CASE("Unit_hipClock64_Check") { + HIPCHECK(hipSetDevice(0)); + int clockRate = 0; // in KHz + HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeClockRate, 0)); + + SECTION("Verify kernel execution time via clock64()") { + CHECK(kernelTimeExecution(CLOCK_MODE_CLOCK64, clockRate, ONESECOND, HALFSECOND)); + } +} + +TEST_CASE("Unit_hipWallClock64_Check") { + HIPCHECK(hipSetDevice(0)); + int clockRate = 0; // in KHz + HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeWallClockRate, 0)); + + if(!clockRate) { + INFO("hipDeviceAttributeWallClockRate has not been supported. Skipped"); + return; + } + + SECTION("Verify kernel execution time via wall_clock64()") { + CHECK(kernelTimeExecution(CLOCK_MODE_WALL_CLOCK64, clockRate, ONESECOND, HALFSECOND)); + } +} diff --git a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp index e6f5c49a9b..4df6bd8600 100644 --- a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp +++ b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp @@ -34,8 +34,8 @@ #include "test_common.h" #include "hip/hip_ext.h" -#define FOURSEC_KERNEL 4999 -#define TWOSEC_KERNEL 2999 +#define FIVESEC_KERNEL 4999 +#define THREESEC_KERNEL 2999 __device__ int globalvar = 1; __global__ void TwoSecKernel(int clockrate) { @@ -163,12 +163,12 @@ bool KernelTimeExecution() { e = hipEventElapsedTime(&time_4sec, start_event1, end_event1); e = hipEventElapsedTime(&time_2sec, start_event2, end_event2); - if ( (time_4sec < static_cast(FOURSEC_KERNEL)) && - (time_2sec < static_cast(TWOSEC_KERNEL))) { + if ( (time_4sec < static_cast(FIVESEC_KERNEL)) && + (time_2sec < static_cast(THREESEC_KERNEL))) { testStatus = true; } else { printf("Expected Vs Actual: Kernel1-<%d Vs %f Kernel2-<%d Vs %f\n", - FOURSEC_KERNEL, time_4sec, TWOSEC_KERNEL, time_2sec); + FIVESEC_KERNEL, time_4sec, THREESEC_KERNEL, time_2sec); testStatus = false; }