Remove non-working non-default-rounded math apis

In ROCm-Device-Libs, they have dropped the non-default-rounded versions of add, sub, mul, div, sqrt and fma. Therefore, ocml has removed the rte, rtp, rtn, and rtz counterparts. This will remove the same math APIs in HIP for _ru, _rd, _rn, and _rz.


[ROCm/hip commit: 2b108a2cfd]
This commit is contained in:
Aaron Enye Shi
2018-11-05 22:34:16 +00:00
parent 2020c337c2
commit 1c6f47ef55
4 changed files with 265 additions and 217 deletions
@@ -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) {
@@ -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);