[clang] ce04d4e - Fix pow and ldexp in HIP header

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2020-07-21T17:39:46-04:00
New Revision: ce04d4e39c936528fa0e0dc49fdf376eaee21f14

URL: https://github.com/llvm/llvm-project/commit/ce04d4e39c936528fa0e0dc49fdf376eaee21f14
DIFF: https://github.com/llvm/llvm-project/commit/ce04d4e39c936528fa0e0dc49fdf376eaee21f14.diff

LOG: Fix pow and ldexp in HIP header

Added: 
    

Modified: 
    clang/lib/Headers/__clang_hip_libdevice_declares.h
    clang/lib/Headers/__clang_hip_math.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index e1cd49a39c65..711040443440 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -78,6 +78,7 @@ __device__ __attribute__((const)) float __ocml_len4_f32(float, float, float,
 __device__ __attribute__((pure)) float __ocml_ncdf_f32(float);
 __device__ __attribute__((pure)) float __ocml_ncdfinv_f32(float);
 __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
+__device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
 __device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
 __device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
 __device__ float __ocml_remquo_f32(float, float,
@@ -205,6 +206,7 @@ __device__ __attribute__((const)) double __ocml_len4_f64(double, double, double,
 __device__ __attribute__((pure)) double __ocml_ncdf_f64(double);
 __device__ __attribute__((pure)) double __ocml_ncdfinv_f64(double);
 __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
+__device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
 __device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
 __device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
 __device__ double __ocml_remquo_f64(double, double,
@@ -290,6 +292,7 @@ __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
 __device__ _Float16 __ocml_sin_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
+__device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
 
 typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
 typedef short __2i16 __attribute__((ext_vector_type(2)));
@@ -320,6 +323,7 @@ __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
 __device__ __2f16 __ocml_sin_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
+__device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
 
 } // extern "C"
 

diff  --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index cf7014b9aefe..47d3c1717559 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -294,6 +294,8 @@ normf(int __dim,
 __DEVICE__
 inline float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
 __DEVICE__
+inline float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }
+__DEVICE__
 inline float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }
 __DEVICE__
 inline float remainderf(float __x, float __y) {
@@ -759,6 +761,8 @@ inline double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }
 __DEVICE__
 inline double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }
 __DEVICE__
+inline double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); }
+__DEVICE__
 inline double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }
 __DEVICE__
 inline double remainder(double __x, double __y) {
@@ -1134,6 +1138,7 @@ __DEF_FUN1(double, trunc);
   __DEVICE__                                                                   \
   inline float __func(float __x, int __y) { return __func##f(__x, __y); }
 __DEF_FLOAT_FUN2I(scalbn)
+__DEF_FLOAT_FUN2I(ldexp)
 
 template <class T> __DEVICE__ inline T min(T __arg1, T __arg2) {
   return (__arg1 < __arg2) ? __arg1 : __arg2;
@@ -1173,6 +1178,17 @@ __host__ inline static int max(int __arg1, int __arg2) {
   return std::max(__arg1, __arg2);
 }
 
+__DEVICE__
+inline float pow(float __base, int __iexp) { return powif(__base, __iexp); }
+
+__DEVICE__
+inline double pow(double __base, int __iexp) { return powi(__base, __iexp); }
+
+__DEVICE__
+inline _Float16 pow(_Float16 __base, int __iexp) {
+  return __ocml_pown_f16(__base, __iexp);
+}
+
 #pragma pop_macro("__DEF_FUN1")
 #pragma pop_macro("__DEF_FUN2")
 #pragma pop_macro("__DEF_FUNI")