Index: clang/lib/Headers/__clang_hip_libdevice_declares.h =================================================================== --- clang/lib/Headers/__clang_hip_libdevice_declares.h +++ clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -137,23 +137,6 @@ __device__ __attribute__((const)) float __ocml_fma_rtn_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); - -__device__ inline __attribute__((const)) float -__llvm_amdgcn_cos_f32(float __x) { - return __builtin_amdgcn_cosf(__x); -} -__device__ inline __attribute__((const)) float -__llvm_amdgcn_rcp_f32(float __x) { - return __builtin_amdgcn_rcpf(__x); -} -__device__ inline __attribute__((const)) float -__llvm_amdgcn_rsq_f32(float __x) { - return __builtin_amdgcn_rsqf(__x); -} -__device__ inline __attribute__((const)) float -__llvm_amdgcn_sin_f32(float __x) { - return __builtin_amdgcn_sinf(__x); -} // END INTRINSICS // END FLOAT @@ -277,15 +260,6 @@ __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, double); -__device__ inline __attribute__((const)) double -__llvm_amdgcn_rcp_f64(double __x) { - return __builtin_amdgcn_rcp(__x); -} -__device__ inline __attribute__((const)) double -__llvm_amdgcn_rsq_f64(double __x) { - return __builtin_amdgcn_rsq(__x); -} - __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); __device__ _Float16 __ocml_cos_f16(_Float16); __device__ __attribute__((const)) _Float16 __ocml_cvtrtn_f16_f32(float); @@ -305,7 +279,6 @@ __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16); __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16); __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16); -__device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16); __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16); __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16); __device__ _Float16 __ocml_sin_f16(_Float16); @@ -332,11 +305,6 @@ __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16); __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16); __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16); -__device__ inline __2f16 -__llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL. -{ - return (__2f16)(__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y)); -} __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16); __device__ __2f16 __ocml_sin_2f16(__2f16); Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -659,7 +659,7 @@ #endif __DEVICE__ -float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); } +float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__