[OpenMP] Use __OPENMP_NVPTX__ instead of _OPENMP in complex wrapper headers.

[fodinabor]

Summary

This is very similar to 7f1e6fcff942, just fixing a left-over.
With this, it should be possible to use both, -x cuda and -fopenmp in the same invocation,
enabling to use both OpenMP, targeting CPU, and CUDA, targeting the GPU.

[jdoerfert]

You need to define & undefine the macro around the includes of __clang_cuda_complex_builtins.h in clang/lib/Headers/openmp_wrappers/complex and clang/lib/Headers/openmp_wrappers/complex.h. (see also rG7f1e6fcff942) That should fix the tests. Assuming the tests pass, LGTM.

For release 11, can you file a bug and prepare a patch that applies to the relevant branch?

[fodinabor]

Add missing macro definitions.

[fodinabor]

Thanks, will land it later.

For the bug see: https://bugs.llvm.org/show_bug.cgi?id=48014
Do I have to create a new phabricator review, too?
I'm currently building the release/11.x branch with the patch cherry-picked (which worked flawlessly).

[jdoerfert]

In D90415#2363056, @fodinabor wrote:

Thanks, will land it later.

For the bug see: https://bugs.llvm.org/show_bug.cgi?id=48014
Do I have to create a new phabricator review, too?
I'm currently building the release/11.x branch with the patch cherry-picked (which worked flawlessly).

CC [e-mail] in the bug and ask for it to be merged. Attach the patch file to the bug. At least that is (one way) how it works I think.

clang/lib/Headers/__clang_cuda_complex_builtins.h

 #define __CLANG_CUDA_COMPLEX_BUILTINS
 
 // This header defines __muldc3, __mulsc3, __divdc3, and __divsc3.  These are
 // libgcc functions that clang assumes are available when compiling c99 complex
 // operations.  (These implementations come from libc++, and have been modified
 // to work with CUDA and OpenMP target offloading [in C and C++ mode].)
 
 #pragma push_macro("__DEVICE__")
-#ifdef _OPENMP
+#ifdef __OPENMP_NVPTX__
 #pragma omp declare target
 #define __DEVICE__ __attribute__((noinline, nothrow, cold, weak))
 #else
 #define __DEVICE__ __device__ inline
 #endif
 
 // To make the algorithms available for C and C++ in CUDA and OpenMP we select
 // different but equivalent function versions. TODO: For OpenMP we currently
 // select the native builtins as the overload support for templates is lacking.
-#if !defined(_OPENMP)
+#if !defined(__OPENMP_NVPTX__)
 #define _ISNANd std::isnan
 #define _ISNANf std::isnan
 #define _ISINFd std::isinf
 #define _ISINFf std::isinf
 #define _ISFINITEd std::isfinite
 #define _ISFINITEf std::isfinite
 #define _COPYSIGNd std::copysign
 #define _COPYSIGNf std::copysign

 #undef _SCALBNf
 #undef _ABSd
 #undef _ABSf
 #undef _LOGBd
 #undef _LOGBf
 #undef _fmaxd
 #undef _fmaxf
 
-#ifdef _OPENMP
+#ifdef __OPENMP_NVPTX__
 #pragma omp end declare target
 #endif
 
 #pragma pop_macro("__DEVICE__")
 
 #endif // __CLANG_CUDA_COMPLEX_BUILTINS

clang/lib/Headers/openmp_wrappers/complex

 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 // We require std::math functions in the complex builtins below.
 #include <cmath>
 
 #define __CUDA__
+#define __OPENMP_NVPTX__
 #include <__clang_cuda_complex_builtins.h>
+#undef __OPENMP_NVPTX__
 #endif
 
 // Grab the host header too.
 #include_next <complex>
 
 
 #ifdef __cplusplus
 

clang/lib/Headers/openmp_wrappers/complex.h

 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 // We require math functions in the complex builtins below.
 #include <math.h>
 
 #define __CUDA__
+#define __OPENMP_NVPTX__
 #include <__clang_cuda_complex_builtins.h>
+#undef __OPENMP_NVPTX__
 #endif
 
 // Grab the host header too.
 #include_next <complex.h>