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 b748835fa9..8742bf10dd 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 @@ -52,10 +52,12 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) { return __builtin_popcountll(input); } -__device__ static inline int __clz(int input) { return __ockl_clz_u32((uint)input); } +__device__ static inline int __clz(int input) { + return input == 0u ? 32 : __builtin_clz((uint)input); +} __device__ static inline int __clzll(long long int input) { - return __ockl_clz_u64((__hip_uint64_t)input); + return input == 0u ? 64 : __builtin_clzl((__hip_uint64_t)input); } __device__ static inline int __ffs(unsigned int input) { diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h index e1dc54f270..39c2e59686 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h @@ -56,11 +56,6 @@ extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint); extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int); extern "C" __device__ __attribute__((const)) uint __ockl_sadd_u32(uint, uint, uint); -extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar); -extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort); -extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint); -extern "C" __device__ __attribute__((const)) __hip_uint64_t __ockl_clz_u64(__hip_uint64_t); - extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float); extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);