Replace usages of __ockl_clz with builtins (#2234)

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) {

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);