yaxunl created this revision.
yaxunl added a reviewer: tra.
yaxunl requested review of this revision.
Remove the dependence on standard C++ header
for overloaded math functions in HIP header
since standard C++ header is not available for hipRTC.
https://reviews.llvm.org/D100794
Files:
clang/lib/Headers/__clang_hip_cmath.h
clang/lib/Headers/__clang_hip_runtime_wrapper.h
clang/test/Headers/hip-header.hip
Index: clang/test/Headers/hip-header.hip
===================================================================
--- clang/test/Headers/hip-header.hip
+++ clang/test/Headers/hip-header.hip
@@ -21,7 +21,21 @@
}
// CHECK-LABEL: define{{.*}} i64 @_Z11test_size_tv
-// CHEC: ret i64 8
+// CHECK: ret i64 8
__device__ size_t test_size_t() {
return sizeof(size_t);
}
+
+// Check there is no ambiguity when calling overloaded math functions.
+
+// CHECK-LABEL: define{{.*}}@_Z10test_floorv
+// CHECK: call {{.*}}double @__ocml_floor_f64(double
+__device__ float test_floor() {
+ return floor(5);
+}
+
+// CHECK-LABEL: define{{.*}}@_Z8test_maxv
+// CHECK: call {{.*}}double @__ocml_fmax_f64(double {{.*}}, double
+__device__ float test_max() {
+ return max(5, 6.0);
+}
Index: clang/lib/Headers/__clang_hip_runtime_wrapper.h
===================================================================
--- clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -72,10 +72,12 @@
#include <__clang_hip_libdevice_declares.h>
#include <__clang_hip_math.h>
-#if !defined(__HIPCC_RTC__)
#if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
+#if !defined(__HIPCC_RTC__)
#include <__clang_cuda_math_forward_declares.h>
+#endif // __HIPCC_RTC__
#include <__clang_hip_cmath.h>
+#if !defined(__HIPCC_RTC__)
#include <__clang_cuda_complex_builtins.h>
#include <algorithm>
Index: clang/lib/Headers/__clang_hip_cmath.h
===================================================================
--- clang/lib/Headers/__clang_hip_cmath.h
+++ clang/lib/Headers/__clang_hip_cmath.h
@@ -22,7 +22,7 @@
#endif
#include <limits.h>
#include <stdint.h>
-#endif // __HIPCC_RTC__
+#endif // !defined(__HIPCC_RTC__)
#pragma push_macro("__DEVICE__")
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
@@ -36,6 +36,9 @@
__DEVICE__ float fma(float __x, float __y, float __z) {
return ::fmaf(__x, __y, __z);
}
+#if !defined(__HIPCC_RTC__)
+// The value returned by fpclassify is platform dependent, therefore it is not
+// supported by hipRTC.
__DEVICE__ int fpclassify(float __x) {
return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
FP_ZERO, __x);
@@ -44,6 +47,8 @@
return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
FP_ZERO, __x);
}
+#endif // !defined(__HIPCC_RTC__)
+
__DEVICE__ float frexp(float __arg, int *__exp) {
return ::frexpf(__arg, __exp);
}
@@ -209,11 +214,115 @@
template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
+namespace __hip {
+template <class _Tp> struct is_integral {
+ enum { value = 0 };
+};
+template <> struct is_integral<bool> {
+ enum { value = 1 };
+};
+template <> struct is_integral<char> {
+ enum { value = 1 };
+};
+template <> struct is_integral<signed char> {
+ enum { value = 1 };
+};
+template <> struct is_integral<unsigned char> {
+ enum { value = 1 };
+};
+template <> struct is_integral<wchar_t> {
+ enum { value = 1 };
+};
+template <> struct is_integral<short> {
+ enum { value = 1 };
+};
+template <> struct is_integral<unsigned short> {
+ enum { value = 1 };
+};
+template <> struct is_integral<int> {
+ enum { value = 1 };
+};
+template <> struct is_integral<unsigned int> {
+ enum { value = 1 };
+};
+template <> struct is_integral<long> {
+ enum { value = 1 };
+};
+template <> struct is_integral<unsigned long> {
+ enum { value = 1 };
+};
+template <> struct is_integral<long long> {
+ enum { value = 1 };
+};
+template <> struct is_integral<unsigned long long> {
+ enum { value = 1 };
+};
+
+template <class _Tp> struct is_arithmetic {
+ enum { value = 0 };
+};
+template <> struct is_arithmetic<bool> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<char> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<signed char> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<unsigned char> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<wchar_t> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<short> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<unsigned short> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<int> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<unsigned int> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<long> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<unsigned long> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<long long> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<unsigned long long> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<float> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<double> {
+ enum { value = 1 };
+};
+template <> struct is_arithmetic<_Float16> {
+ enum { value = 1 };
+};
+
+struct true_type {
+ static const __constant__ bool value = true;
+};
+struct false_type {
+ static const __constant__ bool value = false;
+};
+
+template <typename __T, typename __U> struct is_same : public false_type {};
+template <typename __T> struct is_same<__T, __T> : public true_type {};
+
// decltype is only available in C++11 and above.
#if __cplusplus >= 201103L
// __hip_promote
-namespace __hip {
-
template <class _Tp> struct __numeric_type {
static void __test(...);
static _Float16 __test(_Float16);
@@ -229,8 +338,8 @@
// No support for long double, use double instead.
static double __test(long double);
- typedef decltype(__test(std::declval<_Tp>())) type;
- static const bool value = !std::is_same<type, void>::value;
+ typedef decltype(__test(_Tp{})) type;
+ static const bool value = !is_same<type, void>::value;
};
template <> struct __numeric_type<void> { static const bool value = true; };
@@ -273,18 +382,17 @@
template <class _A1, class _A2 = void, class _A3 = void>
class __promote : public __promote_imp<_A1, _A2, _A3> {};
-
+#endif
} // namespace __hip
-#endif //__cplusplus >= 201103L
// __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) { \
+ __DEVICE__ \
+ typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type \
+ __fn(__T __x) { \
return ::__fn((double)__x); \
}
@@ -295,8 +403,7 @@
#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, \
+ __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, \
typename __hip::__promote<__T1, __T2>::type>::type \
__fn(__T1 __x, __T2 __y) { \
typedef typename __hip::__promote<__T1, __T2>::type __result_type; \
@@ -305,11 +412,10 @@
#else
#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) { \
+ __DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \
+ __hip::is_arithmetic<__T2>::value, \
+ __retty>::type \
+ __fn(__T1 __x, __T2 __y) { \
return __fn((double)__x, (double)__y); \
}
#endif
@@ -337,7 +443,9 @@
__HIP_OVERLOAD2(double, fmax)
__HIP_OVERLOAD2(double, fmin)
__HIP_OVERLOAD2(double, fmod)
+#if !defined(__HIPCC_RTC__)
__HIP_OVERLOAD1(int, fpclassify)
+#endif // !defined(__HIPCC_RTC__)
__HIP_OVERLOAD2(double, hypot)
__HIP_OVERLOAD1(int, ilogb)
__HIP_OVERLOAD1(bool, isfinite)
@@ -383,9 +491,8 @@
#if __cplusplus >= 201103L
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,
+ __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
+ __hip::is_arithmetic<__T3>::value,
typename __hip::__promote<__T1, __T2, __T3>::type>::type
fma(__T1 __x, __T2 __y, __T3 __z) {
typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
@@ -393,33 +500,32 @@
}
#else
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) {
+__DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
+ __hip::is_arithmetic<__T2>::value &&
+ __hip::is_arithmetic<__T3>::value,
+ double>::type
+fma(__T1 __x, __T2 __y, __T3 __z) {
return ::fma((double)__x, (double)__y, (double)__z);
}
#endif
template <typename __T>
__DEVICE__
- typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type
+ typename __hip_enable_if<__hip::is_integral<__T>::value, 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
+ typename __hip_enable_if<__hip::is_integral<__T>::value, 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
+ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
modf(__T __x, double *__exp) {
return ::modf((double)__x, __exp);
}
@@ -427,8 +533,8 @@
#if __cplusplus >= 201103L
template <typename __T1, typename __T2>
__DEVICE__
- typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized &&
- std::numeric_limits<__T2>::is_specialized,
+ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
+ __hip::is_arithmetic<__T2>::value,
typename __hip::__promote<__T1, __T2>::type>::type
remquo(__T1 __x, __T2 __y, int *__quo) {
typedef typename __hip::__promote<__T1, __T2>::type __result_type;
@@ -436,25 +542,24 @@
}
#else
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) {
+__DEVICE__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
+ __hip::is_arithmetic<__T2>::value,
+ double>::type
+remquo(__T1 __x, __T2 __y, int *__quo) {
return ::remquo((double)__x, (double)__y, __quo);
}
#endif
template <typename __T>
__DEVICE__
- typename __hip_enable_if<std::numeric_limits<__T>::is_integer, double>::type
+ typename __hip_enable_if<__hip::is_integral<__T>::value, 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
+ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
scalbn(__T __x, int __exp) {
return ::scalbn((double)__x, __exp);
}
@@ -469,6 +574,7 @@
#endif // defined(__cplusplus)
// Define these overloads inside the namespace our standard library uses.
+#if !defined(__HIPCC_RTC__)
#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
_LIBCPP_BEGIN_NAMESPACE_STD
#else
@@ -624,8 +730,10 @@
#endif
} // namespace std
#endif
+#endif // !defined(__HIPCC_RTC__)
// Define device-side math functions from <ymath.h> on MSVC.
+#if !defined(__HIPCC_RTC__)
#if defined(_MSC_VER)
// Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.
@@ -659,6 +767,7 @@
}
#endif // defined(__cplusplus)
#endif // defined(_MSC_VER)
+#endif // !defined(__HIPCC_RTC__)
#pragma pop_macro("__DEVICE__")
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
