From a72344ad676e00f8f529dd67d076de9ab29b41b5 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Thu, 6 Oct 2022 13:32:16 -0400 Subject: [PATCH] SWDEV-337452 - Adding__ockl_steadyctr_u64 option for wall_clock64. Change-Id: I4ee6fd4d40d500363ef84975078185093dab5ddd [ROCm/clr commit: 5ff27faa042b872a0d02048fc65299aee18ad421] --- .../hipamd/include/hip/amd_detail/amd_device_functions.h | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index 96353b0f54..ab1e1f9e97 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -42,6 +42,8 @@ template static inline __device__ void printf(const char* format, All... all) {} #endif // __HIP_CLANG_ONLY__ +extern "C" __device__ unsigned long long __ockl_steadyctr_u64(); + /* Integer Intrinsics */ @@ -719,12 +721,7 @@ long long int __clock() { return __clock64(); } __device__ inline __attribute__((always_inline)) long long int wall_clock64() { -#if __has_builtin(__builtin_amdgcn_s_memrealtime) - // Exists since gfx8 - return (long long int) __builtin_amdgcn_s_memrealtime(); -#else - return -1; // Negative return means __builtin_amdgcn_s_memrealtime unavailable. -#endif + return (long long int) __ockl_steadyctr_u64(); } __device__