[OpenMP][bugfix] Fix issues with C++ 17 compilation when handling math functions

[gtbercea]

Summary

In OpenMP device offloading we must ensure that unde C++ 17, the inclusion of cstdlib will works correctly.

[Hahnfeld]

lib/Headers/__clang_cuda_device_functions.h

1486

If I recall correctly, __cplusplus is not defined in C mode, so both GCC and Clang will issue a warning with -Wundef.

Maybe this can be solved with something similar to:

#ifdef __cplusplus
#define cpp_version __cplusplus
#else
#define cpp_version 0
#endif

(to avoid repetition of #if defined(__cplusplus) && __cplusplus >= 201703L)

[jdoerfert]

Missing defined(__cplusplus) to avoid warnings and we have to resolve the two new includes. Otherwise, this looks good. So, if you "fix" both go ahead and commit, if not, lets discuss.

lib/Headers/__clang_cuda_device_functions.h

1486

I dislike defining the version and advice to repeating #if defined(__cplusplus) && __cplusplus >= 201703L)

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

24 ↗

(On Diff #199613)

I ask this question again and again, do we need them now? Why?

[tra]

lib/Headers/__clang_cuda_device_functions.h

1486

I agree with @jdoerfert.
If that's needed too often, we can move c++-only functions under one global #if defined(__cplusplus).

1584–1591

This is typically dealt with by defining and using a macro which would expand to noexcept or to nothing.
E.g. something like what libcxx does.

This should reduce the number of needed #if.

[gtbercea]

• Refactor.

[gtbercea]

lib/Headers/__clang_cuda_device_functions.h

1486

@tra is this closer to what you want?

[tra]

LGTM.

lib/Headers/__clang_cuda_cmath.h

42

I think the change is useful for CUDA, too, but I'm OK keeping it OpenMP-specific for now. I can enable it for CUDA later.

[jdoerfert]

lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

24 ↗

(On Diff #199613)

Why did you add these includes?

[gtbercea]

• Remove included headers.

lib/Headers/__clang_cuda_cmath.h

 // std covers all of the known knowns.
 
 #ifdef _OPENMP
 #define __DEVICE__ static __attribute__((always_inline))
 #else
 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
 #endif
 
+// For C++ 17 we need to include noexcept attribute to be compatible
+// with the header-defined version. This may be removed once
+// variant is supported.
+#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
+#define __NOEXCEPT noexcept
+#else
+#define __NOEXCEPT
+#endif
+
 #if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 #endif
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
 __DEVICE__ float exp(float __x) { return ::expf(__x); }
-__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
+__DEVICE__ float fabs(float __x) __NOEXCEPT { return ::fabsf(__x); }
 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
 // TODO: remove when variant is supported
 #ifndef _OPENMP
 __DEVICE__ int fpclassify(float __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
                               FP_ZERO, __x);
 }

 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
 
+#undef __NOEXCEPT
 #undef __DEVICE__
 
 #endif

lib/Headers/__clang_cuda_device_functions.h

 // -ffast-math or -fcuda-approx-transcendentals are in effect.
 #pragma push_macro("__FAST_OR_SLOW")
 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
 #define __FAST_OR_SLOW(fast, slow) fast
 #else
 #define __FAST_OR_SLOW(fast, slow) slow
 #endif
 
+// For C++ 17 we need to include noexcept attribute to be compatible
+// with the header-defined version. This may be removed once
+// variant is supported.
+#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
+#define __NOEXCEPT noexcept
+#else
+#define __NOEXCEPT
+#endif
+
 __DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); }
 __DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); }
 __DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); }
 __DEVICE__ unsigned int __brev(unsigned int __a) { return __nv_brev(__a); }
 __DEVICE__ unsigned long long __brevll(unsigned long long __a) {
   return __nv_brevll(__a);
 }
 #if defined(__cplusplus)

 __DEVICE__ unsigned int __vsubus4(unsigned int __a, unsigned int __b) {
   unsigned int r;
   asm("vsub4.u32.u32.u32.sat %0,%1,%2,%3;"
       : "=r"(r)
       : "r"(__a), "r"(__b), "r"(0));
   return r;
 }
 #endif // CUDA_VERSION >= 9020
-__DEVICE__ int abs(int __a) { return __nv_abs(__a); }
+__DEVICE__ int abs(int __a) __NOEXCEPT { return __nv_abs(__a); }
+__DEVICE__ double fabs(double __a) __NOEXCEPT { return __nv_fabs(__a); }
 __DEVICE__ double acos(double __a) { return __nv_acos(__a); }
 __DEVICE__ float acosf(float __a) { return __nv_acosf(__a); }
 __DEVICE__ double acosh(double __a) { return __nv_acosh(__a); }
 __DEVICE__ float acoshf(float __a) { return __nv_acoshf(__a); }
 __DEVICE__ double asin(double __a) { return __nv_asin(__a); }
 __DEVICE__ float asinf(float __a) { return __nv_asinf(__a); }
 __DEVICE__ double asinh(double __a) { return __nv_asinh(__a); }
 __DEVICE__ float asinhf(float __a) { return __nv_asinhf(__a); }

 __DEVICE__ double exp(double __a) { return __nv_exp(__a); }
 __DEVICE__ double exp10(double __a) { return __nv_exp10(__a); }
 __DEVICE__ float exp10f(float __a) { return __nv_exp10f(__a); }
 __DEVICE__ double exp2(double __a) { return __nv_exp2(__a); }
 __DEVICE__ float exp2f(float __a) { return __nv_exp2f(__a); }
 __DEVICE__ float expf(float __a) { return __nv_expf(__a); }
 __DEVICE__ double expm1(double __a) { return __nv_expm1(__a); }
 __DEVICE__ float expm1f(float __a) { return __nv_expm1f(__a); }
-__DEVICE__ double fabs(double __a) { return __nv_fabs(__a); }
 __DEVICE__ float fabsf(float __a) { return __nv_fabsf(__a); }
 __DEVICE__ double fdim(double __a, double __b) { return __nv_fdim(__a, __b); }
 __DEVICE__ float fdimf(float __a, float __b) { return __nv_fdimf(__a, __b); }
 __DEVICE__ double fdivide(double __a, double __b) { return __a / __b; }
 __DEVICE__ float fdividef(float __a, float __b) {
 #if __FAST_MATH__ && !__CUDA_PREC_DIV
   return __nv_fast_fdividef(__a, __b);
 #else

 __DEVICE__ int ilogbf(float __a) { return __nv_ilogbf(__a); }
 __DEVICE__ double j0(double __a) { return __nv_j0(__a); }
 __DEVICE__ float j0f(float __a) { return __nv_j0f(__a); }
 __DEVICE__ double j1(double __a) { return __nv_j1(__a); }
 __DEVICE__ float j1f(float __a) { return __nv_j1f(__a); }
 __DEVICE__ double jn(int __n, double __a) { return __nv_jn(__n, __a); }
 __DEVICE__ float jnf(int __n, float __a) { return __nv_jnf(__n, __a); }
 #if defined(__LP64__) || defined(_WIN64)
-__DEVICE__ long labs(long __a) { return __nv_llabs(__a); };
+__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_llabs(__a); };
 #else
-__DEVICE__ long labs(long __a) { return __nv_abs(__a); };
+__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_abs(__a); };
 #endif
 __DEVICE__ double ldexp(double __a, int __b) { return __nv_ldexp(__a, __b); }
 __DEVICE__ float ldexpf(float __a, int __b) { return __nv_ldexpf(__a, __b); }
 __DEVICE__ double lgamma(double __a) { return __nv_lgamma(__a); }
 __DEVICE__ float lgammaf(float __a) { return __nv_lgammaf(__a); }
-__DEVICE__ long long llabs(long long __a) { return __nv_llabs(__a); }
+__DEVICE__ long long llabs(long long __a) __NOEXCEPT { return __nv_llabs(__a); }
 __DEVICE__ long long llmax(long long __a, long long __b) {
   return __nv_llmax(__a, __b);
 }
 __DEVICE__ long long llmin(long long __a, long long __b) {
   return __nv_llmin(__a, __b);
 }
 __DEVICE__ long long llrint(double __a) { return __nv_llrint(__a); }
 __DEVICE__ long long llrintf(float __a) { return __nv_llrintf(__a); }

 }
 __DEVICE__ double y0(double __a) { return __nv_y0(__a); }
 __DEVICE__ float y0f(float __a) { return __nv_y0f(__a); }
 __DEVICE__ double y1(double __a) { return __nv_y1(__a); }
 __DEVICE__ float y1f(float __a) { return __nv_y1f(__a); }
 __DEVICE__ double yn(int __a, double __b) { return __nv_yn(__a, __b); }
 __DEVICE__ float ynf(int __a, float __b) { return __nv_ynf(__a, __b); }
 
+#undef __NOEXCEPT
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__FAST_OR_SLOW")
 #endif // __CLANG_CUDA_DEVICE_FUNCTIONS_H__

lib/Headers/__clang_cuda_math_forward_declares.h

 #pragma push_macro("__DEVICE__")
 #ifdef _OPENMP
 #define __DEVICE__ static __inline__ __attribute__((always_inline))
 #else
 #define __DEVICE__                                                             \
   static __inline__ __attribute__((always_inline)) __attribute__((device))
 #endif
 
+// For C++ 17 we need to include noexcept attribute to be compatible
+// with the header-defined version. This may be removed once
+// variant is supported.
+#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
+#define __NOEXCEPT noexcept
+#else
+#define __NOEXCEPT
+#endif
+
 #if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long abs(long);
 __DEVICE__ long long abs(long long);
 #endif
-__DEVICE__ int abs(int);
+__DEVICE__ int abs(int) __NOEXCEPT;
 __DEVICE__ double abs(double);
 __DEVICE__ float abs(float);
 __DEVICE__ double acos(double);
 __DEVICE__ float acos(float);
 __DEVICE__ double acosh(double);
 __DEVICE__ float acosh(float);
 __DEVICE__ double asin(double);
 __DEVICE__ float asin(float);

 __DEVICE__ double erf(double);
 __DEVICE__ float erf(float);
 __DEVICE__ double exp2(double);
 __DEVICE__ float exp2(float);
 __DEVICE__ double exp(double);
 __DEVICE__ float exp(float);
 __DEVICE__ double expm1(double);
 __DEVICE__ float expm1(float);
-__DEVICE__ double fabs(double);
-__DEVICE__ float fabs(float);
+__DEVICE__ double fabs(double) __NOEXCEPT;
+__DEVICE__ float fabs(float) __NOEXCEPT;
 __DEVICE__ double fdim(double, double);
 __DEVICE__ float fdim(float, float);
 __DEVICE__ double floor(double);
 __DEVICE__ float floor(float);
 __DEVICE__ double fma(double, double, double);
 __DEVICE__ float fma(float, float, float);
 __DEVICE__ double fmax(double, double);
 __DEVICE__ float fmax(float, float);

 __DEVICE__ bool isnan(long double);
 #endif
 __DEVICE__ bool isnan(double);
 __DEVICE__ bool isnan(float);
 __DEVICE__ bool isnormal(double);
 __DEVICE__ bool isnormal(float);
 __DEVICE__ bool isunordered(double, double);
 __DEVICE__ bool isunordered(float, float);
-__DEVICE__ long labs(long);
+__DEVICE__ long labs(long) __NOEXCEPT;
 __DEVICE__ double ldexp(double, int);
 __DEVICE__ float ldexp(float, int);
 __DEVICE__ double lgamma(double);
 __DEVICE__ float lgamma(float);
-__DEVICE__ long long llabs(long long);
+__DEVICE__ long long llabs(long long) __NOEXCEPT;
 __DEVICE__ long long llrint(double);
 __DEVICE__ long long llrint(float);
 __DEVICE__ double log10(double);
 __DEVICE__ float log10(float);
 __DEVICE__ double log1p(double);
 __DEVICE__ float log1p(float);
 __DEVICE__ double log2(double);
 __DEVICE__ float log2(float);

 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
 
+#undef __NOEXCEPT
 #pragma pop_macro("__DEVICE__")
 
 #endif

test/Headers/Inputs/include/cstdlib

 #pragma once
 
+#if __cplusplus >= 201703L
+extern int abs (int __x) throw()  __attribute__ ((__const__)) ;
+extern long int labs (long int __x) throw() __attribute__ ((__const__)) ;
+#else
 extern int abs (int __x) __attribute__ ((__const__)) ;
 extern long int labs (long int __x) __attribute__ ((__const__)) ;
+#endif
 
 namespace std
 {
 
 using ::abs;
 
 inline long
 abs(long __i) { return __builtin_labs(__i); }
 
 inline long long
 abs(long long __x) { return __builtin_llabs (__x); }
 }

test/Headers/nvptx_device_cmath_functions_cxx17.cpp

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -std=c++17
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -internal-isystem %S/Inputs/include -include stdlib.h -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -std=c++17 -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cmath>
+#include <cstdlib>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}

test/Headers/nvptx_device_math_functions_cxx17.cpp

+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -std=c++17
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -internal-isystem %S/Inputs/include -include stdlib.h -include limits -include cstdlib -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -std=c++17 -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cstdlib>
+#include <math.h>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}