yaxunl updated this revision to Diff 339009.
yaxunl marked an inline comment as done.
yaxunl added a comment.
fix __hip::__numeric_type
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D100794/new/
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
@@ -12,16 +12,41 @@
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ -std=c++14 | FileCheck %s
-
// expected-no-diagnostics
+struct Number {
+ __device__ Number(float _x) : x(_x) {}
+ float x;
+};
+
+// Check __hip::__numeric_type can be used with a class without default ctor.
+#if __cplusplus >= 201103L
+__device__ void test_numeric_type() {
+ int x = __hip::__numeric_type<Number>::value;
+}
+#endif
+
// CHECK-LABEL: amdgpu_kernel void @_Z4kernPff
__global__ void kern(float *x, float y) {
*x = sin(y);
}
// 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,12 +72,13 @@
#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_hip_cmath.h>
+#else
#include <__clang_cuda_math_forward_declares.h>
#include <__clang_hip_cmath.h>
#include <__clang_cuda_complex_builtins.h>
-
#include <algorithm>
#include <complex>
#include <new>
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,119 @@
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 {};
+
+template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
+
+template <typename __T> typename add_rvalue_reference<__T>::type declval();
+
// 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 +342,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(declval<_Tp>())) type;
+ static const bool value = !is_same<type, void>::value;
};
template <> struct __numeric_type<void> { static const bool value = true; };
@@ -273,18 +386,17 @@
template <class _A1, class _A2 = void, class _A3 = void>
class __promote : public __promote_imp<_A1, _A2, _A3> {};
-
-} // namespace __hip
#endif //__cplusplus >= 201103L
+} // namespace __hip
// __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 +407,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 +416,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 +447,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 +495,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 +504,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 +537,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 +546,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 +578,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 +734,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 +771,7 @@
}
#endif // defined(__cplusplus)
#endif // defined(_MSC_VER)
+#endif // !defined(__HIPCC_RTC__)
#pragma pop_macro("__DEVICE__")
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
