Add pow(float/double/_Float16,int)

Change-Id: Ie65d15cd3df9853a3bbd613d8c7188ae39c327c7


[ROCm/clr commit: b72e8b54e0]

projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h

@@ -44,6 +44,7 @@ extern "C"
     __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
     __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
     __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
     __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
     __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
     __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);

projects/clr/hipamd/include/hip/hcc_detail/math_functions.h

@@ -430,6 +430,9 @@ inline
 float powf(float x, float y) { return __ocml_pow_f32(x, y); }
 __DEVICE__
 inline
+float powif(float base, int iexp) { return __ocml_pown_f32(base, iexp); }
+__DEVICE__
+inline
 float rcbrtf(float x) { return __ocml_rcbrt_f32(x); }
 __DEVICE__
 inline
@@ -985,6 +988,9 @@ inline
 double pow(double x, double y) { return __ocml_pow_f64(x, y); }
 __DEVICE__
 inline
+double powi(double base, int iexp) { return __ocml_pown_f64(base, iexp); }
+__DEVICE__
+inline
 double rcbrt(double x) { return __ocml_rcbrt_f64(x); }
 __DEVICE__
 inline
@@ -1511,6 +1517,22 @@ __host__ inline static int min(int arg1, int arg2) {
 __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);
+}
+
 #endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 
 #pragma pop_macro("__DEF_FLOAT_FUN")

projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h

@@ -243,6 +243,9 @@ __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))
@@ -555,6 +558,9 @@ __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))

projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp

@@ -186,6 +186,14 @@ void test_fp16() {
   CHECK_SIMPLE([]__device__(){ return min<__fp16>(1.0f, 2.0f); }, 1.0f);
 }
 
+void test_pown() {
+  CHECK_SIMPLE([]__device__(){ return powif(2.0f, 2); }, 4.0f);
+  CHECK_SIMPLE([]__device__(){ return powi(2.0, 2); }, 4.0);
+  CHECK_SIMPLE([]__device__(){ return pow(2.0f, 2); }, 4.0f);
+  CHECK_SIMPLE([]__device__(){ return pow(2.0, 2); }, 4.0);
+  CHECK_SIMPLE([]__device__(){ return pow(2.0f16, 2); }, 4.0f16);
+}
+
 int main(int argc, char* argv[]) {
     HipTest::parseStandardArguments(argc, argv, true);
 
@@ -195,5 +203,7 @@ int main(int argc, char* argv[]) {
     
     test_fp16();
 
+    test_pown();
+
     passed();
 }