diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index 794efb8a41..263834877a 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -656,12 +656,17 @@ __device__ void __named_sync(); __device__ inline __attribute((always_inline)) long long int __clock64() { -#if __has_builtin(__builtin_amdgcn_s_memtime) - // Exists on gfx8, gfx9, gfx10.1, gfx10.2, gfx10.3 - return (long long int) __builtin_amdgcn_s_memtime(); -#else - // Subject to change when better solution available - return (long long int) __builtin_readcyclecounter(); +#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 }