AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

EXSWHTEC-304 - Expand and refactor tests for clock device functions (#273)

Browse Source 
- Expand and refactor clock device functions
- Add comments and format code
- Fix doxygen test group comments

[ROCm/hip-tests commit: f73d1b21ff]
This commit is contained in:
nives-vukovic
2023-07-18 09:30:51 +02:00
committed by GitHub
parent d660c0a84d
commit 64167b64c4
3 changed files with 149 additions and 68 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/include/hip_test_defgroups.hh
+6
View File
@@ -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.
*/
projects/hip-tests/catch/unit/CMakeLists.txt
+1 -1
View File
@@ -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()
projects/hip-tests/catch/unit/clock/hipClockCheck.cc
+142 -67
View File
@@ -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 <hip_test_checkers.hh>
#include <hip/hip_ext.h>
#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));
}