diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index f0966ae93b..402625006a 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -135,3 +135,9 @@ THE SOFTWARE. * This section describes tests for the Device Language API. * @} */ + +/** + * @defgroup DeviceLanguageTest Device Language + * @{ + * This section describes tests for the Device Language API. + */ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index e0b1f5a1b2..50e2ca0f58 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -44,7 +44,7 @@ add_subdirectory(executionControl) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) -#add_subdirectory(clock) +add_subdirectory(clock) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) endif() diff --git a/projects/hip-tests/catch/unit/clock/hipClockCheck.cc b/projects/hip-tests/catch/unit/clock/hipClockCheck.cc index 4d817a3764..148ba143c8 100644 --- a/projects/hip-tests/catch/unit/clock/hipClockCheck.cc +++ b/projects/hip-tests/catch/unit/clock/hipClockCheck.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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 @@ -21,40 +21,51 @@ THE SOFTWARE. #include #include -#define ONESECOND 1000 // in ms -#define HALFSECOND 500 // in ms +/** + * @addtogroup clock clock + * @{ + * @ingroup DeviceLanguageTest + * Contains unit tests for clock, clock64 and wall_clock64 APIs + */ -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_c64(int clock_rate, uint64_t wait_t) { + uint64_t start = clock64() / clock_rate, cur = 0; // in ms + do { + cur = clock64() / clock_rate - 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); +__global__ void kernel_c(int clock_rate, uint64_t wait_t) { + uint64_t start = clock() / clock_rate, cur = 0; // in ms + do { + cur = clock() / clock_rate - 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; +__global__ void kernel_wc64(int clock_rate, uint64_t wait_t) { + uint64_t start = wall_clock64() / clock_rate, cur = 0; // in ms + do { + cur = wall_clock64() / clock_rate - start; + } while (cur < wait_t); +} - 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; +bool verify_time_execution(float ratio, float time1, float time2, float expected_time1, + float expected_time2) { + bool test_status = false; + + if (fabs(time1 - expected_time1) < ratio * expected_time1 && + fabs(time2 - expected_time2) < ratio * expected_time2) { + INFO("Succeeded: Expected Vs Actual: Kernel1 - " << expected_time1 << " Vs " << time1 + << ", Kernel2 - " << expected_time2 << " Vs " + << time2); + test_status = true; } else { - FAIL_CHECK("Failed: Expected Vs Actual: Kernel1 -" << expectedTime1 << " Vs " << time1 - << ", Kernel2 - " << expectedTime2 << " Vs " << time2); - testStatus = false; + INFO("Failed: Expected Vs Actual: Kernel1 -" << expected_time1 << " Vs " << time1 + << ", Kernel2 - " << expected_time2 << " Vs " + << time2); + test_status = false; } - return testStatus; + return test_status; } /* @@ -62,55 +73,119 @@ bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2, * 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) { + */ +bool kernel_time_execution(void (*kernel)(int, uint64_t), int clock_rate, uint64_t expected_time1, + uint64_t expected_time2) { 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)); + HIP_CHECK(hipEventCreate(&start_event1)); + HIP_CHECK(hipEventCreate(&end_event1)); + HIP_CHECK(hipEventCreate(&start_event2)); + HIP_CHECK(hipEventCreate(&end_event2)); + HIP_CHECK(hipStreamCreate(&stream)); + hipExtLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, + clock_rate, expected_time1); + hipExtLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, + clock_rate, expected_time2); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipEventElapsedTime(&time1, start_event1, end_event1)); + HIP_CHECK(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)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipEventDestroy(start_event1)); + HIP_CHECK(hipEventDestroy(end_event1)); + HIP_CHECK(hipEventDestroy(start_event2)); + HIP_CHECK(hipEventDestroy(end_event2)); - return verifyTimeExecution(m, time1, time2, expectedTime1, expectedTime2); + float ratio = kernel == kernel_wc64 ? 0.01 : 0.5; + + return verify_time_execution(ratio, time1, time2, expected_time1, expected_time2); } -TEST_CASE("Unit_hipClock64_Check") { - HIPCHECK(hipSetDevice(0)); - int clockRate = 0; // in KHz - HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeClockRate, 0)); +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function clock64. Kernel execution time is calculated through elapsed time between + * the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipClock64_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, 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"); + if (IsGfx11()) { + HipTest::HIP_SKIP_TEST("Issue with clock64() function on gfx11 devices!"); return; } - SECTION("Verify kernel execution time via wall_clock64()") { - CHECK(kernelTimeExecution(CLOCK_MODE_WALL_CLOCK64, clockRate, ONESECOND, HALFSECOND)); - } + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_c64, clock_rate, expected_time1, expected_time2)); +} + +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function clock. Kernel execution time is calculated through elapsed time between + * the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipClock_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); + + if (IsGfx11()) { + HipTest::HIP_SKIP_TEST("Issue with clock() function on gfx11 devices!"); + return; + } + + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_c, clock_rate, expected_time1, expected_time2)); +} + +/** + * Test Description + * ------------------------ + * - Launches two kernels that run for a specified amount of time passed as a kernel argument by + * using device function wall_clock64. Kernel execution time is calculated through elapsed time + * between the start and end event, and calculated time is compared with passed time values. + * Test source + * ------------------------ + * - catch/unit/clock/hipClockCheck.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipWallClock64_Positive_Basic") { + HIP_CHECK(hipSetDevice(0)); + int clock_rate = 0; // in kHz + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeWallClockRate, 0)); + + if (!clock_rate) { + HipTest::HIP_SKIP_TEST("hipDeviceAttributeWallClockRate is not supported"); + return; + } + + const auto expected_time1 = GENERATE(1000, 1500, 2000); + const auto expected_time2 = expected_time1 / 2; + + REQUIRE(kernel_time_execution(kernel_wc64, clock_rate, expected_time1, expected_time2)); }