From 5e61121d880ac48e9880414797cade682120c157 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Fri, 23 Dec 2022 10:22:13 -0500 Subject: [PATCH] SWDEV-1 - Use proper __builtin_amdgcn_rcph builtin Change-Id: Ic0b1b25a5ab7d132cb51425f2a17bc779c762668 [ROCm/clr commit: eacf63673b0fcf15aa7c8c2f298968ff26039179] --- projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h | 6 +++--- .../clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h | 1 - 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 694deb228c..fda272191c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1569,7 +1569,7 @@ THE SOFTWARE. __half hrcp(__half x) { return __half_raw{ - __llvm_amdgcn_rcp_f16(static_cast<__half_raw>(x).data)}; + __builtin_amdgcn_rcph(static_cast<__half_raw>(x).data)}; } inline __device__ @@ -1673,8 +1673,8 @@ THE SOFTWARE. inline __HOST_DEVICE__ __half2 h2rcp(__half2 x) { - return _Float16_2{__llvm_amdgcn_rcp_f16(x.x), - __llvm_amdgcn_rcp_f16(x.y)}; + return _Float16_2{__builtin_amdgcn_rcph(x.x), + __builtin_amdgcn_rcph(x.y)}; } inline __HOST_DEVICE__ diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h index aac0bcca53..caf6ec752c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h @@ -44,7 +44,6 @@ extern "C" __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16); __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16); __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int); - __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);