diff --git a/clang/lib/Headers/__clang_cuda_cmath.h b/clang/lib/Headers/__clang_cuda_cmath.h --- a/clang/lib/Headers/__clang_cuda_cmath.h +++ b/clang/lib/Headers/__clang_cuda_cmath.h @@ -12,7 +12,7 @@ #error "This file is for CUDA compilation only." #endif -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ #include #endif @@ -32,7 +32,7 @@ // implementation. Declaring in the global namespace and pulling into namespace // std covers all of the known knowns. -#ifdef _OPENMP +#ifdef __OPENMP_NVPTX__ #define __DEVICE__ static constexpr __attribute__((always_inline, nothrow)) #else #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) @@ -69,7 +69,7 @@ // Windows. For OpenMP we omit these as some old system headers have // non-conforming `isinf(float)` and `isnan(float)` implementations that return // an `int`. The system versions of these functions should be fine anyway. -#if !defined(_MSC_VER) && !defined(_OPENMP) +#if !defined(_MSC_VER) && !defined(__OPENMP_NVPTX__) __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); } __DEVICE__ bool isinf(double __x) { return ::__isinf(__x); } __DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); } @@ -146,7 +146,7 @@ // libdevice doesn't provide an implementation, and we don't want to be in the // business of implementing tricky libm functions in this header. -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ // Now we've defined everything we promised we'd define in // __clang_cuda_math_forward_declares.h. We need to do two additional things to @@ -463,7 +463,7 @@ } // namespace std #endif -#endif // _OPENMP +#endif // __OPENMP_NVPTX__ #undef __DEVICE__ diff --git a/clang/lib/Headers/__clang_cuda_device_functions.h b/clang/lib/Headers/__clang_cuda_device_functions.h --- a/clang/lib/Headers/__clang_cuda_device_functions.h +++ b/clang/lib/Headers/__clang_cuda_device_functions.h @@ -10,7 +10,7 @@ #ifndef __CLANG_CUDA_DEVICE_FUNCTIONS_H__ #define __CLANG_CUDA_DEVICE_FUNCTIONS_H__ -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ #if CUDA_VERSION < 9000 #error This file is intended to be used with CUDA-9+ only. #endif @@ -20,7 +20,7 @@ // we implement in this file. We need static in order to avoid emitting unused // functions and __forceinline__ helps inlining these wrappers at -O1. #pragma push_macro("__DEVICE__") -#ifdef _OPENMP +#ifdef __OPENMP_NVPTX__ #define __DEVICE__ static __attribute__((always_inline, nothrow)) #else #define __DEVICE__ static __device__ __forceinline__ @@ -1466,14 +1466,14 @@ // For OpenMP we require the user to include as we need to know what // clock_t is on the system. -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ __DEVICE__ /* clock_t= */ int clock() { return __nvvm_read_ptx_sreg_clock(); } #endif __DEVICE__ long long clock64() { return __nvvm_read_ptx_sreg_clock64(); } // These functions shouldn't be declared when including this header // for math function resolution purposes. -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ __DEVICE__ void *memcpy(void *__a, const void *__b, size_t __c) { return __builtin_memcpy(__a, __b, __c); } diff --git a/clang/lib/Headers/__clang_cuda_libdevice_declares.h b/clang/lib/Headers/__clang_cuda_libdevice_declares.h --- a/clang/lib/Headers/__clang_cuda_libdevice_declares.h +++ b/clang/lib/Headers/__clang_cuda_libdevice_declares.h @@ -14,7 +14,7 @@ extern "C" { #endif -#if defined(_OPENMP) +#if defined(__OPENMP_NVPTX__) #define __DEVICE__ #elif defined(__CUDA__) #define __DEVICE__ __device__ diff --git a/clang/lib/Headers/__clang_cuda_math.h b/clang/lib/Headers/__clang_cuda_math.h --- a/clang/lib/Headers/__clang_cuda_math.h +++ b/clang/lib/Headers/__clang_cuda_math.h @@ -12,7 +12,7 @@ #error "This file is for CUDA compilation only." #endif -#ifndef _OPENMP +#ifndef __OPENMP_NVPTX__ #if CUDA_VERSION < 9000 #error This file is intended to be used with CUDA-9+ only. #endif @@ -22,7 +22,7 @@ // we implement in this file. We need static in order to avoid emitting unused // functions and __forceinline__ helps inlining these wrappers at -O1. #pragma push_macro("__DEVICE__") -#ifdef _OPENMP +#ifdef __OPENMP_NVPTX__ #if defined(__cplusplus) #define __DEVICE__ static constexpr __attribute__((always_inline, nothrow)) #else @@ -36,7 +36,7 @@ // because the OpenMP overlay requires constexpr functions here but prior to // c++14 void return functions could not be constexpr. #pragma push_macro("__DEVICE_VOID__") -#ifdef _OPENMP && defined(__cplusplus) && __cplusplus < 201402L +#ifdef __OPENMP_NVPTX__ && defined(__cplusplus) && __cplusplus < 201402L #define __DEVICE_VOID__ static __attribute__((always_inline, nothrow)) #else #define __DEVICE_VOID__ __DEVICE__ diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h --- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h +++ b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h @@ -22,11 +22,15 @@ #endif #define __CUDA__ +#define __OPENMP_NVPTX__ + /// Include declarations for libdevice functions. #include <__clang_cuda_libdevice_declares.h> /// Provide definitions for these functions. #include <__clang_cuda_device_functions.h> + +#undef __OPENMP_NVPTX__ #undef __CUDA__ #ifdef __cplusplus diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath --- a/clang/lib/Headers/openmp_wrappers/cmath +++ b/clang/lib/Headers/openmp_wrappers/cmath @@ -28,7 +28,9 @@ device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) #define __CUDA__ +#define __OPENMP_NVPTX__ #include <__clang_cuda_cmath.h> +#undef __OPENMP_NVPTX__ #undef __CUDA__ // Overloads not provided by the CUDA wrappers but by the CUDA system headers. diff --git a/clang/lib/Headers/openmp_wrappers/math.h b/clang/lib/Headers/openmp_wrappers/math.h --- a/clang/lib/Headers/openmp_wrappers/math.h +++ b/clang/lib/Headers/openmp_wrappers/math.h @@ -41,7 +41,9 @@ device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) #define __CUDA__ +#define __OPENMP_NVPTX__ #include <__clang_cuda_math.h> +#undef __OPENMP_NVPTX__ #undef __CUDA__ #pragma omp end declare variant