diff --git a/include/hip/hcc_detail/hip_fp16_math_fwd.h b/include/hip/hcc_detail/hip_fp16_math_fwd.h index 16d834cca8..53a2c66f9c 100644 --- a/include/hip/hcc_detail/hip_fp16_math_fwd.h +++ b/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); diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 7b9edfffd1..60dc644ecd 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/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") diff --git a/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h index 4c0fde591c..34b38aa723 100644 --- a/include/hip/hcc_detail/math_fwd.h +++ b/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)) diff --git a/tests/src/deviceLib/hipMathFunctions.cpp b/tests/src/deviceLib/hipMathFunctions.cpp index c99d8a91f2..cbc608b72b 100644 --- a/tests/src/deviceLib/hipMathFunctions.cpp +++ b/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(); }