From 320334155157e8800ac59ffa459f6ea1bbb78566 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Wed, 23 Nov 2016 20:01:18 -0600
Subject: [PATCH] added support for rcp for float and double

Change-Id: Ibeba3a9f64494fc0a176bcb4a854fb2f56567b55


[ROCm/hip commit: 7131d0b96145f9f4543da31fec901853ddc3c72b]
---
 .../hip/include/hip/hcc_detail/hip_runtime.h  | 34 +++++++++++++++++++
 .../hipDoublePrecisionIntrinsics.cpp          |  8 ++---
 2 files changed, 38 insertions(+), 4 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h
index 63cfb2ea3c..b1877ed0b3 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h
@@ -510,9 +510,11 @@ __device__ void __hip_fast_sincosf(float,float*,float*);
 extern __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
 __device__ float __hip_fast_tanf(float);
 extern __attribute__((const)) float __hip_fast_fmaf(float,float,float) __asm("llvm.fma.f32");
+extern __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
 
 extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
 extern __attribute__((const)) double __hip_fast_fma(double,double,double) __asm("llvm.fma.f64");
+extern __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
 
 #ifdef HIP_FAST_MATH
 // Single Precision Precise Math when enabled
@@ -648,6 +650,22 @@ __device__ inline float __fmaf_rz(float x, float y, float z) {
   return __hip_fast_fmaf(x, y, z);
 }
 
+__device__ inline float __frcp_rd(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_rn(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_ru(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_rz(float x) {
+  return __hip_fast_frcp(x);
+}
+
 __device__ inline double __dsqrt_rd(double x) {
   return __hip_fast_dsqrt(x);
 }
@@ -680,6 +698,22 @@ __device__ inline double __fma_rz(double x, double y, double z) {
   return __hip_fast_fma(x, y, z);
 }
 
+__device__ inline double __drcp_rd(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_rn(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_ru(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_rz(double x) {
+  return __hip_fast_drcp(x);
+}
+
 /**
  * CUDA 8 device function features
 
diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
index 0b4e0840a4..330c4fc799 100644
--- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
@@ -39,10 +39,10 @@ __device__ void double_precision_intrinsics()
     //__dmul_rn(1.0, 2.0);
     //__dmul_ru(1.0, 2.0);
     //__dmul_rz(1.0, 2.0);
-    //__drcp_rd(2.0);
-    //__drcp_rn(2.0);
-    //__drcp_ru(2.0);
-    //__drcp_rz(2.0);
+    __drcp_rd(2.0);
+    __drcp_rn(2.0);
+    __drcp_ru(2.0);
+    __drcp_rz(2.0);
     __dsqrt_rd(4.0);
     __dsqrt_rn(4.0);
     __dsqrt_ru(4.0);