From 78f3caba0ba5ecfcc3c218201b125aa95e0fada1 Mon Sep 17 00:00:00 2001 From: "Sang, Tao" Date: Tue, 20 May 2025 16:57:17 -0400 Subject: [PATCH] SWDEV-519346 - Fix __clock64() compiling issue in SPIRV (#207) --- .../hip/amd_detail/amd_device_functions.h | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) 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 }