diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index b12e7aca89..8a6091858b 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -514,69 +514,69 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); } __DEVICE__ inline float __expf(float x) { return __ocml_exp_f32(x); } -__DEVICE__ -inline -float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); } -__DEVICE__ -inline -float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } -__DEVICE__ -inline -float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); } -__DEVICE__ -inline -float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } -__DEVICE__ -inline -float __fdiv_rd(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_rn(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_ru(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fdiv_rz(float x, float y) { return x / y; } +// __DEVICE__ +// inline +// float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } __DEVICE__ inline float __fdividef(float x, float y) { return x / y; } -__DEVICE__ -inline -float __fmaf_rd(float x, float y, float z) -{ - return __ocml_fma_rtp_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_rn(float x, float y, float z) -{ - return __ocml_fma_rte_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_ru(float x, float y, float z) -{ - return __ocml_fma_rtn_f32(x, y, z); -} -__DEVICE__ -inline -float __fmaf_rz(float x, float y, float z) -{ - return __ocml_fma_rtz_f32(x, y, z); -} -__DEVICE__ -inline -float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); } -__DEVICE__ -inline -float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } -__DEVICE__ -inline -float __fmul_ru(float x, float y) { return __ocml_mul_rtn_f32(x, y); } -__DEVICE__ -inline -float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fmaf_rd(float x, float y, float z) +// { +// return __ocml_fma_rtn_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_rn(float x, float y, float z) +// { +// return __ocml_fma_rte_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_ru(float x, float y, float z) +// { +// return __ocml_fma_rtp_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmaf_rz(float x, float y, float z) +// { +// return __ocml_fma_rtz_f32(x, y, z); +// } +// __DEVICE__ +// inline +// float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } +// __DEVICE__ +// inline +// 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); } __DEVICE__ inline float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); } @@ -592,30 +592,30 @@ float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); } __DEVICE__ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } -__DEVICE__ -inline -float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); } -__DEVICE__ -inline -float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); } -__DEVICE__ -inline -float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } -__DEVICE__ -inline -float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); } -__DEVICE__ -inline -float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } +// __DEVICE__ +// inline +// float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); } +// __DEVICE__ +// inline +// float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } +// __DEVICE__ +// inline +// float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); } +// __DEVICE__ +// inline +// float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } __DEVICE__ inline float __log10f(float x) { return __ocml_log10_f32(x); } @@ -1034,42 +1034,42 @@ double yn(int n, double x) } // BEGIN INTRINSICS -__DEVICE__ -inline -double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); } -__DEVICE__ -inline -double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } -__DEVICE__ -inline -double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); } -__DEVICE__ -inline -double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } -__DEVICE__ -inline -double __ddiv_rd(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_rn(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_ru(double x, double y) { return x / y; } -__DEVICE__ -inline -double __ddiv_rz(double x, double y) { return x / y; } -__DEVICE__ -inline -double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); } -__DEVICE__ -inline -double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } -__DEVICE__ -inline -double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); } -__DEVICE__ -inline -double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } +// __DEVICE__ +// inline +// 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); } __DEVICE__ inline double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); } @@ -1082,54 +1082,54 @@ double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); } __DEVICE__ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } -__DEVICE__ -inline -double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); } -__DEVICE__ -inline -double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); } -__DEVICE__ -inline -double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } -__DEVICE__ -inline -double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); } -__DEVICE__ -inline -double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } -__DEVICE__ -inline -double __fma_rd(double x, double y, double z) -{ - return __ocml_fma_rtp_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_rn(double x, double y, double z) -{ - return __ocml_fma_rte_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_ru(double x, double y, double z) -{ - return __ocml_fma_rtn_f64(x, y, z); -} -__DEVICE__ -inline -double __fma_rz(double x, double y, double z) -{ - return __ocml_fma_rtz_f64(x, y, z); -} +// __DEVICE__ +// inline +// double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); } +// __DEVICE__ +// inline +// double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } +// __DEVICE__ +// inline +// double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); } +// __DEVICE__ +// inline +// double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); } +// __DEVICE__ +// inline +// double __fma_rd(double x, double y, double z) +// { +// return __ocml_fma_rtn_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_rn(double x, double y, double z) +// { +// return __ocml_fma_rte_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_ru(double x, double y, double z) +// { +// return __ocml_fma_rtp_f64(x, y, z); +// } +// __DEVICE__ +// inline +// double __fma_rz(double x, double y, double z) +// { +// return __ocml_fma_rtz_f64(x, y, z); +// } // END INTRINSICS // END DOUBLE diff --git a/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h index 404c2f81d5..e5594924ba 100644 --- a/include/hip/hcc_detail/math_fwd.h +++ b/include/hip/hcc_detail/math_fwd.h @@ -288,6 +288,30 @@ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); __device__ __attribute__((const)) +float __ocml_div_rte_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtn_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtp_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_div_rtz_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rte_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtn_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtp_f32(float, float); +__device__ +__attribute__((const)) +float __ocml_sqrt_rtz_f32(float, float); +__device__ +__attribute__((const)) float __ocml_fma_rte_f32(float, float, float); __device__ __attribute__((const)) @@ -572,6 +596,30 @@ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); __device__ __attribute__((const)) +double __ocml_div_rte_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtn_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtp_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_div_rtz_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rte_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtn_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtp_f64(double, double); +__device__ +__attribute__((const)) +double __ocml_sqrt_rtz_f64(double, double); +__device__ +__attribute__((const)) double __ocml_fma_rte_f64(double, double, double); __device__ __attribute__((const)) @@ -594,4 +642,4 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); #if defined(__cplusplus) } // extern "C" -#endif \ No newline at end of file +#endif diff --git a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 17cd82c9ab..f6c515c03a 100644 --- a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -34,34 +34,34 @@ THE SOFTWARE. #pragma clang diagnostic ignored "-Wunused-variable" __device__ void double_precision_intrinsics() { - __dadd_rd(0.0, 1.0); - __dadd_rn(0.0, 1.0); - __dadd_ru(0.0, 1.0); - __dadd_rz(0.0, 1.0); - __ddiv_rd(0.0, 1.0); - __ddiv_rn(0.0, 1.0); - __ddiv_ru(0.0, 1.0); - __ddiv_rz(0.0, 1.0); - __dmul_rd(1.0, 2.0); - __dmul_rn(1.0, 2.0); - __dmul_ru(1.0, 2.0); - __dmul_rz(1.0, 2.0); + // __dadd_rd(0.0, 1.0); + // __dadd_rn(0.0, 1.0); + // __dadd_ru(0.0, 1.0); + // __dadd_rz(0.0, 1.0); + // __ddiv_rd(0.0, 1.0); + // __ddiv_rn(0.0, 1.0); + // __ddiv_ru(0.0, 1.0); + // __ddiv_rz(0.0, 1.0); + // __dmul_rd(1.0, 2.0); + // __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); - __dsqrt_rd(4.0); - __dsqrt_rn(4.0); - __dsqrt_ru(4.0); - __dsqrt_rz(4.0); - __dsub_rd(2.0, 1.0); - __dsub_rn(2.0, 1.0); - __dsub_ru(2.0, 1.0); - __dsub_rz(2.0, 1.0); - __fma_rd(1.0, 2.0, 3.0); - __fma_rn(1.0, 2.0, 3.0); - __fma_ru(1.0, 2.0, 3.0); - __fma_rz(1.0, 2.0, 3.0); + // __dsqrt_rd(4.0); + // __dsqrt_rn(4.0); + // __dsqrt_ru(4.0); + // __dsqrt_rz(4.0); + // __dsub_rd(2.0, 1.0); + // __dsub_rn(2.0, 1.0); + // __dsub_ru(2.0, 1.0); + // __dsub_rz(2.0, 1.0); + // __fma_rd(1.0, 2.0, 3.0); + // __fma_rn(1.0, 2.0, 3.0); + // __fma_ru(1.0, 2.0, 3.0); + // __fma_rz(1.0, 2.0, 3.0); } __global__ void compileDoublePrecisionIntrinsics(int ignored) { diff --git a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index f3d2a36931..623ea08a94 100644 --- a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -39,36 +39,36 @@ __device__ void single_precision_intrinsics() { __cosf(0.0f); __exp10f(0.0f); __expf(0.0f); - __fadd_rd(0.0f, 1.0f); - __fadd_rn(0.0f, 1.0f); - __fadd_ru(0.0f, 1.0f); - __fadd_rz(0.0f, 1.0f); - __fdiv_rd(4.0f, 2.0f); - __fdiv_rn(4.0f, 2.0f); - __fdiv_ru(4.0f, 2.0f); - __fdiv_rz(4.0f, 2.0f); + // __fadd_rd(0.0f, 1.0f); + // __fadd_rn(0.0f, 1.0f); + // __fadd_ru(0.0f, 1.0f); + // __fadd_rz(0.0f, 1.0f); + // __fdiv_rd(4.0f, 2.0f); + // __fdiv_rn(4.0f, 2.0f); + // __fdiv_ru(4.0f, 2.0f); + // __fdiv_rz(4.0f, 2.0f); __fdividef(4.0f, 2.0f); - __fmaf_rd(1.0f, 2.0f, 3.0f); - __fmaf_rn(1.0f, 2.0f, 3.0f); - __fmaf_ru(1.0f, 2.0f, 3.0f); - __fmaf_rz(1.0f, 2.0f, 3.0f); - __fmul_rd(1.0f, 2.0f); - __fmul_rn(1.0f, 2.0f); - __fmul_ru(1.0f, 2.0f); - __fmul_rz(1.0f, 2.0f); + // __fmaf_rd(1.0f, 2.0f, 3.0f); + // __fmaf_rn(1.0f, 2.0f, 3.0f); + // __fmaf_ru(1.0f, 2.0f, 3.0f); + // __fmaf_rz(1.0f, 2.0f, 3.0f); + // __fmul_rd(1.0f, 2.0f); + // __fmul_rn(1.0f, 2.0f); + // __fmul_ru(1.0f, 2.0f); + // __fmul_rz(1.0f, 2.0f); __frcp_rd(2.0f); __frcp_rn(2.0f); __frcp_ru(2.0f); __frcp_rz(2.0f); __frsqrt_rn(4.0f); - __fsqrt_rd(4.0f); - __fsqrt_rn(4.0f); - __fsqrt_ru(4.0f); - __fsqrt_rz(4.0f); - __fsub_rd(2.0f, 1.0f); - __fsub_rn(2.0f, 1.0f); - __fsub_ru(2.0f, 1.0f); - __fsub_rz(2.0f, 1.0f); + // __fsqrt_rd(4.0f); + // __fsqrt_rn(4.0f); + // __fsqrt_ru(4.0f); + // __fsqrt_rz(4.0f); + // __fsub_rd(2.0f, 1.0f); + // __fsub_rn(2.0f, 1.0f); + // __fsub_ru(2.0f, 1.0f); + // __fsub_rz(2.0f, 1.0f); __log10f(1.0f); __log2f(1.0f); __logf(1.0f);