Merge pull request #864 from mangupta/swdev-157382

Switch to faster ocml variants for trig instrinsics

[ROCm/hip commit: c85c8959f6]

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); }

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))