From 91d9ec75d75e0634abbcbc234c9b349a749d98bc Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 1 Jun 2018 16:20:21 +0100 Subject: [PATCH] Fix typos / address review comments. --- .../include/hip/hcc_detail/math_functions.h | 117 +++++++++--------- hipamd/include/hip/hcc_detail/ocml_math_fwd.h | 100 +++++++-------- 2 files changed, 110 insertions(+), 107 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 49650987ac..3072fc1d4e 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -341,7 +341,7 @@ float normcdfinvf(float x) { return __ocml_ncdfinv_f32(x); } __device__ inline float normf(int dim, const float* a) -{ +{ // TODO: placeholder until OCML adds support. float r = 0; while (dim--) { r += a[0] * a[0]; ++a; } @@ -389,7 +389,7 @@ float rnorm4df(float x, float y, float z, float w) __device__ inline float rnormf(int dim, const float* a) -{ +{ // TODO: placeholder until OCML adds support. float r = 0; while (dim--) { r += a[0] * a[0]; ++a; } @@ -468,7 +468,8 @@ inline float ynf(int n, float x) { // TODO: we could use Ahmes multiplication and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if - // it'd be beneficial in this case. + // it'd be beneficial in this case. Placeholder until OCML adds + // support. if (n == 0) return y0f(x); if (n == 1) return y1f(x); @@ -495,67 +496,67 @@ inline float __expf(float x) { return __ocml_exp_f32(x); } __device__ inline -float __fadd_rd(float x, float y) { return __llvm_add_rte_f32(x, y); } +float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); } __device__ inline -float __fadd_rn(float x, float y) { return __llvm_add_rtn_f32(x, y); } +float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); } __device__ inline -float __fadd_ru(float x, float y) { return __llvm_add_rtp_f32(x, y); } +float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); } __device__ inline -float __fadd_rz(float x, float y) { return __llvm_add_rtz_f32(x, y); } +float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } __device__ inline -float __fdiv_rd(float x, float y) { return __llvm_div_rte_f32(x, y); } +float __fdiv_rd(float x, float y) { return __ocml_div_rtp_f32(x, y); } __device__ inline -float __fdiv_rn(float x, float y) { return __llvm_div_rtn_f32(x, y); } +float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } __device__ inline -float __fdiv_ru(float x, float y) { return __llvm_div_rtp_f32(x, y); } +float __fdiv_ru(float x, float y) { return __ocml_div_rtn_f32(x, y); } __device__ inline -float __fdiv_rz(float x, float y) { return __llvm_div_rtz_f32(x, y); } +float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } __device__ inline -float __fdividef(float x, float y) { return __llvm_div_rte_f32(x, y); } +float __fdividef(float x, float y) { return __ocml_div_rte_f32(x, y); } __device__ inline float __fmaf_rd(float x, float y, float z) { - return __llvm_fma_rte_f32(x, y, z); + return __ocml_fma_rtp_f32(x, y, z); } __device__ inline float __fmaf_rn(float x, float y, float z) { - return __llvm_fma_rtn_f32(x, y, z); + return __ocml_fma_rte_f32(x, y, z); } __device__ inline float __fmaf_ru(float x, float y, float z) { - return __llvm_fma_rtp_f32(x, y, z); + return __ocml_fma_rtn_f32(x, y, z); } __device__ inline float __fmaf_rz(float x, float y, float z) { - return __llvm_fma_rtz_f32(x, y, z); + return __ocml_fma_rtz_f32(x, y, z); } __device__ inline -float __fmul_rd(float x, float y) { return __llvm_mul_rte_f32(x, y); } +float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); } __device__ inline -float __fmul_rn(float x, float y) { return __llvm_mul_rtn_f32(x, y); } +float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); } __device__ inline -float __fmul_ru(float x, float y) { return __llvm_mul_rtp_f32(x, y); } +float __fmul_ru(float x, float y) { return __ocml_mul_rtn_f32(x, y); } __device__ inline -float __fmul_rz(float x, float y) { return __llvm_mul_rtz_f32(x, y); } +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); } @@ -573,28 +574,28 @@ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } __device__ inline -float __fsqrt_rd(float x) { return __llvm_sqrt_rte_f32(x); } +float __fsqrt_rd(float x) { return __ocml_sqrt_rtp_f32(x); } __device__ inline -float __fsqrt_rn(float x) { return __llvm_sqrt_rtn_f32(x); } +float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } __device__ inline -float __fsqrt_ru(float x) { return __llvm_sqrt_rtp_f32(x); } +float __fsqrt_ru(float x) { return __ocml_sqrt_rtn_f32(x); } __device__ inline -float __fsqrt_rz(float x) { return __llvm_sqrt_rtz_f32(x); } +float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } __device__ inline -float __fsub_rd(float x, float y) { return __llvm_sub_rte_f32(x, y); } +float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); } __device__ inline -float __fsub_rn(float x, float y) { return __llvm_sub_rtn_f32(x, y); } +float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); } __device__ inline -float __fsub_ru(float x, float y) { return __llvm_sub_rtp_f32(x, y); } +float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); } __device__ inline -float __fsub_rz(float x, float y) { return __llvm_sub_rtz_f32(x, y); } +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); } @@ -617,7 +618,7 @@ void __sincosf(float x, float* sptr, float* cptr) float tmp; *sptr = - __ocml_sincos_f32(x, (__attribute__((address_space(5))) float*) &tmp); + __ocml_sincos_f32(x, (__attribute__((address_space(5))) float*) &tmp); *cptr = tmp; } __device__ @@ -763,7 +764,8 @@ inline double jn(int n, double x) { // TODO: we could use Ahmes multiplication and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if - // it'd be beneficial in this case. + // it'd be beneficial in this case. Placeholder until OCML adds + // support. if (n == 0) return j0f(x); if (n == 1) return j1f(x); @@ -853,7 +855,7 @@ double nextafter(double x, double y) { return __ocml_nextafter_f64(x, y); } __device__ inline double norm(int dim, const double* a) -{ +{ // TODO: placeholder until OCML adds support. double r = 0; while (dim--) { r += a[0] * a[0]; ++a; } @@ -906,7 +908,7 @@ double rint(double x) { return __ocml_rint_f64(x); } __device__ inline double rnorm(int dim, const double* a) -{ +{ // TODO: placeholder until OCML adds support. double r = 0; while (dim--) { r += a[0] * a[0]; ++a; } @@ -995,7 +997,8 @@ inline double yn(int n, double x) { // TODO: we could use Ahmes multiplication and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if - // it'd be beneficial in this case. + // it'd be beneficial in this case. Placeholder until OCML adds + // support. if (n == 0) return j0f(x); if (n == 1) return j1f(x); @@ -1013,40 +1016,40 @@ double yn(int n, double x) // BEGIN INTRINSICS __device__ inline -double __dadd_rd(double x, double y) { return __llvm_add_rtp_f64(x, y); } +double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); } __device__ inline -double __dadd_rn(double x, double y) { return __llvm_add_rte_f64(x, y); } +double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); } __device__ inline -double __dadd_ru(double x, double y) { return __llvm_add_rtn_f64(x, y); } +double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); } __device__ inline -double __dadd_rz(double x, double y) { return __llvm_add_rtz_f64(x, y); } +double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } __device__ inline -double __ddiv_rd(double x, double y) { return __llvm_div_rtp_f64(x, y); } +double __ddiv_rd(double x, double y) { return __ocml_div_rtp_f64(x, y); } __device__ inline -double __ddiv_rn(double x, double y) { return __llvm_div_rte_f64(x, y); } +double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } __device__ inline -double __ddiv_ru(double x, double y) { return __llvm_div_rtn_f64(x, y); } +double __ddiv_ru(double x, double y) { return __ocml_div_rtn_f64(x, y); } __device__ inline -double __ddiv_rz(double x, double y) { return __llvm_div_rtz_f64(x, y); } +double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } __device__ inline -double __dmul_rd(double x, double y) { return __llvm_mul_rtp_f64(x, y); } +double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); } __device__ inline -double __dmul_rn(double x, double y) { return __llvm_mul_rte_f64(x, y); } +double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); } __device__ inline -double __dmul_ru(double x, double y) { return __llvm_mul_rtn_f64(x, y); } +double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); } __device__ inline -double __dmul_rz(double x, double y) { return __llvm_mul_rtz_f64(x, y); } +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); } @@ -1061,51 +1064,51 @@ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } __device__ inline -double __dsqrt_rd(double x) { return __llvm_sqrt_rtp_f64(x); } +double __dsqrt_rd(double x) { return __ocml_sqrt_rtp_f64(x); } __device__ inline -double __dsqrt_rn(double x) { return __llvm_sqrt_rte_f64(x); } +double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } __device__ inline -double __dsqrt_ru(double x) { return __llvm_sqrt_rtn_f64(x); } +double __dsqrt_ru(double x) { return __ocml_sqrt_rtn_f64(x); } __device__ inline -double __dsqrt_rz(double x) { return __llvm_sqrt_rtz_f64(x); } +double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } __device__ inline -double __dsub_rd(double x, double y) { return __llvm_sub_rtp_f64(x, y); } +double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); } __device__ inline -double __dsub_rn(double x, double y) { return __llvm_sub_rte_f64(x, y); } +double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); } __device__ inline -double __dsub_ru(double x, double y) { return __llvm_sub_rtn_f64(x, y); } +double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); } __device__ inline -double __dsub_rz(double x, double y) { return __llvm_sub_rtz_f64(x, y); } +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 __llvm_fma_rtp_f64(x, y, z); + return __ocml_fma_rtp_f64(x, y, z); } __device__ inline double __fma_rn(double x, double y, double z) { - return __llvm_fma_rte_f64(x, y, z); + return __ocml_fma_rte_f64(x, y, z); } __device__ inline double __fma_ru(double x, double y, double z) { - return __llvm_fma_rtn_f64(x, y, z); + return __ocml_fma_rtn_f64(x, y, z); } __device__ inline double __fma_rz(double x, double y, double z) { - return __llvm_fma_rtz_f64(x, y, z); + return __ocml_fma_rtz_f64(x, y, z); } // END INTRINSICS // END DOUBLE diff --git a/hipamd/include/hip/hcc_detail/ocml_math_fwd.h b/hipamd/include/hip/hcc_detail/ocml_math_fwd.h index 773fe0bf9b..ec349010a9 100644 --- a/hipamd/include/hip/hcc_detail/ocml_math_fwd.h +++ b/hipamd/include/hip/hcc_detail/ocml_math_fwd.h @@ -169,53 +169,53 @@ float __ocml_y1_f32(float); // BEGIN INTRINSICS __attribute__((const)) -float __llvm_add_rte_f32(float, float); +float __ocml_add_rte_f32(float, float); __attribute__((const)) -float __llvm_add_rtn_f32(float, float); +float __ocml_add_rtn_f32(float, float); __attribute__((const)) -float __llvm_add_rtp_f32(float, float); +float __ocml_add_rtp_f32(float, float); __attribute__((const)) -float __llvm_add_rtz_f32(float, float); +float __ocml_add_rtz_f32(float, float); __attribute__((const)) -float __llvm_sub_rte_f32(float, float); +float __ocml_sub_rte_f32(float, float); __attribute__((const)) -float __llvm_sub_rtn_f32(float, float); +float __ocml_sub_rtn_f32(float, float); __attribute__((const)) -float __llvm_sub_rtp_f32(float, float); +float __ocml_sub_rtp_f32(float, float); __attribute__((const)) -float __llvm_sub_rtz_f32(float, float); +float __ocml_sub_rtz_f32(float, float); __attribute__((const)) -float __llvm_mul_rte_f32(float, float); +float __ocml_mul_rte_f32(float, float); __attribute__((const)) -float __llvm_mul_rtn_f32(float, float); +float __ocml_mul_rtn_f32(float, float); __attribute__((const)) -float __llvm_mul_rtp_f32(float, float); +float __ocml_mul_rtp_f32(float, float); __attribute__((const)) -float __llvm_mul_rtz_f32(float, float); +float __ocml_mul_rtz_f32(float, float); __attribute__((const)) -float __llvm_div_rte_f32(float, float); +float __ocml_div_rte_f32(float, float); __attribute__((const)) -float __llvm_div_rtn_f32(float, float); +float __ocml_div_rtn_f32(float, float); __attribute__((const)) -float __llvm_div_rtp_f32(float, float); +float __ocml_div_rtp_f32(float, float); __attribute__((const)) -float __llvm_div_rtz_f32(float, float); +float __ocml_div_rtz_f32(float, float); __attribute__((const)) -float __llvm_sqrt_rte_f32(float); +float __ocml_sqrt_rte_f32(float); __attribute__((const)) -float __llvm_sqrt_rtn_f32(float); +float __ocml_sqrt_rtn_f32(float); __attribute__((const)) -float __llvm_sqrt_rtp_f32(float); +float __ocml_sqrt_rtp_f32(float); __attribute__((const)) -float __llvm_sqrt_rtz_f32(float); +float __ocml_sqrt_rtz_f32(float); __attribute__((const)) -float __llvm_fma_rte_f32(float, float, float); +float __ocml_fma_rte_f32(float, float, float); __attribute__((const)) -float __llvm_fma_rtn_f32(float, float, float); +float __ocml_fma_rtn_f32(float, float, float); __attribute__((const)) -float __llvm_fma_rtp_f32(float, float, float); +float __ocml_fma_rtp_f32(float, float, float); __attribute__((const)) -float __llvm_fma_rtz_f32(float, float, float); +float __ocml_fma_rtz_f32(float, float, float); __attribute__((const)) float __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); @@ -371,58 +371,58 @@ double __ocml_y1_f64(double); // BEGIN INTRINSICS __attribute__((const)) -double __llvm_add_rte_f64(double, double); +double __ocml_add_rte_f64(double, double); __attribute__((const)) -double __llvm_add_rtn_f64(double, double); +double __ocml_add_rtn_f64(double, double); __attribute__((const)) -double __llvm_add_rtp_f64(double, double); +double __ocml_add_rtp_f64(double, double); __attribute__((const)) -double __llvm_add_rtz_f64(double, double); +double __ocml_add_rtz_f64(double, double); __attribute__((const)) -double __llvm_sub_rte_f64(double, double); +double __ocml_sub_rte_f64(double, double); __attribute__((const)) -double __llvm_sub_rtn_f64(double, double); +double __ocml_sub_rtn_f64(double, double); __attribute__((const)) -double __llvm_sub_rtp_f64(double, double); +double __ocml_sub_rtp_f64(double, double); __attribute__((const)) -double __llvm_sub_rtz_f64(double, double); +double __ocml_sub_rtz_f64(double, double); __attribute__((const)) -double __llvm_mul_rte_f64(double, double); +double __ocml_mul_rte_f64(double, double); __attribute__((const)) -double __llvm_mul_rtn_f64(double, double); +double __ocml_mul_rtn_f64(double, double); __attribute__((const)) -double __llvm_mul_rtp_f64(double, double); +double __ocml_mul_rtp_f64(double, double); __attribute__((const)) -double __llvm_mul_rtz_f64(double, double); +double __ocml_mul_rtz_f64(double, double); __attribute__((const)) -double __llvm_div_rte_f64(double, double); +double __ocml_div_rte_f64(double, double); __attribute__((const)) -double __llvm_div_rtn_f64(double, double); +double __ocml_div_rtn_f64(double, double); __attribute__((const)) -double __llvm_div_rtp_f64(double, double); +double __ocml_div_rtp_f64(double, double); __attribute__((const)) -double __llvm_div_rtz_f64(double, double); +double __ocml_div_rtz_f64(double, double); __attribute__((const)) -double __llvm_sqrt_rte_f64(double); +double __ocml_sqrt_rte_f64(double); __attribute__((const)) -double __llvm_sqrt_rtn_f64(double); +double __ocml_sqrt_rtn_f64(double); __attribute__((const)) -double __llvm_sqrt_rtp_f64(double); +double __ocml_sqrt_rtp_f64(double); __attribute__((const)) -double __llvm_sqrt_rtz_f64(double); +double __ocml_sqrt_rtz_f64(double); __attribute__((const)) -double __llvm_fma_rte_f64(float, float, float); +double __ocml_fma_rte_f64(double, double, double); __attribute__((const)) -double __llvm_fma_rtn_f64(float, float, float); +double __ocml_fma_rtn_f64(double, double, double); __attribute__((const)) -double __llvm_fma_rtp_f64(float, float, float); +double __ocml_fma_rtp_f64(double, double, double); __attribute__((const)) -double __llvm_fma_rtz_f64(float, float, float); +double __ocml_fma_rtz_f64(double, double, double); __attribute__((const)) -double __llvm_amdgcn_rcp_f64(double); +double __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); __attribute__((const)) -double __llvm_amdgcn_rsq_f64(double); +double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); // END INTRINSICS // END DOUBLE