From 1aae121e864b7c19be35ccb4bfb393fddbeff0f7 Mon Sep 17 00:00:00 2001 From: "Sang, Tao" Date: Tue, 22 Jul 2025 01:11:23 -0400 Subject: [PATCH] SWDEV-533964 - use __builtin_readcyclecounter() only in clock64 (#368) [ROCm/clr commit: 5fc6208b197f7bf51f049f4e329c98139ecbdcb3] --- .../include/hip/amd_detail/amd_device_functions.h | 13 +------------ 1 file changed, 1 insertion(+), 12 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 38731a9548..07f3debd6b 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 @@ -656,18 +656,7 @@ __device__ void __named_sync(); __device__ inline __attribute((always_inline)) long long int __clock64() { -#if __has_builtin(__builtin_amdgcn_is_invocable) // ZCFS - if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_memtime)) - return (long long int)__builtin_amdgcn_s_memtime(); - else - return (long long int)__builtin_readcyclecounter(); -#else // LEGACY BRANCH FOR COMPAT - #if __has_builtin(__builtin_amdgcn_s_memtime) && !defined(__SPIRV__) - return (long long int)__builtin_amdgcn_s_memtime(); - #else - return (long long int)__builtin_readcyclecounter(); - #endif -#endif + return (long long int)__builtin_readcyclecounter(); } __device__