[OpenMP][bugfix] Add missing math functions variants for log and abs.

[gtbercea]

Summary

When including the random header in C++, some of the math functions it relies on are not present in the CUDA headers. We include this variants in this case.

[tra]

lib/Headers/__clang_cuda_cmath.h

55–56

Where do these functions come from? https://en.cppreference.com/w/cpp/numeric/math/fabs does not seem to show any abs functions with const args.

lib/Headers/__clang_cuda_device_functions.h

1604 ↗

(On Diff #199937)

Should the return type also be long double ?
In general long double is not supported by CUDA. I'd provide a declaration here, if it's absolutely needed, but no definition, so if someone attempts to actually use it, the application will fail to link.

Silently reducing precision by falling back to double is the wrong thing to do, IMO.

[gtbercea]

lib/Headers/__clang_cuda_cmath.h

55–56

What's happening is that when including the random header file, this header file uses abs with const arguments:

const double __n = _M_nd(__urng);
const double __y = -std::abs(__n) * __param._M_sm - 1;

And without there functions here the error I get is:

In file included from test.c:8:
In file included from /usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/random:52:
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/bits/random.tcc:1476:27: error: call to 'abs' is ambiguous
                    const double __y = -std::abs(__n) * __param._M_sm - 1;
                                        ^~~~~~~~
/usr/include/stdlib.h:770:12: note: candidate function
extern int abs (int __x) __THROW __attribute__ ((__const__)) __wur;
           ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:166:3: note: candidate function
  abs(long __i) { return __builtin_labs(__i); }
  ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:174:3: note: candidate function
  abs(long long __x) { return __builtin_llabs (__x); }
  ^
/usr/lib/gcc/ppc64le-redhat-linux/4.8.5/../../../../include/c++/4.8.5/cstdlib:179:3: note: candidate function
  abs(__int128 __x) { return __x >= 0 ? __x : -__x; }
  ^

[gtbercea]

lib/Headers/__clang_cuda_device_functions.h

1604 ↗

(On Diff #199937)

Agreed.

[efriedma]

lib/Headers/__clang_cuda_cmath.h

55–56

Overloading ignores the "const"; float abs(float) is the same thing. So probably clearer to declare it without the extra "const" modifiers, and then clean up the redundant definitions.

That said, I'm not sure how you managed to pick up the declarations from cstdlib, but not cmath. What files are getting included when this fails?

(Actually, in C++17, both cmath and cstdlib are supposed to declare all the float and integer overloads of abs, but that's only implemented in newer versions of libstdc++.)

[jdoerfert]

lib/Headers/__clang_cuda_cmath.h

55–56

That said, I'm not sure how you managed to pick up the declarations from cstdlib, but not cmath. What files are getting included when this fails?

Until we get support for OpenMP omp declare variant begin/end support in clang (X), we will only provide the CUDA declarations & definitions for math functions, not the host ones. It is one of these problems that arises when the single source language needs both, device and host definitions. (In CUDA the __device__ overload is the way out, in OpenMP it will be variant with an appropriate device context.)

(X) I'm right now working on integrating this into OpenMP 5.1 but we can implement it as soon as the design is decided on.

[jdoerfert]

This LGTM me, @tra or @efriedma any objections?

[gtbercea]

• update patch

[gtbercea]

@tra I eliminated the long double definition for log and only left the declaration.

[tra]

I'd add a comment with a brief explanation for the const variant and a TODO() to remove it.

Looks OK to me otherwise.

[JonChesterfield]

I'm not sure about this diff. I think it's breaking <iostream> and <algorithm>. Raised bug https://bugs.llvm.org/show_bug.cgi?id=42972

lib/Headers/__clang_cuda_cmath.h

 #endif
 
 #if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
 #endif
+// TODO: remove once variat is supported.
+#if defined(_OPENMP) && defined(__cplusplus)
+__DEVICE__ const float abs(const float __x) { return ::fabsf((float)__x); }
+__DEVICE__ const double abs(const double __x) { return ::fabs((double)__x); }
+#endif
 __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); }

lib/Headers/__clang_cuda_math_forward_declares.h

 #endif
 
 #if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long abs(long);
 __DEVICE__ long long abs(long long);
 __DEVICE__ double abs(double);
 __DEVICE__ float abs(float);
 #endif
+// While providing the CUDA declarations and definitions for math functions,
+// we may manually define additional functions.
+// TODO: Once variant is supported the additional functions will have
+// to be removed.
+#if defined(_OPENMP) && defined(__cplusplus)
+__DEVICE__ const double abs(const double);
+__DEVICE__ const float abs(const float);
+#endif
 __DEVICE__ int abs(int) __NOEXCEPT;
 __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 asinh(double);

 __DEVICE__ double log10(double);
 __DEVICE__ float log10(float);
 __DEVICE__ double log1p(double);
 __DEVICE__ float log1p(float);
 __DEVICE__ double log2(double);
 __DEVICE__ float log2(float);
 __DEVICE__ double logb(double);
 __DEVICE__ float logb(float);
+#if defined(_OPENMP) && defined(__cplusplus)
+__DEVICE__ long double log(long double);
+#endif
 __DEVICE__ double log(double);
 __DEVICE__ float log(float);
 __DEVICE__ long lrint(double);
 __DEVICE__ long lrint(float);
 __DEVICE__ long lround(double);
 __DEVICE__ long lround(float);
 __DEVICE__ long long llround(float); // No llround(double).
 __DEVICE__ double modf(double, double *);