From 5db57f36d7743b0e1b7368e4b3db3fa0b5ffa972 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 14 Jan 2019 13:01:57 +0530 Subject: [PATCH 1/2] Switch to faster ocml variants for trig instrinsics Change-Id: If62821e2fe1b0da91ad1b8c5580ebf1a009405e9 [ROCm/hip commit: 2b9d52b21025954f97a349fec474bf234fa41ffe] --- projects/hip/include/hip/hcc_detail/math_functions.h | 11 ++++------- projects/hip/include/hip/hcc_detail/math_fwd.h | 4 ++++ 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 63e48fab29..8ac9ca068c 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index e5594924ba..ab57ebb909 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); @@ -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__ From 09fe2d6bc5e2ed3b4d4e99c3a197685865c5dd49 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 15 Jan 2019 19:40:37 +0530 Subject: [PATCH 2/2] 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))