From 2b9d52b21025954f97a349fec474bf234fa41ffe Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 14 Jan 2019 13:01:57 +0530
Subject: [PATCH] Switch to faster ocml variants for trig instrinsics

Change-Id: If62821e2fe1b0da91ad1b8c5580ebf1a009405e9
---
 include/hip/hcc_detail/math_functions.h | 11 ++++-------
 include/hip/hcc_detail/math_fwd.h       |  4 ++++
 2 files changed, 8 insertions(+), 7 deletions(-)

diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h
index 63e48fab29..8ac9ca068c 100644
--- a/include/hip/hcc_detail/math_functions.h
+++ b/include/hip/hcc_detail/math_functions.h
@@ -507,7 +507,7 @@ 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); }
@@ -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/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h
index e5594924ba..ab57ebb909 100644
--- a/include/hip/hcc_detail/math_fwd.h
+++ b/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);
@@ -227,6 +229,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__