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

SWDEV-450735 - Add compiler option for using clock64

Browse Source 
Change-Id: I9efed88d691ee1b2b4465286b3340b820f7bf627


[ROCm/hip-tests commit: 5b67a2a4a0]
This commit is contained in:
Vladana Stojiljkovic
2024-03-15 17:54:34 +01:00
parent 56440ded48
commit e057b9e471
11 changed files with 41 additions and 29 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/CMakeLists.txt
+6
View File
@@ -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}")
projects/hip-tests/catch/include/hip_test_common.hh
+6
View File
@@ -35,6 +35,12 @@ THE SOFTWARE.
#include <cstdlib>
#include <thread>
#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
projects/hip-tests/catch/include/utils.hh
+2 -2
View File
@@ -140,8 +140,8 @@ template <typename T> __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
projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+2 -2
View File
@@ -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
projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc
+2 -2
View File
@@ -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);
}
projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc
+4 -4
View File
@@ -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);
}
projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc
+6 -6
View File
@@ -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
}
projects/hip-tests/catch/unit/memory/hipMemPoolApi.cc
+3 -3
View File
@@ -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
}
projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc
+3 -3
View File
@@ -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
}
projects/hip-tests/catch/unit/memory/mempool_common.hh
+3 -3
View File
@@ -81,14 +81,14 @@ template <typename T> __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
projects/hip-tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc
+4 -4
View File
@@ -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