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

Add clock() and clock64()

Browse Source 
[ROCm/hip commit: 1f74a66bea]
This commit is contained in:
Yaxun Sam Liu
2018-08-28 19:30:20 -04:00
parent ae41c374fe
commit c1ab349fc7
2 changed files with 101 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/device_functions.h
+32 -5
View File
@@ -611,16 +611,43 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) { return
#ifdef __HCC_OR_HIP_CLANG__
// Clock functions
__device__ long long int __clock64();
__device__ long long int __clock();
__device__ long long int clock64();
__device__ long long int clock();
// hip.amdgcn.bc - named sync
__device__ void __named_sync(int a, int b);
#ifdef __HIP_DEVICE_COMPILE__
// Clock functions
__device__
inline
long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
#if __HCC__
extern "C" uint64_t __clock_u64() __HC__;
#endif
__device__
inline
long long int __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
inline __attribute((always_inline))
long long int __clock64() {
// ToDo: Unify HCC and HIP implementation.
#if __HCC__
return (long long int) __clock_u64();
#else
return (long long int) __builtin_amdgcn_s_memrealtime();
#endif
}
__device__
inline __attribute((always_inline))
long long int __clock() { return __clock64(); }
__device__
inline __attribute__((always_inline))
long long int clock64() { return __clock64(); }
__device__
inline __attribute__((always_inline))
long long int clock() { return __clock(); }
// hip.amdgcn.bc - named sync
__device__