diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 63e48fab29..3e9db0d1f6 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -507,13 +507,13 @@ float ynf(int n, float x) // BEGIN INTRINSICS __DEVICE__ inline -float __cosf(float x) { return __ocml_cos_f32(x); } +float __cosf(float x) { return __ocml_native_cos_f32(x); } __DEVICE__ inline -float __exp10f(float x) { return __ocml_exp10_f32(x); } +float __exp10f(float x) { return __ocml_native_exp10_f32(x); } __DEVICE__ inline -float __expf(float x) { return __ocml_exp_f32(x); } +float __expf(float x) { return __ocml_native_exp_f32(x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline @@ -613,7 +613,7 @@ float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); } #endif __DEVICE__ inline -float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); } +float __fsqrt_rn(float x) { return __ocml_native_sqrt_f32(x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ inline @@ -638,13 +638,13 @@ float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); } #endif __DEVICE__ inline -float __log10f(float x) { return __ocml_log10_f32(x); } +float __log10f(float x) { return __ocml_native_log10_f32(x); } __DEVICE__ inline -float __log2f(float x) { return __ocml_log2_f32(x); } +float __log2f(float x) { return __ocml_native_log2_f32(x); } __DEVICE__ inline -float __logf(float x) { return __ocml_log_f32(x); } +float __logf(float x) { return __ocml_native_log_f32(x); } __DEVICE__ inline float __powf(float x, float y) { return __ocml_pow_f32(x, y); } @@ -655,15 +655,12 @@ __DEVICE__ inline void __sincosf(float x, float* sptr, float* cptr) { - float tmp; - - *sptr = - __ocml_sincos_f32(x, (__attribute__((address_space(5))) float*) &tmp); - *cptr = tmp; + *sptr = __ocml_native_sin_f32(x); + *cptr = __ocml_native_cos_f32(x); } __DEVICE__ inline -float __sinf(float x) { return __ocml_sin_f32(x); } +float __sinf(float x) { return __ocml_native_sin_f32(x); } __DEVICE__ inline float __tanf(float x) { return __ocml_tan_f32(x); } diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index e5594924ba..cf8aeb9c6c 100644 --- a/projects/hip/include/hip/hcc_detail/math_fwd.h +++ b/projects/hip/include/hip/hcc_detail/math_fwd.h @@ -63,6 +63,8 @@ float __ocml_copysign_f32(float, float); __device__ float __ocml_cos_f32(float); __device__ +float __ocml_native_cos_f32(float); +__device__ __attribute__((pure)) __device__ float __ocml_cosh_f32(float); @@ -92,12 +94,18 @@ __attribute__((pure)) float __ocml_exp10_f32(float); __device__ __attribute__((pure)) +float __ocml_native_exp10_f32(float); +__device__ +__attribute__((pure)) float __ocml_exp2_f32(float); __device__ __attribute__((pure)) float __ocml_exp_f32(float); __device__ __attribute__((pure)) +float __ocml_native_exp_f32(float); +__device__ +__attribute__((pure)) float __ocml_expm1_f32(float); __device__ __attribute__((const)) @@ -152,17 +160,26 @@ __attribute__((pure)) float __ocml_log10_f32(float); __device__ __attribute__((pure)) +float __ocml_native_log10_f32(float); +__device__ +__attribute__((pure)) float __ocml_log1p_f32(float); __device__ __attribute__((pure)) float __ocml_log2_f32(float); __device__ +__attribute__((pure)) +float __ocml_native_log2_f32(float); +__device__ __attribute__((const)) float __ocml_logb_f32(float); __device__ __attribute__((pure)) float __ocml_log_f32(float); __device__ +__attribute__((pure)) +float __ocml_native_log_f32(float); +__device__ float __ocml_modf_f32(float, __attribute__((address_space(5))) float*); __device__ __attribute__((const)) @@ -227,6 +244,8 @@ float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float*); __device__ float __ocml_sin_f32(float); __device__ +float __ocml_native_sin_f32(float); +__device__ __attribute__((pure)) float __ocml_sinh_f32(float); __device__ @@ -235,6 +254,9 @@ __device__ __attribute__((const)) float __ocml_sqrt_f32(float); __device__ +__attribute__((const)) +float __ocml_native_sqrt_f32(float); +__device__ float __ocml_tan_f32(float); __device__ __attribute__((pure))