SWDEV-1 - Use proper __builtin_amdgcn_rcph builtin

Change-Id: Ic0b1b25a5ab7d132cb51425f2a17bc779c762668

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__

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