diff --git a/docs/markdown/hip-math-api.md b/docs/markdown/hip-math-api.md index b3698ff2b3..9b8a3f2f11 100644 --- a/docs/markdown/hip-math-api.md +++ b/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 diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 08be321d68..557257b2b0 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/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); } diff --git a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 295fd83708..939bdae743 100644 --- a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/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); diff --git a/tests/src/deviceLib/hipFloatMath.cpp b/tests/src/deviceLib/hipFloatMath.cpp index ee83309f28..c6a07e26a9 100644 --- a/tests/src/deviceLib/hipFloatMath.cpp +++ b/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]); diff --git a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index db60099558..b216b3cb54 100644 --- a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/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);