diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h --- a/clang/lib/Headers/__clang_hip_libdevice_declares.h +++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -138,14 +138,22 @@ __device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); -__device__ __attribute__((const)) float -__llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); -__device__ __attribute__((const)) float -__llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32"); -__device__ __attribute__((const)) float -__llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32"); -__device__ __attribute__((const)) float -__llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32"); +__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 @@ -269,10 +277,14 @@ __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, double); -__device__ __attribute__((const)) double -__llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); -__device__ __attribute__((const)) double -__llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); +__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);