Add *_rn functions back into HIP intrinsics

Add back the round-to-nearest-even intrinsics back to HIP math
intrinsics as it caused regression.

tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp

@@ -36,31 +36,45 @@ THE SOFTWARE.
 __device__ void double_precision_intrinsics() {
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dadd_rd(0.0, 1.0);
+#endif
     __dadd_rn(0.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dadd_ru(0.0, 1.0);
     __dadd_rz(0.0, 1.0);
     __ddiv_rd(0.0, 1.0);
+#endif
     __ddiv_rn(0.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __ddiv_ru(0.0, 1.0);
     __ddiv_rz(0.0, 1.0);
     __dmul_rd(1.0, 2.0);
+#endif
     __dmul_rn(1.0, 2.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dmul_ru(1.0, 2.0);
     __dmul_rz(1.0, 2.0);
     __drcp_rd(2.0);
+#endif
     __drcp_rn(2.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __drcp_ru(2.0);
     __drcp_rz(2.0);
     __dsqrt_rd(4.0);
+#endif
     __dsqrt_rn(4.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dsqrt_ru(4.0);
     __dsqrt_rz(4.0);
     __dsub_rd(2.0, 1.0);
+#endif
     __dsub_rn(2.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dsub_ru(2.0, 1.0);
     __dsub_rz(2.0, 1.0);
     __fma_rd(1.0, 2.0, 3.0);
+#endif
     __fma_rn(1.0, 2.0, 3.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fma_ru(1.0, 2.0, 3.0);
     __fma_rz(1.0, 2.0, 3.0);
 #endif

tests/src/deviceLib/hipFloatMath.cpp

@@ -38,10 +38,12 @@ __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]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __fsqrt_rd(Out[tid]);
+#endif
     Out[tid] = __fsqrt_rn(Out[tid]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __fsqrt_ru(Out[tid]);
     Out[tid] = __fsqrt_rz(Out[tid]);
 #endif

tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp

@@ -41,35 +41,51 @@ __device__ void single_precision_intrinsics() {
     __expf(0.0f);
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fadd_rd(0.0f, 1.0f);
+#endif
     __fadd_rn(0.0f, 1.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fadd_ru(0.0f, 1.0f);
     __fadd_rz(0.0f, 1.0f);
     __fdiv_rd(4.0f, 2.0f);
+#endif
     __fdiv_rn(4.0f, 2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fdiv_ru(4.0f, 2.0f);
     __fdiv_rz(4.0f, 2.0f);
 #endif
     __fdividef(4.0f, 2.0f);
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmaf_rd(1.0f, 2.0f, 3.0f);
+#endif
     __fmaf_rn(1.0f, 2.0f, 3.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmaf_ru(1.0f, 2.0f, 3.0f);
     __fmaf_rz(1.0f, 2.0f, 3.0f);
     __fmul_rd(1.0f, 2.0f);
+#endif
     __fmul_rn(1.0f, 2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmul_ru(1.0f, 2.0f);
     __fmul_rz(1.0f, 2.0f);
     __frcp_rd(2.0f);
+#endif
     __frcp_rn(2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __frcp_ru(2.0f);
     __frcp_rz(2.0f);
+#endif
     __frsqrt_rn(4.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsqrt_rd(4.0f);
+#endif
     __fsqrt_rn(4.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsqrt_ru(4.0f);
     __fsqrt_rz(4.0f);
     __fsub_rd(2.0f, 1.0f);
+#endif
     __fsub_rn(2.0f, 1.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsub_ru(2.0f, 1.0f);
     __fsub_rz(2.0f, 1.0f);
 #endif