[clang] aa2b593 - [HIP] Restructure hip headers to add cmath
Aaron En Ye Shi via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 6 07:49:33 PDT 2020
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>
More information about the cfe-commits
mailing list