Author: Aaron En Ye Shi Date: 2020-10-06T14:48:53Z New Revision: aa2b593f1495a972a4a592952760ec9d5f7c01f1
URL: https://github.com/llvm/llvm-project/commit/aa2b593f1495a972a4a592952760ec9d5f7c01f1 DIFF: https://github.com/llvm/llvm-project/commit/aa2b593f1495a972a4a592952760ec9d5f7c01f1.diff LOG: [HIP] Restructure hip headers to add cmath Separate __clang_hip_math.h header into __clang_hip_cmath.h and __clang_hip_math.h. Improve the math function definition, and add missing definitions or declarations. Add missing overloads. Reviewed By: tra, JonChesterfield Differential Review: https://reviews.llvm.org/D88837 Added: clang/lib/Headers/__clang_hip_cmath.h Modified: clang/lib/Headers/CMakeLists.txt clang/lib/Headers/__clang_hip_libdevice_declares.h clang/lib/Headers/__clang_hip_math.h clang/lib/Headers/__clang_hip_runtime_wrapper.h Removed: ################################################################################ diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index 95047e7069e7..533ff4506ffe 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -47,6 +47,7 @@ set(files __clang_cuda_math_forward_declares.h __clang_cuda_runtime_wrapper.h __clang_hip_libdevice_declares.h + __clang_hip_cmath.h __clang_hip_math.h __clang_hip_runtime_wrapper.h cetintrin.h diff --git a/clang/lib/Headers/__clang_hip_cmath.h b/clang/lib/Headers/__clang_hip_cmath.h new file mode 100644 index 000000000000..067c7e6c9d1b --- /dev/null +++ b/clang/lib/Headers/__clang_hip_cmath.h @@ -0,0 +1,521 @@ +/*===---- __clang_hip_cmath.h - HIP cmath decls -----------------------------=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#ifndef __CLANG_HIP_CMATH_H__ +#define __CLANG_HIP_CMATH_H__ + +#if !defined(__HIP__) +#error "This file is for HIP and OpenMP AMDGCN device compilation only." +#endif + +#if defined(__cplusplus) +#include <limits> +#endif +#include <limits.h> +#include <stdint.h> + +#pragma push_macro("__DEVICE__") +#define __DEVICE__ static __device__ inline __attribute__((always_inline)) + +// Start with functions that cannot be defined by DEF macros below. +#if defined(__cplusplus) +__DEVICE__ double abs(double __x) { return ::fabs(__x); } +__DEVICE__ float abs(float __x) { return ::fabsf(__x); } +__DEVICE__ long long abs(long long __n) { return ::llabs(__n); } +__DEVICE__ long abs(long __n) { return ::labs(__n); } +__DEVICE__ float fma(float __x, float __y, float __z) { + return ::fmaf(__x, __y, __z); +} +__DEVICE__ int fpclassify(float __x) { + return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, + FP_ZERO, __x); +} +__DEVICE__ int fpclassify(double __x) { + return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, + FP_ZERO, __x); +} +__DEVICE__ float frexp(float __arg, int *__exp) { + return ::frexpf(__arg, __exp); +} +__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } +__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); } +__DEVICE__ bool isgreater(float __x, float __y) { + return __builtin_isgreater(__x, __y); +} +__DEVICE__ bool isgreater(double __x, double __y) { + return __builtin_isgreater(__x, __y); +} +__DEVICE__ bool isgreaterequal(float __x, float __y) { + return __builtin_isgreaterequal(__x, __y); +} +__DEVICE__ bool isgreaterequal(double __x, double __y) { + return __builtin_isgreaterequal(__x, __y); +} +__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } +__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } +__DEVICE__ bool isless(float __x, float __y) { + return __builtin_isless(__x, __y); +} +__DEVICE__ bool isless(double __x, double __y) { + return __builtin_isless(__x, __y); +} +__DEVICE__ bool islessequal(float __x, float __y) { + return __builtin_islessequal(__x, __y); +} +__DEVICE__ bool islessequal(double __x, double __y) { + return __builtin_islessequal(__x, __y); +} +__DEVICE__ bool islessgreater(float __x, float __y) { + return __builtin_islessgreater(__x, __y); +} +__DEVICE__ bool islessgreater(double __x, double __y) { + return __builtin_islessgreater(__x, __y); +} +__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); } +__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); } +__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); } +__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); } +__DEVICE__ bool isunordered(float __x, float __y) { + return __builtin_isunordered(__x, __y); +} +__DEVICE__ bool isunordered(double __x, double __y) { + return __builtin_isunordered(__x, __y); +} +__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } +__DEVICE__ float pow(float __base, int __iexp) { + return ::powif(__base, __iexp); +} +__DEVICE__ double pow(double __base, int __iexp) { + return ::powi(__base, __iexp); +} +__DEVICE__ float remquo(float __x, float __y, int *__quo) { + return ::remquof(__x, __y, __quo); +} +__DEVICE__ float scalbln(float __x, long int __n) { + return ::scalblnf(__x, __n); +} +__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); } +__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); } + +// Notably missing above is nexttoward. We omit it because +// ocml doesn't provide an implementation, and we don't want to be in the +// business of implementing tricky libm functions in this header. + +// Other functions. +__DEVICE__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) { + return __ocml_fma_f16(__x, __y, __z); +} +__DEVICE__ _Float16 pow(_Float16 __base, int __iexp) { + return __ocml_pown_f16(__base, __iexp); +} + +// BEGIN DEF_FUN and HIP_OVERLOAD + +// BEGIN DEF_FUN + +#pragma push_macro("__DEF_FUN1") +#pragma push_macro("__DEF_FUN2") +#pragma push_macro("__DEF_FUN2_FI") + +// Define cmath functions with float argument and returns __retty. +#define __DEF_FUN1(__retty, __func) \ + __DEVICE__ \ + __retty __func(float __x) { return __func##f(__x); } + +// Define cmath functions with two float arguments and returns __retty. +#define __DEF_FUN2(__retty, __func) \ + __DEVICE__ \ + __retty __func(float __x, float __y) { return __func##f(__x, __y); } + +// Define cmath functions with a float and an int argument and returns __retty. +#define __DEF_FUN2_FI(__retty, __func) \ + __DEVICE__ \ + __retty __func(float __x, int __y) { return __func##f(__x, __y); } + +__DEF_FUN1(float, acos) +__DEF_FUN1(float, acosh) +__DEF_FUN1(float, asin) +__DEF_FUN1(float, asinh) +__DEF_FUN1(float, atan) +__DEF_FUN2(float, atan2) +__DEF_FUN1(float, atanh) +__DEF_FUN1(float, cbrt) +__DEF_FUN1(float, ceil) +__DEF_FUN2(float, copysign) +__DEF_FUN1(float, cos) +__DEF_FUN1(float, cosh) +__DEF_FUN1(float, erf) +__DEF_FUN1(float, erfc) +__DEF_FUN1(float, exp) +__DEF_FUN1(float, exp2) +__DEF_FUN1(float, expm1) +__DEF_FUN1(float, fabs) +__DEF_FUN2(float, fdim) +__DEF_FUN1(float, floor) +__DEF_FUN2(float, fmax) +__DEF_FUN2(float, fmin) +__DEF_FUN2(float, fmod) +__DEF_FUN2(float, hypot) +__DEF_FUN1(int, ilogb) +__DEF_FUN2_FI(float, ldexp) +__DEF_FUN1(float, lgamma) +__DEF_FUN1(float, log) +__DEF_FUN1(float, log10) +__DEF_FUN1(float, log1p) +__DEF_FUN1(float, log2) +__DEF_FUN1(float, logb) +__DEF_FUN1(long long, llrint) +__DEF_FUN1(long long, llround) +__DEF_FUN1(long, lrint) +__DEF_FUN1(long, lround) +__DEF_FUN1(float, nearbyint) +__DEF_FUN2(float, nextafter) +__DEF_FUN2(float, pow) +__DEF_FUN2(float, remainder) +__DEF_FUN1(float, rint) +__DEF_FUN1(float, round) +__DEF_FUN2_FI(float, scalbn) +__DEF_FUN1(float, sin) +__DEF_FUN1(float, sinh) +__DEF_FUN1(float, sqrt) +__DEF_FUN1(float, tan) +__DEF_FUN1(float, tanh) +__DEF_FUN1(float, tgamma) +__DEF_FUN1(float, trunc) + +#pragma pop_macro("__DEF_FUN1") +#pragma pop_macro("__DEF_FUN2") +#pragma pop_macro("__DEF_FUN2_FI") + +// END DEF_FUN + +// BEGIN HIP_OVERLOAD + +#pragma push_macro("__HIP_OVERLOAD1") +#pragma push_macro("__HIP_OVERLOAD2") + +// __hip_enable_if::type is a type function which returns __T if __B is true. +template <bool __B, class __T = void> struct __hip_enable_if {}; + +template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; }; + +// __HIP_OVERLOAD1 is used to resolve function calls with integer argument to +// avoid compilation error due to ambibuity. e.g. floor(5) is resolved with +// floor(double). +#define __HIP_OVERLOAD1(__retty, __fn) \ + template <typename __T> \ + __DEVICE__ typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \ + __retty>::type \ + __fn(__T __x) { \ + return ::__fn((double)__x); \ + } + +// __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double +// or integer argument to avoid compilation error due to ambibuity. e.g. +// max(5.0f, 6.0) is resolved with max(double, double). +#define __HIP_OVERLOAD2(__retty, __fn) \ + template <typename __T1, typename __T2> \ + __DEVICE__ \ + typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && \ + std::numeric_limits<__T2>::is_specialized, \ + __retty>::type \ + __fn(__T1 __x, __T2 __y) { \ + return __fn((double)__x, (double)__y); \ + } + +__HIP_OVERLOAD1(double, abs) +__HIP_OVERLOAD1(double, acos) +__HIP_OVERLOAD1(double, acosh) +__HIP_OVERLOAD1(double, asin) +__HIP_OVERLOAD1(double, asinh) +__HIP_OVERLOAD1(double, atan) +__HIP_OVERLOAD2(double, atan2) +__HIP_OVERLOAD1(double, atanh) +__HIP_OVERLOAD1(double, cbrt) +__HIP_OVERLOAD1(double, ceil) +__HIP_OVERLOAD2(double, copysign) +__HIP_OVERLOAD1(double, cos) +__HIP_OVERLOAD1(double, cosh) +__HIP_OVERLOAD1(double, erf) +__HIP_OVERLOAD1(double, erfc) +__HIP_OVERLOAD1(double, exp) +__HIP_OVERLOAD1(double, exp2) +__HIP_OVERLOAD1(double, expm1) +__HIP_OVERLOAD1(double, fabs) +__HIP_OVERLOAD2(double, fdim) +__HIP_OVERLOAD1(double, floor) +__HIP_OVERLOAD2(double, fmax) +__HIP_OVERLOAD2(double, fmin) +__HIP_OVERLOAD2(double, fmod) +__HIP_OVERLOAD1(int, fpclassify) +__HIP_OVERLOAD2(double, hypot) +__HIP_OVERLOAD1(int, ilogb) +__HIP_OVERLOAD1(bool, isfinite) +__HIP_OVERLOAD2(bool, isgreater) +__HIP_OVERLOAD2(bool, isgreaterequal) +__HIP_OVERLOAD1(bool, isinf) +__HIP_OVERLOAD2(bool, isless) +__HIP_OVERLOAD2(bool, islessequal) +__HIP_OVERLOAD2(bool, islessgreater) +__HIP_OVERLOAD1(bool, isnan) +__HIP_OVERLOAD1(bool, isnormal) +__HIP_OVERLOAD2(bool, isunordered) +__HIP_OVERLOAD1(double, lgamma) +__HIP_OVERLOAD1(double, log) +__HIP_OVERLOAD1(double, log10) +__HIP_OVERLOAD1(double, log1p) +__HIP_OVERLOAD1(double, log2) +__HIP_OVERLOAD1(double, logb) +__HIP_OVERLOAD1(long long, llrint) +__HIP_OVERLOAD1(long long, llround) +__HIP_OVERLOAD1(long, lrint) +__HIP_OVERLOAD1(long, lround) +__HIP_OVERLOAD1(double, nearbyint) +__HIP_OVERLOAD2(double, nextafter) +__HIP_OVERLOAD2(double, pow) +__HIP_OVERLOAD2(double, remainder) +__HIP_OVERLOAD1(double, rint) +__HIP_OVERLOAD1(double, round) +__HIP_OVERLOAD1(bool, signbit) +__HIP_OVERLOAD1(double, sin) +__HIP_OVERLOAD1(double, sinh) +__HIP_OVERLOAD1(double, sqrt) +__HIP_OVERLOAD1(double, tan) +__HIP_OVERLOAD1(double, tanh) +__HIP_OVERLOAD1(double, tgamma) +__HIP_OVERLOAD1(double, trunc) + +// Overload these but don't add them to std, they are not part of cmath. +__HIP_OVERLOAD2(double, max) +__HIP_OVERLOAD2(double, min) + +// Additional Overloads that don't quite match HIP_OVERLOAD. +template <typename __T1, typename __T2, typename __T3> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && + std::numeric_limits<__T2>::is_specialized && + std::numeric_limits<__T3>::is_specialized, + double>::type + fma(__T1 __x, __T2 __y, __T3 __z) { + return ::fma((double)__x, (double)__y, (double)__z); +} + +template <typename __T> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type + frexp(__T __x, int *__exp) { + return ::frexp((double)__x, __exp); +} + +template <typename __T> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type + ldexp(__T __x, int __exp) { + return ::ldexp((double)__x, __exp); +} + +template <typename __T> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type + modf(__T __x, double *__exp) { + return ::modf((double)__x, __exp); +} + +template <typename __T1, typename __T2> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && + std::numeric_limits<__T2>::is_specialized, + double>::type + remquo(__T1 __x, __T2 __y, int *__quo) { + return ::remquo((double)__x, (double)__y, __quo); +} + +template <typename __T> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type + scalbln(__T __x, long int __exp) { + return ::scalbln((double)__x, __exp); +} + +template <typename __T> +__DEVICE__ + typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type + scalbn(__T __x, int __exp) { + return ::scalbn((double)__x, __exp); +} + +#pragma pop_macro("__HIP_OVERLOAD1") +#pragma pop_macro("__HIP_OVERLOAD2") + +// END HIP_OVERLOAD + +// END DEF_FUN and HIP_OVERLOAD + +#endif // defined(__cplusplus) + +// Define these overloads inside the namespace our standard library uses. +#ifdef _LIBCPP_BEGIN_NAMESPACE_STD +_LIBCPP_BEGIN_NAMESPACE_STD +#else +namespace std { +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_BEGIN_NAMESPACE_VERSION +#endif +#endif + +// Pull the new overloads we defined above into namespace std. +// using ::abs; +using ::acos; +using ::acosh; +using ::asin; +using ::asinh; +using ::atan; +using ::atan2; +using ::atanh; +using ::cbrt; +using ::ceil; +using ::copysign; +using ::cos; +using ::cosh; +using ::erf; +using ::erfc; +using ::exp; +using ::exp2; +using ::expm1; +using ::fabs; +using ::fdim; +using ::floor; +using ::fma; +using ::fmax; +using ::fmin; +using ::fmod; +using ::fpclassify; +using ::frexp; +using ::hypot; +using ::ilogb; +using ::isfinite; +using ::isgreater; +using ::isgreaterequal; +using ::isless; +using ::islessequal; +using ::islessgreater; +using ::isnormal; +using ::isunordered; +using ::ldexp; +using ::lgamma; +using ::llrint; +using ::llround; +using ::log; +using ::log10; +using ::log1p; +using ::log2; +using ::logb; +using ::lrint; +using ::lround; +using ::modf; +// using ::nan; +// using ::nanf; +// using ::nanl; - This is not yet defined. +using ::nearbyint; +using ::nextafter; +// using ::nexttoward; - Omit this since we do not have a definition. +using ::pow; +using ::remainder; +using ::remquo; +using ::rint; +using ::round; +using ::scalbln; +using ::scalbn; +using ::signbit; +using ::sin; +using ::sinh; +using ::sqrt; +using ::tan; +using ::tanh; +using ::tgamma; +using ::trunc; + +// Well this is fun: We need to pull these symbols in for libc++, but we can't +// pull them in with libstdc++, because its ::isinf and ::isnan are diff erent +// than its std::isinf and std::isnan. +#ifndef __GLIBCXX__ +using ::isinf; +using ::isnan; +#endif + +// Finally, pull the "foobarf" functions that HIP defines into std. +using ::acosf; +using ::acoshf; +using ::asinf; +using ::asinhf; +using ::atan2f; +using ::atanf; +using ::atanhf; +using ::cbrtf; +using ::ceilf; +using ::copysignf; +using ::cosf; +using ::coshf; +using ::erfcf; +using ::erff; +using ::exp2f; +using ::expf; +using ::expm1f; +using ::fabsf; +using ::fdimf; +using ::floorf; +using ::fmaf; +using ::fmaxf; +using ::fminf; +using ::fmodf; +using ::frexpf; +using ::hypotf; +using ::ilogbf; +using ::ldexpf; +using ::lgammaf; +using ::llrintf; +using ::llroundf; +using ::log10f; +using ::log1pf; +using ::log2f; +using ::logbf; +using ::logf; +using ::lrintf; +using ::lroundf; +using ::modff; +using ::nearbyintf; +using ::nextafterf; +// using ::nexttowardf; - Omit this since we do not have a definition. +using ::powf; +using ::remainderf; +using ::remquof; +using ::rintf; +using ::roundf; +using ::scalblnf; +using ::scalbnf; +using ::sinf; +using ::sinhf; +using ::sqrtf; +using ::tanf; +using ::tanhf; +using ::tgammaf; +using ::truncf; + +#ifdef _LIBCPP_END_NAMESPACE_STD +_LIBCPP_END_NAMESPACE_STD +#else +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_END_NAMESPACE_VERSION +#endif +} // namespace std +#endif + +#pragma pop_macro("__DEVICE__") + +#endif // __CLANG_HIP_CMATH_H__ diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h index de2f82cd8eca..ac98907ad5de 100644 --- a/clang/lib/Headers/__clang_hip_libdevice_declares.h +++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -129,10 +129,10 @@ __device__ __attribute__((const)) float __ocml_div_rte_f32(float, float); __device__ __attribute__((const)) float __ocml_div_rtn_f32(float, float); __device__ __attribute__((const)) float __ocml_div_rtp_f32(float, float); __device__ __attribute__((const)) float __ocml_div_rtz_f32(float, float); -__device__ __attribute__((const)) float __ocml_sqrt_rte_f32(float, float); -__device__ __attribute__((const)) float __ocml_sqrt_rtn_f32(float, float); -__device__ __attribute__((const)) float __ocml_sqrt_rtp_f32(float, float); -__device__ __attribute__((const)) float __ocml_sqrt_rtz_f32(float, float); +__device__ __attribute__((const)) float __ocml_sqrt_rte_f32(float); +__device__ __attribute__((const)) float __ocml_sqrt_rtn_f32(float); +__device__ __attribute__((const)) float __ocml_sqrt_rtp_f32(float); +__device__ __attribute__((const)) float __ocml_sqrt_rtz_f32(float); __device__ __attribute__((const)) float __ocml_fma_rte_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtn_f32(float, float, float); __device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float); @@ -256,10 +256,10 @@ __device__ __attribute__((const)) double __ocml_div_rte_f64(double, double); __device__ __attribute__((const)) double __ocml_div_rtn_f64(double, double); __device__ __attribute__((const)) double __ocml_div_rtp_f64(double, double); __device__ __attribute__((const)) double __ocml_div_rtz_f64(double, double); -__device__ __attribute__((const)) double __ocml_sqrt_rte_f64(double, double); -__device__ __attribute__((const)) double __ocml_sqrt_rtn_f64(double, double); -__device__ __attribute__((const)) double __ocml_sqrt_rtp_f64(double, double); -__device__ __attribute__((const)) double __ocml_sqrt_rtz_f64(double, double); +__device__ __attribute__((const)) double __ocml_sqrt_rte_f64(double); +__device__ __attribute__((const)) double __ocml_sqrt_rtn_f64(double); +__device__ __attribute__((const)) double __ocml_sqrt_rtp_f64(double); +__device__ __attribute__((const)) double __ocml_sqrt_rtz_f64(double); __device__ __attribute__((const)) double __ocml_fma_rte_f64(double, double, double); __device__ __attribute__((const)) double __ocml_fma_rtn_f64(double, double, diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index b72bb40ccdb6..f2365e8844fe 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -1,4 +1,4 @@ -/*===---- __clang_hip_math.h - HIP math decls -------------------------------=== +/*===---- __clang_hip_math.h - Device-side HIP math support ----------------=== * * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. * See https://llvm.org/LICENSE.txt for license information. @@ -6,25 +6,32 @@ * *===-----------------------------------------------------------------------=== */ - #ifndef __CLANG_HIP_MATH_H__ #define __CLANG_HIP_MATH_H__ +#if !defined(__HIP__) +#error "This file is for HIP and OpenMP AMDGCN device compilation only." +#endif + +#if defined(__cplusplus) #include <algorithm> +#endif #include <limits.h> -#include <limits> #include <stdint.h> -#include <assert.h> #pragma push_macro("__DEVICE__") -#pragma push_macro("__RETURN_TYPE") +#define __DEVICE__ static __device__ inline __attribute__((always_inline)) -// to be consistent with __clang_cuda_math_forward_declares -#define __DEVICE__ static __device__ +// A few functions return bool type starting only in C++11. +#pragma push_macro("__RETURN_TYPE") +#if defined(__cplusplus) #define __RETURN_TYPE bool +#else +#define __RETURN_TYPE int +#endif #if defined (__cplusplus) && __cplusplus < 201103L -//emulate static_assert on type sizes +// emulate static_assert on type sizes template<bool> struct __compare_result{}; template<> @@ -33,26 +40,23 @@ struct __compare_result<true> { }; __DEVICE__ -inline void __suppress_unused_warning(bool b) {}; -template<unsigned int S, unsigned int T> -__DEVICE__ -inline void __static_assert_equal_size() { - __suppress_unused_warning(__compare_result<S==T>::valid); +void __suppress_unused_warning(bool b){}; +template <unsigned int S, unsigned int T> +__DEVICE__ void __static_assert_equal_size() { + __suppress_unused_warning(__compare_result<S == T>::valid); } #define __static_assert_type_size_equal(A, B) \ __static_assert_equal_size<A,B>() #else - #define __static_assert_type_size_equal(A,B) \ static_assert((A) == (B), "") #endif - __DEVICE__ -inline uint64_t __make_mantissa_base8(const char *__tagp) { +uint64_t __make_mantissa_base8(const char *__tagp) { uint64_t __r = 0; while (__tagp) { char __tmp = *__tagp; @@ -69,7 +73,7 @@ inline uint64_t __make_mantissa_base8(const char *__tagp) { } __DEVICE__ -inline uint64_t __make_mantissa_base10(const char *__tagp) { +uint64_t __make_mantissa_base10(const char *__tagp) { uint64_t __r = 0; while (__tagp) { char __tmp = *__tagp; @@ -86,7 +90,7 @@ inline uint64_t __make_mantissa_base10(const char *__tagp) { } __DEVICE__ -inline uint64_t __make_mantissa_base16(const char *__tagp) { +uint64_t __make_mantissa_base16(const char *__tagp) { uint64_t __r = 0; while (__tagp) { char __tmp = *__tagp; @@ -107,7 +111,7 @@ inline uint64_t __make_mantissa_base16(const char *__tagp) { } __DEVICE__ -inline uint64_t __make_mantissa(const char *__tagp) { +uint64_t __make_mantissa(const char *__tagp) { if (!__tagp) return 0u; @@ -124,80 +128,124 @@ inline uint64_t __make_mantissa(const char *__tagp) { } // BEGIN FLOAT -#ifdef __cplusplus +#if defined(__cplusplus) +__DEVICE__ +int abs(int __x) { + int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1); + return (__x ^ __sgn) - __sgn; +} +__DEVICE__ +long labs(long __x) { + long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1); + return (__x ^ __sgn) - __sgn; +} __DEVICE__ -inline float abs(float __x) { return __ocml_fabs_f32(__x); } +long long llabs(long long __x) { + long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1); + return (__x ^ __sgn) - __sgn; +} #endif + __DEVICE__ -inline float acosf(float __x) { return __ocml_acos_f32(__x); } +float acosf(float __x) { return __ocml_acos_f32(__x); } + __DEVICE__ -inline float acoshf(float __x) { return __ocml_acosh_f32(__x); } +float acoshf(float __x) { return __ocml_acosh_f32(__x); } + __DEVICE__ -inline float asinf(float __x) { return __ocml_asin_f32(__x); } +float asinf(float __x) { return __ocml_asin_f32(__x); } + __DEVICE__ -inline float asinhf(float __x) { return __ocml_asinh_f32(__x); } +float asinhf(float __x) { return __ocml_asinh_f32(__x); } + __DEVICE__ -inline float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); } +float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); } + __DEVICE__ -inline float atanf(float __x) { return __ocml_atan_f32(__x); } +float atanf(float __x) { return __ocml_atan_f32(__x); } + __DEVICE__ -inline float atanhf(float __x) { return __ocml_atanh_f32(__x); } +float atanhf(float __x) { return __ocml_atanh_f32(__x); } + __DEVICE__ -inline float cbrtf(float __x) { return __ocml_cbrt_f32(__x); } +float cbrtf(float __x) { return __ocml_cbrt_f32(__x); } + __DEVICE__ -inline float ceilf(float __x) { return __ocml_ceil_f32(__x); } +float ceilf(float __x) { return __ocml_ceil_f32(__x); } + __DEVICE__ -inline float copysignf(float __x, float __y) { - return __ocml_copysign_f32(__x, __y); -} +float copysignf(float __x, float __y) { return __ocml_copysign_f32(__x, __y); } + __DEVICE__ -inline float cosf(float __x) { return __ocml_cos_f32(__x); } +float cosf(float __x) { return __ocml_cos_f32(__x); } + __DEVICE__ -inline float coshf(float __x) { return __ocml_cosh_f32(__x); } +float coshf(float __x) { return __ocml_cosh_f32(__x); } + __DEVICE__ -inline float cospif(float __x) { return __ocml_cospi_f32(__x); } +float cospif(float __x) { return __ocml_cospi_f32(__x); } + __DEVICE__ -inline float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); } +float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); } + __DEVICE__ -inline float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); } +float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); } + __DEVICE__ -inline float erfcf(float __x) { return __ocml_erfc_f32(__x); } +float erfcf(float __x) { return __ocml_erfc_f32(__x); } + __DEVICE__ -inline float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); } +float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); } + __DEVICE__ -inline float erfcxf(float __x) { return __ocml_erfcx_f32(__x); } +float erfcxf(float __x) { return __ocml_erfcx_f32(__x); } + __DEVICE__ -inline float erff(float __x) { return __ocml_erf_f32(__x); } +float erff(float __x) { return __ocml_erf_f32(__x); } + __DEVICE__ -inline float erfinvf(float __x) { return __ocml_erfinv_f32(__x); } +float erfinvf(float __x) { return __ocml_erfinv_f32(__x); } + __DEVICE__ -inline float exp10f(float __x) { return __ocml_exp10_f32(__x); } +float exp10f(float __x) { return __ocml_exp10_f32(__x); } + __DEVICE__ -inline float exp2f(float __x) { return __ocml_exp2_f32(__x); } +float exp2f(float __x) { return __ocml_exp2_f32(__x); } + __DEVICE__ -inline float expf(float __x) { return __ocml_exp_f32(__x); } +float expf(float __x) { return __ocml_exp_f32(__x); } + __DEVICE__ -inline float expm1f(float __x) { return __ocml_expm1_f32(__x); } +float expm1f(float __x) { return __ocml_expm1_f32(__x); } + __DEVICE__ -inline float fabsf(float __x) { return __ocml_fabs_f32(__x); } +float fabsf(float __x) { return __ocml_fabs_f32(__x); } + __DEVICE__ -inline float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); } +float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); } + __DEVICE__ -inline float fdividef(float __x, float __y) { return __x / __y; } +float fdividef(float __x, float __y) { return __x / __y; } + __DEVICE__ -inline float floorf(float __x) { return __ocml_floor_f32(__x); } +float floorf(float __x) { return __ocml_floor_f32(__x); } + __DEVICE__ -inline float fmaf(float __x, float __y, float __z) { +float fmaf(float __x, float __y, float __z) { return __ocml_fma_f32(__x, __y, __z); } + __DEVICE__ -inline float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); } +float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); } + __DEVICE__ -inline float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); } +float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); } + __DEVICE__ -inline float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } +float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } + __DEVICE__ -inline float frexpf(float __x, int *__nptr) { +float frexpf(float __x, int *__nptr) { int __tmp; float __r = __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp); @@ -205,24 +253,31 @@ inline float frexpf(float __x, int *__nptr) { return __r; } + __DEVICE__ -inline float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); } +float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); } + __DEVICE__ -inline int ilogbf(float __x) { return __ocml_ilogb_f32(__x); } +int ilogbf(float __x) { return __ocml_ilogb_f32(__x); } + __DEVICE__ -inline __RETURN_TYPE isfinite(float __x) { return __ocml_isfinite_f32(__x); } +__RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); } + __DEVICE__ -inline __RETURN_TYPE isinf(float __x) { return __ocml_isinf_f32(__x); } +__RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); } + __DEVICE__ -inline __RETURN_TYPE isnan(float __x) { return __ocml_isnan_f32(__x); } +__RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); } + __DEVICE__ -inline float j0f(float __x) { return __ocml_j0_f32(__x); } +float j0f(float __x) { return __ocml_j0_f32(__x); } + __DEVICE__ -inline float j1f(float __x) { return __ocml_j1_f32(__x); } +float j1f(float __x) { return __ocml_j1_f32(__x); } + __DEVICE__ -inline float jnf(int __n, - float __x) { // TODO: we could use Ahmes multiplication - // and the Miller & Brown algorithm +float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication + // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. if (__n == 0) @@ -240,46 +295,58 @@ inline float jnf(int __n, return __x1; } + __DEVICE__ -inline float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); } +float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); } + __DEVICE__ -inline float lgammaf(float __x) { return __ocml_lgamma_f32(__x); } +float lgammaf(float __x) { return __ocml_lgamma_f32(__x); } + __DEVICE__ -inline long long int llrintf(float __x) { return __ocml_rint_f32(__x); } +long long int llrintf(float __x) { return __ocml_rint_f32(__x); } + __DEVICE__ -inline long long int llroundf(float __x) { return __ocml_round_f32(__x); } +long long int llroundf(float __x) { return __ocml_round_f32(__x); } + __DEVICE__ -inline float log10f(float __x) { return __ocml_log10_f32(__x); } +float log10f(float __x) { return __ocml_log10_f32(__x); } + __DEVICE__ -inline float log1pf(float __x) { return __ocml_log1p_f32(__x); } +float log1pf(float __x) { return __ocml_log1p_f32(__x); } + __DEVICE__ -inline float log2f(float __x) { return __ocml_log2_f32(__x); } +float log2f(float __x) { return __ocml_log2_f32(__x); } + __DEVICE__ -inline float logbf(float __x) { return __ocml_logb_f32(__x); } +float logbf(float __x) { return __ocml_logb_f32(__x); } + __DEVICE__ -inline float logf(float __x) { return __ocml_log_f32(__x); } +float logf(float __x) { return __ocml_log_f32(__x); } + __DEVICE__ -inline long int lrintf(float __x) { return __ocml_rint_f32(__x); } +long int lrintf(float __x) { return __ocml_rint_f32(__x); } + __DEVICE__ -inline long int lroundf(float __x) { return __ocml_round_f32(__x); } +long int lroundf(float __x) { return __ocml_round_f32(__x); } + __DEVICE__ -inline float modff(float __x, float *__iptr) { +float modff(float __x, float *__iptr) { float __tmp; float __r = __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp); *__iptr = __tmp; - return __r; } + __DEVICE__ -inline float nanf(const char *__tagp) { +float nanf(const char *__tagp) { union { float val; struct ieee_float { - uint32_t mantissa : 22; - uint32_t quiet : 1; - uint32_t exponent : 8; - uint32_t sign : 1; + unsigned int mantissa : 22; + unsigned int quiet : 1; + unsigned int exponent : 8; + unsigned int sign : 1; } bits; } __tmp; __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits)); @@ -291,28 +358,34 @@ inline float nanf(const char *__tagp) { return __tmp.val; } + __DEVICE__ -inline float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); } +float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); } + __DEVICE__ -inline float nextafterf(float __x, float __y) { +float nextafterf(float __x, float __y) { return __ocml_nextafter_f32(__x, __y); } + __DEVICE__ -inline float norm3df(float __x, float __y, float __z) { +float norm3df(float __x, float __y, float __z) { return __ocml_len3_f32(__x, __y, __z); } + __DEVICE__ -inline float norm4df(float __x, float __y, float __z, float __w) { +float norm4df(float __x, float __y, float __z, float __w) { return __ocml_len4_f32(__x, __y, __z, __w); } + __DEVICE__ -inline float normcdff(float __x) { return __ocml_ncdf_f32(__x); } +float normcdff(float __x) { return __ocml_ncdf_f32(__x); } + __DEVICE__ -inline float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); } +float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); } + __DEVICE__ -inline float -normf(int __dim, - const float *__a) { // TODO: placeholder until OCML adds support. +float normf(int __dim, + const float *__a) { // TODO: placeholder until OCML adds support. float __r = 0; while (__dim--) { __r += __a[0] * __a[0]; @@ -321,18 +394,23 @@ normf(int __dim, return __ocml_sqrt_f32(__r); } + __DEVICE__ -inline float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); } +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); } +float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); } + __DEVICE__ -inline float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); } +float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); } + __DEVICE__ -inline float remainderf(float __x, float __y) { +float remainderf(float __x, float __y) { return __ocml_remainder_f32(__x, __y); } + __DEVICE__ -inline float remquof(float __x, float __y, int *__quo) { +float remquof(float __x, float __y, int *__quo) { int __tmp; float __r = __ocml_remquo_f32( __x, __y, (__attribute__((address_space(5))) int *)&__tmp); @@ -340,25 +418,26 @@ inline float remquof(float __x, float __y, int *__quo) { return __r; } + __DEVICE__ -inline float rhypotf(float __x, float __y) { - return __ocml_rhypot_f32(__x, __y); -} +float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); } + __DEVICE__ -inline float rintf(float __x) { return __ocml_rint_f32(__x); } +float rintf(float __x) { return __ocml_rint_f32(__x); } + __DEVICE__ -inline float rnorm3df(float __x, float __y, float __z) { +float rnorm3df(float __x, float __y, float __z) { return __ocml_rlen3_f32(__x, __y, __z); } __DEVICE__ -inline float rnorm4df(float __x, float __y, float __z, float __w) { +float rnorm4df(float __x, float __y, float __z, float __w) { return __ocml_rlen4_f32(__x, __y, __z, __w); } + __DEVICE__ -inline float -rnormf(int __dim, - const float *__a) { // TODO: placeholder until OCML adds support. +float rnormf(int __dim, + const float *__a) { // TODO: placeholder until OCML adds support. float __r = 0; while (__dim--) { __r += __a[0] * __a[0]; @@ -367,59 +446,74 @@ rnormf(int __dim, return __ocml_rsqrt_f32(__r); } + __DEVICE__ -inline float roundf(float __x) { return __ocml_round_f32(__x); } +float roundf(float __x) { return __ocml_round_f32(__x); } + __DEVICE__ -inline float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); } +float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); } + __DEVICE__ -inline float scalblnf(float __x, long int __n) { +float scalblnf(float __x, long int __n) { return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n) : __ocml_scalb_f32(__x, __n); } + __DEVICE__ -inline float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); } +float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); } + __DEVICE__ -inline __RETURN_TYPE signbit(float __x) { return __ocml_signbit_f32(__x); } +__RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); } + __DEVICE__ -inline void sincosf(float __x, float *__sinptr, float *__cosptr) { +void sincosf(float __x, float *__sinptr, float *__cosptr) { float __tmp; - *__sinptr = __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp); *__cosptr = __tmp; } + __DEVICE__ -inline void sincospif(float __x, float *__sinptr, float *__cosptr) { +void sincospif(float __x, float *__sinptr, float *__cosptr) { float __tmp; - *__sinptr = __ocml_sincospi_f32( __x, (__attribute__((address_space(5))) float *)&__tmp); *__cosptr = __tmp; } + __DEVICE__ -inline float sinf(float __x) { return __ocml_sin_f32(__x); } +float sinf(float __x) { return __ocml_sin_f32(__x); } + __DEVICE__ -inline float sinhf(float __x) { return __ocml_sinh_f32(__x); } +float sinhf(float __x) { return __ocml_sinh_f32(__x); } + __DEVICE__ -inline float sinpif(float __x) { return __ocml_sinpi_f32(__x); } +float sinpif(float __x) { return __ocml_sinpi_f32(__x); } + __DEVICE__ -inline float sqrtf(float __x) { return __ocml_sqrt_f32(__x); } +float sqrtf(float __x) { return __ocml_sqrt_f32(__x); } + __DEVICE__ -inline float tanf(float __x) { return __ocml_tan_f32(__x); } +float tanf(float __x) { return __ocml_tan_f32(__x); } + __DEVICE__ -inline float tanhf(float __x) { return __ocml_tanh_f32(__x); } +float tanhf(float __x) { return __ocml_tanh_f32(__x); } + __DEVICE__ -inline float tgammaf(float __x) { return __ocml_tgamma_f32(__x); } +float tgammaf(float __x) { return __ocml_tgamma_f32(__x); } + __DEVICE__ -inline float truncf(float __x) { return __ocml_trunc_f32(__x); } +float truncf(float __x) { return __ocml_trunc_f32(__x); } + __DEVICE__ -inline float y0f(float __x) { return __ocml_y0_f32(__x); } +float y0f(float __x) { return __ocml_y0_f32(__x); } + __DEVICE__ -inline float y1f(float __x) { return __ocml_y1_f32(__x); } +float y1f(float __x) { return __ocml_y1_f32(__x); } + __DEVICE__ -inline float ynf(int __n, - float __x) { // TODO: we could use Ahmes multiplication - // and the Miller & Brown algorithm +float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication + // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. @@ -440,292 +534,328 @@ inline float ynf(int __n, } // BEGIN INTRINSICS + __DEVICE__ -inline float __cosf(float __x) { return __ocml_native_cos_f32(__x); } +float __cosf(float __x) { return __ocml_native_cos_f32(__x); } + __DEVICE__ -inline float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); } +float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); } + __DEVICE__ -inline float __expf(float __x) { return __ocml_native_exp_f32(__x); } +float __expf(float __x) { return __ocml_native_exp_f32(__x); } + #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fadd_rd(float __x, float __y) { - return __ocml_add_rtn_f32(__x, __y); -} +float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); } #endif __DEVICE__ -inline float __fadd_rn(float __x, float __y) { return __x + __y; } +float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fadd_ru(float __x, float __y) { - return __ocml_add_rtp_f32(__x, __y); -} +float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); } + __DEVICE__ -inline float __fadd_rz(float __x, float __y) { - return __ocml_add_rtz_f32(__x, __y); -} +float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); } + __DEVICE__ -inline float __fdiv_rd(float __x, float __y) { - return __ocml_div_rtn_f32(__x, __y); -} +float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); } #endif __DEVICE__ -inline float __fdiv_rn(float __x, float __y) { return __x / __y; } +float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fdiv_ru(float __x, float __y) { - return __ocml_div_rtp_f32(__x, __y); -} +float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); } + __DEVICE__ -inline float __fdiv_rz(float __x, float __y) { - return __ocml_div_rtz_f32(__x, __y); -} +float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); } #endif __DEVICE__ -inline float __fdividef(float __x, float __y) { return __x / __y; } +float __fdividef(float __x, float __y) { return __x / __y; } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fmaf_rd(float __x, float __y, float __z) { +float __fmaf_rd(float __x, float __y, float __z) { return __ocml_fma_rtn_f32(__x, __y, __z); } #endif __DEVICE__ -inline float __fmaf_rn(float __x, float __y, float __z) { - return __ocml_fma_f32(__x, __y, __z); +float __fmaf_rn(float __x, float __y, float __z) { + return __ocml_fma_rte_f32(__x, __y, __z); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fmaf_ru(float __x, float __y, float __z) { +float __fmaf_ru(float __x, float __y, float __z) { return __ocml_fma_rtp_f32(__x, __y, __z); } + __DEVICE__ -inline float __fmaf_rz(float __x, float __y, float __z) { +float __fmaf_rz(float __x, float __y, float __z) { return __ocml_fma_rtz_f32(__x, __y, __z); } + __DEVICE__ -inline float __fmul_rd(float __x, float __y) { - return __ocml_mul_rtn_f32(__x, __y); -} +float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); } #endif __DEVICE__ -inline float __fmul_rn(float __x, float __y) { return __x * __y; } +float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fmul_ru(float __x, float __y) { - return __ocml_mul_rtp_f32(__x, __y); -} +float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); } + __DEVICE__ -inline float __fmul_rz(float __x, float __y) { - return __ocml_mul_rtz_f32(__x, __y); -} +float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); } + __DEVICE__ -inline float __frcp_rd(float __x) { return __llvm_amdgcn_rcp_f32(__x); } +float __frcp_rd(float __x) { return __llvm_amdgcn_rcp_f32(__x); } #endif __DEVICE__ -inline float __frcp_rn(float __x) { return __llvm_amdgcn_rcp_f32(__x); } +float __frcp_rn(float __x) { return __llvm_amdgcn_rcp_f32(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __frcp_ru(float __x) { return __llvm_amdgcn_rcp_f32(__x); } +float __frcp_ru(float __x) { return __llvm_amdgcn_rcp_f32(__x); } + __DEVICE__ -inline float __frcp_rz(float __x) { return __llvm_amdgcn_rcp_f32(__x); } +float __frcp_rz(float __x) { return __llvm_amdgcn_rcp_f32(__x); } #endif __DEVICE__ -inline float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); } +float __frsqrt_rn(float __x) { return __llvm_amdgcn_rsq_f32(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); } +float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); } #endif __DEVICE__ -inline float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); } +float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); } +float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); } + __DEVICE__ -inline float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); } +float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); } + __DEVICE__ -inline float __fsub_rd(float __x, float __y) { - return __ocml_sub_rtn_f32(__x, __y); -} +float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); } #endif __DEVICE__ -inline float __fsub_rn(float __x, float __y) { return __x - __y; } +float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline float __fsub_ru(float __x, float __y) { - return __ocml_sub_rtp_f32(__x, __y); -} +float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); } + __DEVICE__ -inline float __fsub_rz(float __x, float __y) { - return __ocml_sub_rtz_f32(__x, __y); -} +float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); } #endif __DEVICE__ -inline float __log10f(float __x) { return __ocml_native_log10_f32(__x); } +float __log10f(float __x) { return __ocml_native_log10_f32(__x); } + __DEVICE__ -inline float __log2f(float __x) { return __ocml_native_log2_f32(__x); } +float __log2f(float __x) { return __ocml_native_log2_f32(__x); } + __DEVICE__ -inline float __logf(float __x) { return __ocml_native_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); } +float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); } + __DEVICE__ -inline float __saturatef(float __x) { - return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); -} +float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); } + __DEVICE__ -inline void __sincosf(float __x, float *__sinptr, float *__cosptr) { +void __sincosf(float __x, float *__sinptr, float *__cosptr) { *__sinptr = __ocml_native_sin_f32(__x); *__cosptr = __ocml_native_cos_f32(__x); } + __DEVICE__ -inline float __sinf(float __x) { return __ocml_native_sin_f32(__x); } +float __sinf(float __x) { return __ocml_native_sin_f32(__x); } + __DEVICE__ -inline float __tanf(float __x) { return __ocml_tan_f32(__x); } +float __tanf(float __x) { return __ocml_tan_f32(__x); } // END INTRINSICS // END FLOAT // BEGIN DOUBLE -#ifdef __cplusplus -__DEVICE__ -inline double abs(double __x) { return __ocml_fabs_f64(__x); } -#endif __DEVICE__ -inline double acos(double __x) { return __ocml_acos_f64(__x); } +double acos(double __x) { return __ocml_acos_f64(__x); } + __DEVICE__ -inline double acosh(double __x) { return __ocml_acosh_f64(__x); } +double acosh(double __x) { return __ocml_acosh_f64(__x); } + __DEVICE__ -inline double asin(double __x) { return __ocml_asin_f64(__x); } +double asin(double __x) { return __ocml_asin_f64(__x); } + __DEVICE__ -inline double asinh(double __x) { return __ocml_asinh_f64(__x); } +double asinh(double __x) { return __ocml_asinh_f64(__x); } + __DEVICE__ -inline double atan(double __x) { return __ocml_atan_f64(__x); } +double atan(double __x) { return __ocml_atan_f64(__x); } + __DEVICE__ -inline double atan2(double __x, double __y) { - return __ocml_atan2_f64(__x, __y); -} +double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); } + __DEVICE__ -inline double atanh(double __x) { return __ocml_atanh_f64(__x); } +double atanh(double __x) { return __ocml_atanh_f64(__x); } + __DEVICE__ -inline double cbrt(double __x) { return __ocml_cbrt_f64(__x); } +double cbrt(double __x) { return __ocml_cbrt_f64(__x); } + __DEVICE__ -inline double ceil(double __x) { return __ocml_ceil_f64(__x); } +double ceil(double __x) { return __ocml_ceil_f64(__x); } + __DEVICE__ -inline double copysign(double __x, double __y) { +double copysign(double __x, double __y) { return __ocml_copysign_f64(__x, __y); } + __DEVICE__ -inline double cos(double __x) { return __ocml_cos_f64(__x); } +double cos(double __x) { return __ocml_cos_f64(__x); } + __DEVICE__ -inline double cosh(double __x) { return __ocml_cosh_f64(__x); } +double cosh(double __x) { return __ocml_cosh_f64(__x); } + __DEVICE__ -inline double cospi(double __x) { return __ocml_cospi_f64(__x); } +double cospi(double __x) { return __ocml_cospi_f64(__x); } + __DEVICE__ -inline double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); } +double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); } + __DEVICE__ -inline double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); } +double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); } + __DEVICE__ -inline double erf(double __x) { return __ocml_erf_f64(__x); } +double erf(double __x) { return __ocml_erf_f64(__x); } + __DEVICE__ -inline double erfc(double __x) { return __ocml_erfc_f64(__x); } +double erfc(double __x) { return __ocml_erfc_f64(__x); } + __DEVICE__ -inline double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); } +double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); } + __DEVICE__ -inline double erfcx(double __x) { return __ocml_erfcx_f64(__x); } +double erfcx(double __x) { return __ocml_erfcx_f64(__x); } + __DEVICE__ -inline double erfinv(double __x) { return __ocml_erfinv_f64(__x); } +double erfinv(double __x) { return __ocml_erfinv_f64(__x); } + __DEVICE__ -inline double exp(double __x) { return __ocml_exp_f64(__x); } +double exp(double __x) { return __ocml_exp_f64(__x); } + __DEVICE__ -inline double exp10(double __x) { return __ocml_exp10_f64(__x); } +double exp10(double __x) { return __ocml_exp10_f64(__x); } + __DEVICE__ -inline double exp2(double __x) { return __ocml_exp2_f64(__x); } +double exp2(double __x) { return __ocml_exp2_f64(__x); } + __DEVICE__ -inline double expm1(double __x) { return __ocml_expm1_f64(__x); } +double expm1(double __x) { return __ocml_expm1_f64(__x); } + __DEVICE__ -inline double fabs(double __x) { return __ocml_fabs_f64(__x); } +double fabs(double __x) { return __ocml_fabs_f64(__x); } + __DEVICE__ -inline double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); } +double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); } + __DEVICE__ -inline double floor(double __x) { return __ocml_floor_f64(__x); } +double floor(double __x) { return __ocml_floor_f64(__x); } + __DEVICE__ -inline double fma(double __x, double __y, double __z) { +double fma(double __x, double __y, double __z) { return __ocml_fma_f64(__x, __y, __z); } + __DEVICE__ -inline double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); } +double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); } + __DEVICE__ -inline double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); } +double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); } + __DEVICE__ -inline double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } +double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } + __DEVICE__ -inline double frexp(double __x, int *__nptr) { +double frexp(double __x, int *__nptr) { int __tmp; double __r = __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp); *__nptr = __tmp; - return __r; } + __DEVICE__ -inline double hypot(double __x, double __y) { - return __ocml_hypot_f64(__x, __y); -} +double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); } + __DEVICE__ -inline int ilogb(double __x) { return __ocml_ilogb_f64(__x); } +int ilogb(double __x) { return __ocml_ilogb_f64(__x); } + __DEVICE__ -inline __RETURN_TYPE isfinite(double __x) { return __ocml_isfinite_f64(__x); } +__RETURN_TYPE __finite(double __x) { return __ocml_isfinite_f64(__x); } + __DEVICE__ -inline __RETURN_TYPE isinf(double __x) { return __ocml_isinf_f64(__x); } +__RETURN_TYPE __isinf(double __x) { return __ocml_isinf_f64(__x); } + __DEVICE__ -inline __RETURN_TYPE isnan(double __x) { return __ocml_isnan_f64(__x); } +__RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); } + __DEVICE__ -inline double j0(double __x) { return __ocml_j0_f64(__x); } +double j0(double __x) { return __ocml_j0_f64(__x); } + __DEVICE__ -inline double j1(double __x) { return __ocml_j1_f64(__x); } +double j1(double __x) { return __ocml_j1_f64(__x); } + __DEVICE__ -inline double jn(int __n, - double __x) { // TODO: we could use Ahmes multiplication - // and the Miller & Brown algorithm +double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication + // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. if (__n == 0) - return j0f(__x); + return j0(__x); if (__n == 1) - return j1f(__x); + return j1(__x); - double __x0 = j0f(__x); - double __x1 = j1f(__x); + double __x0 = j0(__x); + double __x1 = j1(__x); for (int __i = 1; __i < __n; ++__i) { double __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; __x1 = __x2; } - return __x1; } + __DEVICE__ -inline double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); } +double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); } + __DEVICE__ -inline double lgamma(double __x) { return __ocml_lgamma_f64(__x); } +double lgamma(double __x) { return __ocml_lgamma_f64(__x); } + __DEVICE__ -inline long long int llrint(double __x) { return __ocml_rint_f64(__x); } +long long int llrint(double __x) { return __ocml_rint_f64(__x); } + __DEVICE__ -inline long long int llround(double __x) { return __ocml_round_f64(__x); } +long long int llround(double __x) { return __ocml_round_f64(__x); } + __DEVICE__ -inline double log(double __x) { return __ocml_log_f64(__x); } +double log(double __x) { return __ocml_log_f64(__x); } + __DEVICE__ -inline double log10(double __x) { return __ocml_log10_f64(__x); } +double log10(double __x) { return __ocml_log10_f64(__x); } + __DEVICE__ -inline double log1p(double __x) { return __ocml_log1p_f64(__x); } +double log1p(double __x) { return __ocml_log1p_f64(__x); } + __DEVICE__ -inline double log2(double __x) { return __ocml_log2_f64(__x); } +double log2(double __x) { return __ocml_log2_f64(__x); } + __DEVICE__ -inline double logb(double __x) { return __ocml_logb_f64(__x); } +double logb(double __x) { return __ocml_logb_f64(__x); } + __DEVICE__ -inline long int lrint(double __x) { return __ocml_rint_f64(__x); } +long int lrint(double __x) { return __ocml_rint_f64(__x); } + __DEVICE__ -inline long int lround(double __x) { return __ocml_round_f64(__x); } +long int lround(double __x) { return __ocml_round_f64(__x); } + __DEVICE__ -inline double modf(double __x, double *__iptr) { +double modf(double __x, double *__iptr) { double __tmp; double __r = __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp); @@ -733,8 +863,9 @@ inline double modf(double __x, double *__iptr) { return __r; } + __DEVICE__ -inline double nan(const char *__tagp) { +double nan(const char *__tagp) { #if !_WIN32 union { double val; @@ -755,21 +886,23 @@ inline double nan(const char *__tagp) { return __tmp.val; #else __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double)); - uint64_t val = __make_mantissa(__tagp); - val |= 0xFFF << 51; - return *reinterpret_cast<double *>(&val); + uint64_t __val = __make_mantissa(__tagp); + __val |= 0xFFF << 51; + return *reinterpret_cast<double *>(&__val); #endif } + __DEVICE__ -inline double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); } +double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); } + __DEVICE__ -inline double nextafter(double __x, double __y) { +double nextafter(double __x, double __y) { return __ocml_nextafter_f64(__x, __y); } + __DEVICE__ -inline double -norm(int __dim, - const double *__a) { // TODO: placeholder until OCML adds support. +double norm(int __dim, + const double *__a) { // TODO: placeholder until OCML adds support. double __r = 0; while (__dim--) { __r += __a[0] * __a[0]; @@ -778,30 +911,39 @@ norm(int __dim, return __ocml_sqrt_f64(__r); } + __DEVICE__ -inline double norm3d(double __x, double __y, double __z) { +double norm3d(double __x, double __y, double __z) { return __ocml_len3_f64(__x, __y, __z); } + __DEVICE__ -inline double norm4d(double __x, double __y, double __z, double __w) { +double norm4d(double __x, double __y, double __z, double __w) { return __ocml_len4_f64(__x, __y, __z, __w); } + __DEVICE__ -inline double normcdf(double __x) { return __ocml_ncdf_f64(__x); } +double normcdf(double __x) { return __ocml_ncdf_f64(__x); } + __DEVICE__ -inline double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); } +double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); } + __DEVICE__ -inline double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); } +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); } +double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); } + __DEVICE__ -inline double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); } +double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); } + __DEVICE__ -inline double remainder(double __x, double __y) { +double remainder(double __x, double __y) { return __ocml_remainder_f64(__x, __y); } + __DEVICE__ -inline double remquo(double __x, double __y, int *__quo) { +double remquo(double __x, double __y, int *__quo) { int __tmp; double __r = __ocml_remquo_f64( __x, __y, (__attribute__((address_space(5))) int *)&__tmp); @@ -809,16 +951,16 @@ inline double remquo(double __x, double __y, int *__quo) { return __r; } + __DEVICE__ -inline double rhypot(double __x, double __y) { - return __ocml_rhypot_f64(__x, __y); -} +double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); } + __DEVICE__ -inline double rint(double __x) { return __ocml_rint_f64(__x); } +double rint(double __x) { return __ocml_rint_f64(__x); } + __DEVICE__ -inline double -rnorm(int __dim, - const double *__a) { // TODO: placeholder until OCML adds support. +double rnorm(int __dim, + const double *__a) { // TODO: placeholder until OCML adds support. double __r = 0; while (__dim--) { __r += __a[0] * __a[0]; @@ -827,77 +969,93 @@ rnorm(int __dim, return __ocml_rsqrt_f64(__r); } + __DEVICE__ -inline double rnorm3d(double __x, double __y, double __z) { +double rnorm3d(double __x, double __y, double __z) { return __ocml_rlen3_f64(__x, __y, __z); } + __DEVICE__ -inline double rnorm4d(double __x, double __y, double __z, double __w) { +double rnorm4d(double __x, double __y, double __z, double __w) { return __ocml_rlen4_f64(__x, __y, __z, __w); } + __DEVICE__ -inline double round(double __x) { return __ocml_round_f64(__x); } +double round(double __x) { return __ocml_round_f64(__x); } + __DEVICE__ -inline double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); } +double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); } + __DEVICE__ -inline double scalbln(double __x, long int __n) { +double scalbln(double __x, long int __n) { return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n) : __ocml_scalb_f64(__x, __n); } __DEVICE__ -inline double scalbn(double __x, int __n) { - return __ocml_scalbn_f64(__x, __n); -} +double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); } + __DEVICE__ -inline __RETURN_TYPE signbit(double __x) { return __ocml_signbit_f64(__x); } +__RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); } + __DEVICE__ -inline double sin(double __x) { return __ocml_sin_f64(__x); } +double sin(double __x) { return __ocml_sin_f64(__x); } + __DEVICE__ -inline void sincos(double __x, double *__sinptr, double *__cosptr) { +void sincos(double __x, double *__sinptr, double *__cosptr) { double __tmp; *__sinptr = __ocml_sincos_f64( __x, (__attribute__((address_space(5))) double *)&__tmp); *__cosptr = __tmp; } + __DEVICE__ -inline void sincospi(double __x, double *__sinptr, double *__cosptr) { +void sincospi(double __x, double *__sinptr, double *__cosptr) { double __tmp; *__sinptr = __ocml_sincospi_f64( __x, (__attribute__((address_space(5))) double *)&__tmp); *__cosptr = __tmp; } + __DEVICE__ -inline double sinh(double __x) { return __ocml_sinh_f64(__x); } +double sinh(double __x) { return __ocml_sinh_f64(__x); } + __DEVICE__ -inline double sinpi(double __x) { return __ocml_sinpi_f64(__x); } +double sinpi(double __x) { return __ocml_sinpi_f64(__x); } + __DEVICE__ -inline double sqrt(double __x) { return __ocml_sqrt_f64(__x); } +double sqrt(double __x) { return __ocml_sqrt_f64(__x); } + __DEVICE__ -inline double tan(double __x) { return __ocml_tan_f64(__x); } +double tan(double __x) { return __ocml_tan_f64(__x); } + __DEVICE__ -inline double tanh(double __x) { return __ocml_tanh_f64(__x); } +double tanh(double __x) { return __ocml_tanh_f64(__x); } + __DEVICE__ -inline double tgamma(double __x) { return __ocml_tgamma_f64(__x); } +double tgamma(double __x) { return __ocml_tgamma_f64(__x); } + __DEVICE__ -inline double trunc(double __x) { return __ocml_trunc_f64(__x); } +double trunc(double __x) { return __ocml_trunc_f64(__x); } + __DEVICE__ -inline double y0(double __x) { return __ocml_y0_f64(__x); } +double y0(double __x) { return __ocml_y0_f64(__x); } + __DEVICE__ -inline double y1(double __x) { return __ocml_y1_f64(__x); } +double y1(double __x) { return __ocml_y1_f64(__x); } + __DEVICE__ -inline double yn(int __n, - double __x) { // TODO: we could use Ahmes multiplication - // and the Miller & Brown algorithm +double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication + // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. if (__n == 0) - return j0f(__x); + return y0(__x); if (__n == 1) - return j1f(__x); + return y1(__x); - double __x0 = j0f(__x); - double __x1 = j1f(__x); + double __x0 = y0(__x); + double __x1 = y1(__x); for (int __i = 1; __i < __n; ++__i) { double __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; @@ -910,297 +1068,168 @@ inline double yn(int __n, // BEGIN INTRINSICS #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __dadd_rd(double __x, double __y) { +double __dadd_rd(double __x, double __y) { return __ocml_add_rtn_f64(__x, __y); } #endif __DEVICE__ -inline double __dadd_rn(double __x, double __y) { return __x + __y; } +double __dadd_rn(double __x, double __y) { + return __ocml_add_rte_f64(__x, __y); +} #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __dadd_ru(double __x, double __y) { +double __dadd_ru(double __x, double __y) { return __ocml_add_rtp_f64(__x, __y); } + __DEVICE__ -inline double __dadd_rz(double __x, double __y) { +double __dadd_rz(double __x, double __y) { return __ocml_add_rtz_f64(__x, __y); } + __DEVICE__ -inline double __ddiv_rd(double __x, double __y) { +double __ddiv_rd(double __x, double __y) { return __ocml_div_rtn_f64(__x, __y); } #endif __DEVICE__ -inline double __ddiv_rn(double __x, double __y) { return __x / __y; } +double __ddiv_rn(double __x, double __y) { + return __ocml_div_rte_f64(__x, __y); +} #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __ddiv_ru(double __x, double __y) { +double __ddiv_ru(double __x, double __y) { return __ocml_div_rtp_f64(__x, __y); } + __DEVICE__ -inline double __ddiv_rz(double __x, double __y) { +double __ddiv_rz(double __x, double __y) { return __ocml_div_rtz_f64(__x, __y); } + __DEVICE__ -inline double __dmul_rd(double __x, double __y) { +double __dmul_rd(double __x, double __y) { return __ocml_mul_rtn_f64(__x, __y); } #endif __DEVICE__ -inline double __dmul_rn(double __x, double __y) { return __x * __y; } +double __dmul_rn(double __x, double __y) { + return __ocml_mul_rte_f64(__x, __y); +} #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __dmul_ru(double __x, double __y) { +double __dmul_ru(double __x, double __y) { return __ocml_mul_rtp_f64(__x, __y); } + __DEVICE__ -inline double __dmul_rz(double __x, double __y) { +double __dmul_rz(double __x, double __y) { return __ocml_mul_rtz_f64(__x, __y); } + __DEVICE__ -inline double __drcp_rd(double __x) { return __llvm_amdgcn_rcp_f64(__x); } +double __drcp_rd(double __x) { return __llvm_amdgcn_rcp_f64(__x); } #endif __DEVICE__ -inline double __drcp_rn(double __x) { return __llvm_amdgcn_rcp_f64(__x); } +double __drcp_rn(double __x) { return __llvm_amdgcn_rcp_f64(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __drcp_ru(double __x) { return __llvm_amdgcn_rcp_f64(__x); } +double __drcp_ru(double __x) { return __llvm_amdgcn_rcp_f64(__x); } + __DEVICE__ -inline double __drcp_rz(double __x) { return __llvm_amdgcn_rcp_f64(__x); } +double __drcp_rz(double __x) { return __llvm_amdgcn_rcp_f64(__x); } + __DEVICE__ -inline double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); } +double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); } #endif __DEVICE__ -inline double __dsqrt_rn(double __x) { return __ocml_sqrt_f64(__x); } +double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); } +double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); } + __DEVICE__ -inline double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); } +double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); } + __DEVICE__ -inline double __dsub_rd(double __x, double __y) { +double __dsub_rd(double __x, double __y) { return __ocml_sub_rtn_f64(__x, __y); } #endif __DEVICE__ -inline double __dsub_rn(double __x, double __y) { return __x - __y; } +double __dsub_rn(double __x, double __y) { + return __ocml_sub_rte_f64(__x, __y); +} #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __dsub_ru(double __x, double __y) { +double __dsub_ru(double __x, double __y) { return __ocml_sub_rtp_f64(__x, __y); } + __DEVICE__ -inline double __dsub_rz(double __x, double __y) { +double __dsub_rz(double __x, double __y) { return __ocml_sub_rtz_f64(__x, __y); } + __DEVICE__ -inline double __fma_rd(double __x, double __y, double __z) { +double __fma_rd(double __x, double __y, double __z) { return __ocml_fma_rtn_f64(__x, __y, __z); } #endif __DEVICE__ -inline double __fma_rn(double __x, double __y, double __z) { - return __ocml_fma_f64(__x, __y, __z); +double __fma_rn(double __x, double __y, double __z) { + return __ocml_fma_rte_f64(__x, __y, __z); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ -inline double __fma_ru(double __x, double __y, double __z) { +double __fma_ru(double __x, double __y, double __z) { return __ocml_fma_rtp_f64(__x, __y, __z); } + __DEVICE__ -inline double __fma_rz(double __x, double __y, double __z) { +double __fma_rz(double __x, double __y, double __z) { return __ocml_fma_rtz_f64(__x, __y, __z); } #endif // END INTRINSICS // END DOUBLE -// BEGIN INTEGER -__DEVICE__ -inline int abs(int __x) { - int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1); - return (__x ^ __sgn) - __sgn; -} -__DEVICE__ -inline long labs(long __x) { - long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1); - return (__x ^ __sgn) - __sgn; -} -__DEVICE__ -inline long long llabs(long long __x) { - long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1); - return (__x ^ __sgn) - __sgn; -} +// C only macros +#if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L +#define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x) +#define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x) +#define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x) +#define signbit(__x) \ + _Generic((__x), float : __signbitf, double : __signbit)(__x) +#endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L #if defined(__cplusplus) -__DEVICE__ -inline long abs(long __x) { return labs(__x); } -__DEVICE__ -inline long long abs(long long __x) { return llabs(__x); } -#endif -// END INTEGER - -__DEVICE__ -inline _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) { - return __ocml_fma_f16(__x, __y, __z); -} - -__DEVICE__ -inline float fma(float __x, float __y, float __z) { - return fmaf(__x, __y, __z); -} - -#pragma push_macro("__DEF_FUN1") -#pragma push_macro("__DEF_FUN2") -#pragma push_macro("__DEF_FUNI") -#pragma push_macro("__DEF_FLOAT_FUN2I") -#pragma push_macro("__HIP_OVERLOAD1") -#pragma push_macro("__HIP_OVERLOAD2") - -// __hip_enable_if::type is a type function which returns __T if __B is true. -template <bool __B, class __T = void> struct __hip_enable_if {}; - -template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; }; - -// __HIP_OVERLOAD1 is used to resolve function calls with integer argument to -// avoid compilation error due to ambibuity. e.g. floor(5) is resolved with -// floor(double). -#define __HIP_OVERLOAD1(__retty, __fn) \ - template <typename __T> \ - __DEVICE__ typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \ - __retty>::type \ - __fn(__T __x) { \ - return ::__fn((double)__x); \ - } - -// __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double -// or integer argument to avoid compilation error due to ambibuity. e.g. -// max(5.0f, 6.0) is resolved with max(double, double). -#define __HIP_OVERLOAD2(__retty, __fn) \ - template <typename __T1, typename __T2> \ - __DEVICE__ \ - typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && \ - std::numeric_limits<__T2>::is_specialized, \ - __retty>::type \ - __fn(__T1 __x, __T2 __y) { \ - return __fn((double)__x, (double)__y); \ - } - -// Define cmath functions with float argument and returns float. -#define __DEF_FUN1(__retty, __func) \ - __DEVICE__ \ - inline float __func(float __x) { return __func##f(__x); } \ - __HIP_OVERLOAD1(__retty, __func) - -// Define cmath functions with float argument and returns __retty. -#define __DEF_FUNI(__retty, __func) \ - __DEVICE__ \ - inline __retty __func(float __x) { return __func##f(__x); } \ - __HIP_OVERLOAD1(__retty, __func) - -// define cmath functions with two float arguments. -#define __DEF_FUN2(__retty, __func) \ - __DEVICE__ \ - inline float __func(float __x, float __y) { return __func##f(__x, __y); } \ - __HIP_OVERLOAD2(__retty, __func) - -__DEF_FUN1(double, acos) -__DEF_FUN1(double, acosh) -__DEF_FUN1(double, asin) -__DEF_FUN1(double, asinh) -__DEF_FUN1(double, atan) -__DEF_FUN2(double, atan2); -__DEF_FUN1(double, atanh) -__DEF_FUN1(double, cbrt) -__DEF_FUN1(double, ceil) -__DEF_FUN2(double, copysign); -__DEF_FUN1(double, cos) -__DEF_FUN1(double, cosh) -__DEF_FUN1(double, erf) -__DEF_FUN1(double, erfc) -__DEF_FUN1(double, exp) -__DEF_FUN1(double, exp2) -__DEF_FUN1(double, expm1) -__DEF_FUN1(double, fabs) -__DEF_FUN2(double, fdim); -__DEF_FUN1(double, floor) -__DEF_FUN2(double, fmax); -__DEF_FUN2(double, fmin); -__DEF_FUN2(double, fmod); -//__HIP_OVERLOAD1(int, fpclassify) -__DEF_FUN2(double, hypot); -__DEF_FUNI(int, ilogb) -__HIP_OVERLOAD1(bool, isfinite) -__HIP_OVERLOAD2(bool, isgreater); -__HIP_OVERLOAD2(bool, isgreaterequal); -__HIP_OVERLOAD1(bool, isinf); -__HIP_OVERLOAD2(bool, isless); -__HIP_OVERLOAD2(bool, islessequal); -__HIP_OVERLOAD2(bool, islessgreater); -__HIP_OVERLOAD1(bool, isnan); -//__HIP_OVERLOAD1(bool, isnormal) -__HIP_OVERLOAD2(bool, isunordered); -__DEF_FUN1(double, lgamma) -__DEF_FUN1(double, log) -__DEF_FUN1(double, log10) -__DEF_FUN1(double, log1p) -__DEF_FUN1(double, log2) -__DEF_FUN1(double, logb) -__DEF_FUNI(long long, llrint) -__DEF_FUNI(long long, llround) -__DEF_FUNI(long, lrint) -__DEF_FUNI(long, lround) -__DEF_FUN1(double, nearbyint); -__DEF_FUN2(double, nextafter); -__DEF_FUN2(double, pow); -__DEF_FUN2(double, remainder); -__DEF_FUN1(double, rint); -__DEF_FUN1(double, round); -__HIP_OVERLOAD1(bool, signbit) -__DEF_FUN1(double, sin) -__DEF_FUN1(double, sinh) -__DEF_FUN1(double, sqrt) -__DEF_FUN1(double, tan) -__DEF_FUN1(double, tanh) -__DEF_FUN1(double, tgamma) -__DEF_FUN1(double, trunc); - -// define cmath functions with a float and an integer argument. -#define __DEF_FLOAT_FUN2I(__func) \ - __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) { +template <class T> __DEVICE__ T min(T __arg1, T __arg2) { return (__arg1 < __arg2) ? __arg1 : __arg2; } -template <class T> __DEVICE__ inline T max(T __arg1, T __arg2) { +template <class T> __DEVICE__ T max(T __arg1, T __arg2) { return (__arg1 > __arg2) ? __arg1 : __arg2; } -__DEVICE__ inline int min(int __arg1, int __arg2) { +__DEVICE__ int min(int __arg1, int __arg2) { return (__arg1 < __arg2) ? __arg1 : __arg2; } -__DEVICE__ inline int max(int __arg1, int __arg2) { +__DEVICE__ int max(int __arg1, int __arg2) { return (__arg1 > __arg2) ? __arg1 : __arg2; } __DEVICE__ -inline float max(float __x, float __y) { return fmaxf(__x, __y); } +float max(float __x, float __y) { return fmaxf(__x, __y); } __DEVICE__ -inline double max(double __x, double __y) { return fmax(__x, __y); } +double max(double __x, double __y) { return fmax(__x, __y); } __DEVICE__ -inline float min(float __x, float __y) { return fminf(__x, __y); } +float min(float __x, float __y) { return fminf(__x, __y); } __DEVICE__ -inline double min(double __x, double __y) { return fmin(__x, __y); } - -__HIP_OVERLOAD2(double, max) -__HIP_OVERLOAD2(double, min) +double min(double __x, double __y) { return fmin(__x, __y); } __host__ inline static int min(int __arg1, int __arg2) { return std::min(__arg1, __arg2); @@ -1209,47 +1238,8 @@ __host__ inline static int min(int __arg1, int __arg2) { __host__ inline static int max(int __arg1, int __arg2) { return std::max(__arg1, __arg2); } - -#ifdef __cplusplus -__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); -} - -__DEVICE__ -inline float remquo(float __x, float __y, int *__quo) { - return remquof(__x, __y, __quo); -} - -template <typename __T1, typename __T2> -__DEVICE__ - typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized && - std::numeric_limits<__T2>::is_specialized, - double>::type - remquo(__T1 __x, __T2 __y, int *__quo) { - return remquo((double)__x, (double)__y, __quo); -} - -__DEVICE__ -inline float frexp(float __x, int *__nptr) { return frexpf(__x, __nptr); } - -__DEVICE__ -inline float modf(float __x, float *__iptr) { return modff(__x, __iptr); } - #endif -#pragma pop_macro("__DEF_FUN1") -#pragma pop_macro("__DEF_FUN2") -#pragma pop_macro("__DEF_FUNI") -#pragma pop_macro("__DEF_FLOAT_FUN2I") -#pragma pop_macro("__HIP_OVERLOAD1") -#pragma pop_macro("__HIP_OVERLOAD2") #pragma pop_macro("__DEVICE__") #pragma pop_macro("__RETURN_TYPE") diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index 007b5f12ffbd..81a16a265ae8 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -55,6 +55,7 @@ static inline __device__ void *free(void *__ptr) { #if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ #include <__clang_cuda_math_forward_declares.h> +#include <__clang_hip_cmath.h> #include <__clang_cuda_complex_builtins.h> #include <algorithm> _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
