Guard rcp rounded implementation as well

Since rcp implementations of non-default rounded versions are not correct or supported in OCML, guard them using the same macro OCML_BASIC_ROUNDED_OPERATIONS. Also update the docs and tests.


[ROCm/clr commit: 7b3bbc85c5]

projects/clr/hipamd/docs/markdown/hip-math-api.md

@@ -1569,7 +1569,7 @@ __device__ static  float __fmul_rz(float x, float y);
 __device__  float __frcp_rd(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_rn
@@ -1577,7 +1577,7 @@ __device__  float __frcp_rd(float x);
 __device__  float __frcp_rn(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_ru
@@ -1585,7 +1585,7 @@ __device__  float __frcp_rn(float x);
 __device__  float __frcp_ru(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frcp_rz
@@ -1593,7 +1593,7 @@ __device__  float __frcp_ru(float x);
 __device__  float __frcp_rz(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __frsqrt_rn
@@ -1601,7 +1601,7 @@ __device__  float __frcp_rz(float x);
 __device__  float __frsqrt_rn(float x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __fsqrt_rd
@@ -1833,7 +1833,7 @@ __device__ static  double __dmul_rz(double x, double y);
 __device__  double __drcp_rd(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_rn
@@ -1841,7 +1841,7 @@ __device__  double __drcp_rd(double x);
 __device__  double __drcp_rn(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_ru
@@ -1849,7 +1849,7 @@ __device__  double __drcp_rn(double x);
 __device__  double __drcp_ru(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __drcp_rz
@@ -1857,7 +1857,7 @@ __device__  double __drcp_ru(double x);
 __device__  double __drcp_rz(double x);
 
 ```
-**Description:**  Supported
+**Description:**  Unsupported
 
 
 ### __dsqrt_rd

projects/clr/hipamd/include/hip/hcc_detail/math_functions.h

@@ -580,7 +580,6 @@ float __fmul_ru(float x, float y)  { return __ocml_mul_rtp_f32(x, y); }
 __DEVICE__
 inline
 float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); }
-#endif
 __DEVICE__
 inline
 float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); }
@@ -596,7 +595,6 @@ float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); }
 __DEVICE__
 inline
 float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); }
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
@@ -1077,7 +1075,6 @@ double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
 __DEVICE__
 inline
 double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); }
-#endif
 __DEVICE__
 inline
 double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); }
@@ -1090,7 +1087,6 @@ double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); }
 __DEVICE__
 inline
 double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); }
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); }

projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp

@@ -47,12 +47,10 @@ __device__ void double_precision_intrinsics() {
     __dmul_rn(1.0, 2.0);
     __dmul_ru(1.0, 2.0);
     __dmul_rz(1.0, 2.0);
-#endif
     __drcp_rd(2.0);
     __drcp_rn(2.0);
     __drcp_ru(2.0);
     __drcp_rz(2.0);
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dsqrt_rd(4.0);
     __dsqrt_rn(4.0);
     __dsqrt_ru(4.0);

projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp

@@ -38,11 +38,13 @@ __global__ void floatMath(float* In, float* Out) {
     Out[tid] = __cosf(In[tid]);
     Out[tid] = __exp10f(Out[tid]);
     Out[tid] = __expf(Out[tid]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __frsqrt_rn(Out[tid]);
-    //Out[tid] = __fsqrt_rd(Out[tid]);
-    //Out[tid] = __fsqrt_rn(Out[tid]);
-    //Out[tid] = __fsqrt_ru(Out[tid]);
-    //Out[tid] = __fsqrt_rz(Out[tid]);
+    Out[tid] = __fsqrt_rd(Out[tid]);
+    Out[tid] = __fsqrt_rn(Out[tid]);
+    Out[tid] = __fsqrt_ru(Out[tid]);
+    Out[tid] = __fsqrt_rz(Out[tid]);
+#endif
     Out[tid] = __log10f(Out[tid]);
     Out[tid] = __log2f(Out[tid]);
     Out[tid] = __logf(Out[tid]);

projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp

@@ -59,13 +59,11 @@ __device__ void single_precision_intrinsics() {
     __fmul_rn(1.0f, 2.0f);
     __fmul_ru(1.0f, 2.0f);
     __fmul_rz(1.0f, 2.0f);
-#endif
     __frcp_rd(2.0f);
     __frcp_rn(2.0f);
     __frcp_ru(2.0f);
     __frcp_rz(2.0f);
     __frsqrt_rn(4.0f);
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsqrt_rd(4.0f);
     __fsqrt_rn(4.0f);
     __fsqrt_ru(4.0f);