[HIP] Support overloaded math functions for hipRTC

[yaxunl]

Summary

Remove the dependence on standard C++ header
for overloaded math functions in HIP header
since standard C++ header is not available for hipRTC.

[tra]

LGTM overall.

@jlebar: I could use your opinion here.

clang/lib/Headers/__clang_hip_cmath.h

343–344

@jlebar : Should we expect any observable surprises here?

clang/lib/Headers/__clang_hip_runtime_wrapper.h

76–86

I'd add an empty line separators between #if/#endif blocks.

Or, perhaps, it should be restructured as

#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>
#endif // __HIPCC_RTC__

[yaxunl]

clang/lib/Headers/__clang_hip_runtime_wrapper.h

76–86

will restructure it

[yaxunl]

revised by Artem's comments

[tra]

clang/lib/Headers/__clang_hip_cmath.h

343–344

_Tp{} may break if _Tp does not have a default constructor. At least it does so with C++20, while the old code did not: https://godbolt.org/z/PGEff5Mvv

[yaxunl]

clang/lib/Headers/__clang_hip_cmath.h

343–344

Right. It occurred to me that users may overload these math functions with arguments without default ctor, then we will hit the issue.

I will fix this.

[jlebar]

Other than the declval issue discussed above, this looks reasonable to me. I don't notice anything wrong with the standard library reimplementations here.

[yaxunl]

fix hip::numeric_type

[tra]

LGTM.

clang/lib/Headers/__clang_hip_cmath.h

345

Nit: }; // namespace __hip

[yaxunl]

clang/lib/Headers/__clang_hip_cmath.h

345

will do

[yaxunl]

clang/lib/Headers/__clang_hip_cmath.h

345

actually this is end of struct __numeric_type. end of namespace __hip is at line 390

[yaxunl]

clang/lib/Headers/__clang_hip_cmath.h

309

This changes the overloading behavior of math functions with _Float16 arguments since previously we use std::numeric_limits<T>::is_specialized to determine whether a type is an arithmetic, which is not specialized for _Float16.

To make this patch NFC for non-hipRTC, I would like to split this part to another patch. This facilitates debugging if issues arise

[yaxunl]

make it NFC for non-hipRTC

[tra]

clang/lib/Headers/__clang_hip_cmath.h

586–587

Nit: I'd add // <condition> here for consistency.

[ashi1]

clang/lib/Headers/__clang_hip_cmath.h

20

we may not need to #include limits and/or type_traits if we replaced them with __hip::is_integral and __hip::is_arithmetic.

587

I don't think this #endif is for !defined(__HIPCC_RTC__). It should be for _LIBCPP_BEGIN_NAMESPACE_STD.

[yaxunl]

clang/lib/Headers/__clang_hip_cmath.h

20

right. but we may need to move them to HIP headers first if HIP headers need them.

587

You are right. fixed.

clang/lib/Headers/__clang_hip_cmath.h

 
 #if !defined(__HIP__)
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
 #if !defined(__HIPCC_RTC__)
 #if defined(__cplusplus)
 #include <limits>
 #include <type_traits>
 #include <utility>
 #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))
 
 // 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);
 }
+#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);
 }
 __DEVICE__ int fpclassify(double __x) {
   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);
 }
 __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);
 }

 #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; };
 
+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 };
+};
+
+// ToDo: specializes is_arithmetic<_Float16>
+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 };
+};
+
+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);
   static float __test(float);
   static double __test(char);
   static double __test(int);
   static double __test(unsigned);
   static double __test(long);
   static double __test(unsigned long);
   static double __test(long long);
   static double __test(unsigned long long);
   static double __test(double);
   // 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; };
 
 template <class _A1, class _A2 = void, class _A3 = void,
           bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
               &&__numeric_type<_A3>::value>
 class __promote_imp {
 public:

 template <class _A1> class __promote_imp<_A1, void, void, true> {
 public:
   typedef typename __numeric_type<_A1>::type type;
   static const bool value = true;
 };
 
 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);                                                \
   }
 
 // __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).
 #if __cplusplus >= 201103L
 #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;         \
     return __fn((__result_type)__x, (__result_type)__y);                       \
   }
 #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
 
 __HIP_OVERLOAD1(double, acos)
 __HIP_OVERLOAD1(double, acosh)
 __HIP_OVERLOAD1(double, asin)
 __HIP_OVERLOAD1(double, asinh)

 __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)
+#if !defined(__HIPCC_RTC__)
 __HIP_OVERLOAD1(int, fpclassify)
+#endif // !defined(__HIPCC_RTC__)
 __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)

 // 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.
 #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;
   return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
 }
 #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,
+__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);
 }
 
 #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;
   return ::remquo((__result_type)__x, (__result_type)__y, __quo);
 }
 #else
 template <typename __T1, typename __T2>
-__DEVICE__
-    typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized &&
-                                 std::numeric_limits<__T2>::is_specialized,
+__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);
 }
 
 #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.
+#if !defined(__HIPCC_RTC__)
 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
 _LIBCPP_BEGIN_NAMESPACE_STD
 #else
 namespace std {
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_BEGIN_NAMESPACE_VERSION
 #endif
-#endif
+#endif // !defined(__HIPCC_RTC__)
 
 // Pull the new overloads we defined above into namespace std.
 // using ::abs; - This may be considered for C++.
 using ::acos;
 using ::acosh;
 using ::asin;
 using ::asinh;
 using ::atan;

 #ifdef _LIBCPP_END_NAMESPACE_STD
 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #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.
 // But, from VS2019, it's only included in `<complex>`. Need to include
 // `<ymath.h>` here to ensure C functions declared there won't be markded as
 // `__host__` and `__device__` through `<complex>` wrapper.
 #include <ymath.h>
 

 }
 __DEVICE__ __attribute__((overloadable)) float _FSinh(float x, float y) {
   return sinhf(x) * y;
 }
 #if defined(__cplusplus)
 }
 #endif // defined(__cplusplus)
 #endif // defined(_MSC_VER)
+#endif // !defined(__HIPCC_RTC__)
 
 #pragma pop_macro("__DEVICE__")
 
 #endif // __CLANG_HIP_CMATH_H__

clang/lib/Headers/__clang_hip_runtime_wrapper.h

static inline __device__ void *free(void *__ptr) {
   __builtin_trap();
   return nullptr;
 }
 #endif
 
 #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>
 #endif // __HIPCC_RTC__
 #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
 
 #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
 #if defined(__HIPCC_RTC__)
 #pragma pop_macro("NULL")
 #pragma pop_macro("uint32_t")
 #pragma pop_macro("uint64_t")
 #pragma pop_macro("CHAR_BIT")
 #pragma pop_macro("INT_MAX")
 #endif // __HIPCC_RTC__
 #endif // __HIP__
 #endif // __CLANG_HIP_RUNTIME_WRAPPER_H__

clang/test/Headers/hip-header.hip

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
 // RUN:   -internal-isystem %S/Inputs/include \
 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
 // RUN:   -D__HIPCC_RTC__ | FileCheck %s
 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
 // RUN:   -internal-isystem %S/Inputs/include \
 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
-// RUN:   -D__HIPCC_RTC__ -std=c++14 | FileCheck %s
-
+// RUN:   -D__HIPCC_RTC__ -std=c++14 | FileCheck -check-prefixes=CHECK,CXX14 %s
 
 // expected-no-diagnostics
 
+struct Number {
+  __device__ Number(float _x) : x(_x) {}
+  float x;
+};
+
+#if __cplusplus >= 201103L
+// Check __hip::__numeric_type can be used with a class without default ctor.
+__device__ void test_numeric_type() {
+  int x = __hip::__numeric_type<Number>::value;
+}
+
+// ToDo: Fix __clang_hip_cmake.h to specialize __hip::is_arithmetic<_Float16>
+// to resolve fma(_Float16, _Float16, int) to fma(double, double, double)
+// instead of fma(_Float16, _Float16, _Float16).
+
+// CXX14-LABEL: define{{.*}}@_Z8test_fma
+// CXX14: call {{.*}}@__ocml_fma_f16
+__device__ double test_fma(_Float16 h, int i) {
+  return fma(h, h, i);
+}
+
+#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);
+}