From e057b9e471379ef018461464cad024b1c596daff Mon Sep 17 00:00:00 2001 From: Vladana Stojiljkovic Date: Fri, 15 Mar 2024 17:54:34 +0100 Subject: [PATCH] SWDEV-450735 - Add compiler option for using clock64 Change-Id: I9efed88d691ee1b2b4465286b3340b820f7bf627 [ROCm/hip-tests commit: 5b67a2a4a02221eb73e70b7b36400fbc6118d624] --- projects/hip-tests/catch/CMakeLists.txt | 6 ++++++ projects/hip-tests/catch/include/hip_test_common.hh | 6 ++++++ projects/hip-tests/catch/include/utils.hh | 4 ++-- .../catch/multiproc/hipMemCoherencyTstMProc.cc | 4 ++-- .../unit/cooperativeGrps/hipCGGridGroupType_old.cc | 4 ++-- .../cooperativeGrps/hipCGMultiGridGroupType_old.cc | 8 ++++---- .../hipLaunchCooperativeKernelMultiDevice_old.cc | 12 ++++++------ .../hip-tests/catch/unit/memory/hipMemPoolApi.cc | 6 +++--- .../catch/unit/memory/hipMemoryAllocateCoherent.cc | 6 +++--- .../hip-tests/catch/unit/memory/mempool_common.hh | 6 +++--- .../unit/streamperthread/hipStreamPerThrdTsts.cc | 8 ++++---- 11 files changed, 41 insertions(+), 29 deletions(-) diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index ecd12b3608..b249a8a431 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -8,6 +8,12 @@ project(hiptests) option(ENABLE_ADDRESS_SANITIZER "Option to enable ASAN build" OFF) option(BUILD_SHARED_LIBS "Option for testing shared libraries" ON) + +option(TEST_CLOCK_CYCLE "Option to use clock64" OFF) +if (TEST_CLOCK_CYCLE) + add_definitions(-DTEST_CLOCK_CYCLE) +endif() + # flag to generate standalone exe per src file. message(STATUS "STANDALONE_TESTS : ${STANDALONE_TESTS}") diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 97a16f7be8..7fec627179 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -35,6 +35,12 @@ THE SOFTWARE. #include #include +#ifdef TEST_CLOCK_CYCLE +#define clock_function() clock64() +#else +#define clock_function() wall_clock64() +#endif + #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); // Not thread-safe diff --git a/projects/hip-tests/catch/include/utils.hh b/projects/hip-tests/catch/include/utils.hh index c195dbd292..74c12d0a42 100644 --- a/projects/hip-tests/catch/include/utils.hh +++ b/projects/hip-tests/catch/include/utils.hh @@ -140,8 +140,8 @@ template __global__ void VectorSet(T* const vec, const T value, siz static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { while (interval--) { #if HT_AMD - uint64_t start = wall_clock64(); - while (wall_clock64() - start < ticks_per_ms) { + uint64_t start = clock_function(); + while (clock_function() - start < ticks_per_ms) { __builtin_amdgcn_s_sleep(10); } #endif diff --git a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc index c80068ca36..dec36e6dec 100644 --- a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -55,10 +55,10 @@ __global__ void CoherentTst_gfx11(int *ptr, int PeakClk) { #if HT_AMD // Incrementing the value by 1 int64_t GpuFrq = int64_t(PeakClk) * 1000; - int64_t StrtTck = wall_clock64(); + int64_t StrtTck = clock_function(); atomicAdd(ptr, 1); // The following while loop checks the value in ptr for around 3-4 seconds - while ((wall_clock64() - StrtTck) <= (3 * GpuFrq)) { + while ((clock_function() - StrtTck) <= (3 * GpuFrq)) { if (atomicCAS(ptr, 3, 4) == 3) break; } #endif diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc index a1f6cdaa87..2b70b6719f 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc @@ -187,9 +187,9 @@ __global__ void test_kernel_gfx11(unsigned int* atomic_val, unsigned int* array, // reach the atomicInc, but everyone will have only hit the atomic once. if (rank == (grid.size() - 1)) { long long time_diff = 0; - long long last_clock = wall_clock64(); + long long last_clock = clock_function(); do { - long long cur_clock = wall_clock64(); + long long cur_clock = clock_function(); if (cur_clock > last_clock) { time_diff += (cur_clock - last_clock); } diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc index 33bf5b7849..2556eb8900 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc @@ -251,9 +251,9 @@ __global__ void test_kernel_gfx11(unsigned int* atomic_val, unsigned int* global // reach the atomicInc, but everyone will have only hit the atomic once. if (rank == (grid.size() - 1)) { long long time_diff = 0; - long long last_clock = wall_clock64(); + long long last_clock = clock_function(); do { - long long cur_clock = wall_clock64(); + long long cur_clock = clock_function(); if (cur_clock > last_clock) { time_diff += (cur_clock - last_clock); } @@ -274,9 +274,9 @@ __global__ void test_kernel_gfx11(unsigned int* atomic_val, unsigned int* global // and multiplies will not be aligned between to the two GPUs. if (global_rank == (mgrid.size() - 1)) { long long time_diff = 0; - long long last_clock = wall_clock64(); + long long last_clock = clock_function(); do { - long long cur_clock = wall_clock64(); + long long cur_clock = clock_function(); if (cur_clock > last_clock) { time_diff += (cur_clock - last_clock); } diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc index 5cde00aa34..afc63c9de6 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc @@ -165,9 +165,9 @@ __global__ void test_coop_kernel_gfx11(unsigned int loops, long long* array, int for (int i = 0; i < loops; i++) { long long time_diff = 0; - long long last_clock = wall_clock64(); + long long last_clock = clock_function(); do { - long long cur_clock = wall_clock64(); + long long cur_clock = clock_function(); if (cur_clock > last_clock) { time_diff += (cur_clock - last_clock); } @@ -175,7 +175,7 @@ __global__ void test_coop_kernel_gfx11(unsigned int loops, long long* array, int // So just ignore those slipped cycles. last_clock = cur_clock; } while (time_diff < 1000000); - array[rank] += wall_clock64(); + array[rank] += clock_function(); } #endif } @@ -205,9 +205,9 @@ __global__ void test_kernel_gfx11(uint32_t loops, unsigned long long* array) { for (int i = 0; i < loops; i++) { long long time_diff = 0; - long long last_clock = wall_clock64(); + long long last_clock = clock_function(); do { - long long cur_clock = wall_clock64(); + long long cur_clock = clock_function(); if (cur_clock > last_clock) { time_diff += (cur_clock - last_clock); } @@ -215,7 +215,7 @@ __global__ void test_kernel_gfx11(uint32_t loops, unsigned long long* array) { // So just ignore those slipped cycles. last_clock = cur_clock; } while (time_diff < 1000000); - array[rank] += wall_clock64(); + array[rank] += clock_function(); } #endif } diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc index f0a8b6cd17..da9b6ad6a2 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc @@ -126,11 +126,11 @@ __global__ void kernel500ms_gfx11(float* hostRes, int clkRate) { hostRes[tid] = tid + 1; __threadfence_system(); // expecting that the data is getting flushed to host here! - uint64_t start = wall_clock64()/clkRate, cur; + uint64_t start = clock_function()/clkRate, cur; if (clkRate > 1) { - do { cur = wall_clock64()/clkRate-start;}while (cur < wait_ms); + do { cur = clock_function()/clkRate-start;}while (cur < wait_ms); } else { - do { cur = wall_clock64()/start;}while (cur < wait_ms); + do { cur = clock_function()/start;}while (cur < wait_ms); } #endif } diff --git a/projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc b/projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc index e9fed1fc84..8ba85043ff 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc @@ -47,11 +47,11 @@ __global__ void Kernel_gfx11(float* hostRes, int clkRate) { hostRes[tid] = tid + 1; __threadfence_system(); // expecting that the data is getting flushed to host here! - uint64_t start = wall_clock64()/clkRate, cur; + uint64_t start = clock_function()/clkRate, cur; if (clkRate > 1) { - do { cur = wall_clock64()/clkRate-start;}while (cur < wait_sec); + do { cur = clock_function()/clkRate-start;}while (cur < wait_sec); } else { - do { cur = wall_clock64()/start;}while (cur < wait_sec); + do { cur = clock_function()/start;}while (cur < wait_sec); } #endif } diff --git a/projects/hip-tests/catch/unit/memory/mempool_common.hh b/projects/hip-tests/catch/unit/memory/mempool_common.hh index 7e255984ec..21479005a4 100644 --- a/projects/hip-tests/catch/unit/memory/mempool_common.hh +++ b/projects/hip-tests/catch/unit/memory/mempool_common.hh @@ -81,14 +81,14 @@ template __global__ void kernel_500ms_gfx11(T* host_res, int clk_ra host_res[tid] = tid + 1; __threadfence_system(); // expecting that the data is getting flushed to host here! - uint64_t start = wall_clock64() / clk_rate, cur; + uint64_t start = clock_function() / clk_rate, cur; if (clk_rate > 1) { do { - cur = wall_clock64() / clk_rate - start; + cur = clock_function() / clk_rate - start; } while (cur < wait_ms); } else { do { - cur = wall_clock64() / start; + cur = clock_function() / start; } while (cur < wait_ms); } #endif diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index 0b432ee1b5..f940bf1766 100644 --- a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -104,10 +104,10 @@ __global__ void StreamPerThrd_gfx11(int *Ad, int *Ad1, size_t n, int Pk_Clk, } if (Wait) { int64_t GpuFrq = (Pk_Clk * 1000); - int64_t StrtTck = wall_clock64(); + int64_t StrtTck = clock_function(); if (index == 0) { // The following while loop checks the value in ptr for around 4 seconds - while ((wall_clock64() - StrtTck) <= (6 * GpuFrq)) { + while ((clock_function() - StrtTck) <= (6 * GpuFrq)) { } if (WaitEvnt == 1) { *Ad1 = 1; @@ -120,9 +120,9 @@ __global__ void StreamPerThrd_gfx11(int *Ad, int *Ad1, size_t n, int Pk_Clk, __global__ void StreamPerThrd1_gfx11(int *A, int Pk_Clk) { #if HT_AMD int64_t GpuFrq = (Pk_Clk * 1000); - int64_t StrtTck = wall_clock64(); + int64_t StrtTck = clock_function(); // The following while loop checks the value in ptr for around 3-4 seconds - while ((wall_clock64() - StrtTck) <= (3 * GpuFrq)) { + while ((clock_function() - StrtTck) <= (3 * GpuFrq)) { } *A = 1; #endif