From 09fe2d6bc5e2ed3b4d4e99c3a197685865c5dd49 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Tue, 15 Jan 2019 19:40:37 +0530
Subject: [PATCH] Map more instrincis to ocml native

Change-Id: I54e55d5902931bafdabd0956e4e8c1d7b39f7173


[ROCm/hip commit: dd9bea7d25f2417b6f55159f592488614ffe7970]
---
 .../include/hip/hcc_detail/math_functions.h    | 12 ++++++------
 projects/hip/include/hip/hcc_detail/math_fwd.h | 18 ++++++++++++++++++
 2 files changed, 24 insertions(+), 6 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h
index 8ac9ca068c..3e9db0d1f6 100644
--- a/projects/hip/include/hip/hcc_detail/math_functions.h
+++ b/projects/hip/include/hip/hcc_detail/math_functions.h
@@ -510,10 +510,10 @@ inline
 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); }
diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h
index ab57ebb909..cf8aeb9c6c 100644
--- a/projects/hip/include/hip/hcc_detail/math_fwd.h
+++ b/projects/hip/include/hip/hcc_detail/math_fwd.h
@@ -94,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))
@@ -154,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))
@@ -239,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))