[OpenMP] Provide math functions in OpenMP device code via OpenMP variants

[jdoerfert]

Summary

For OpenMP target regions to piggy back on the CUDA/AMDGPU/... implementation of math functions,
we include the appropriate definitions inside of an omp begin/end declare variant match(device={arch(nvptx)}) scope.
This way, the vendor specific math functions will become specialized versions of the system math functions.
When a system math function is called and specialized version is available the selection logic introduced in D75779
instead call the specialized version. In contrast to the code path we used so far, the system header is actually included.
This means functions without specialized versions are available and so are macro definitions.

This should address PR42061, PR42798, and PR42799.

[JonChesterfield]

That's less invasive than I feared. Nicely done.

It may worth keeping the openmp header wrapper to do architecture dispatch. Something like:

#ifndef __CLANG_OPENMP_MATH_DECLARES_H__
#define __CLANG_OPENMP_MATH_DECLARES_H__

#ifndef _OPENMP
#error "This file is for OpenMP compilation only."
#endif

#if defined(__AMDGCN__)
#pragma omp begin declare variant match(device = {arch(amdgcn)})
#include "equivalent_header.h"
#pragma omp end declare variant
#endif // __AMDGCN__

#if defined(__NVPTX__)
#define __CUDA__
#pragma omp begin declare variant match(device = {arch(nvptx)})

#if defined(__cplusplus)
#include <__clang_cuda_math_forward_declares.h>
#endif

/// Include declarations for libdevice functions.
#include <__clang_cuda_libdevice_declares.h>
/// Provide definitions for these functions.
#include <__clang_cuda_device_functions.h>

#pragma omp end declare variant
#undef __CUDA__
#endif // __NVPTX__


#endif // __CLANG_OPENMP_MATH_DECLARES_H__

clang/lib/Headers/cuda_wrappers/new

36 ↗

(On Diff #248864)

macros look off here - should it be #define DEVICE, or the following uses __DEVICE__?

[jdoerfert]

In D75788#1910743, @JonChesterfield wrote:

That's less invasive than I feared. Nicely done.

We need to run some more tests to make sure it works as expected but I hope we can completely piggy back on the underlying "language" support.

It may worth keeping the openmp header wrapper to do architecture dispatch. Something like:

We can do that or adjust the pipeline based on the target, either is fine with me.

clang/lib/Headers/cuda_wrappers/new

36 ↗

(On Diff #248864)

Yes.

Furthermore I think I want to introduce the effect of __device__ as an attribute, basically match(device={arch(nvptx)} on a single function. That would make the declare variant go away and allow us to piggy back on the __DEVICE__ directly.

[jdoerfert]

Add complex test case

[tra]

Couple of nits below. LGTM for CUDA headers otherwise.

clang/lib/Headers/__clang_cuda_math_forward_declares.h

42–43

Shouldn't float and double abs variants be preserved?

clang/lib/Headers/cuda_wrappers/new

107 ↗

(On Diff #248937)

You may want to push/pop __DEVICE__ macro here and in other headers where it's not done.

[jdoerfert]

Adjust to new scheme, tested locally with some math functions, seems to work

[jdoerfert]

clang/lib/AST/ASTContext.cpp

1621 ↗

(On Diff #250035)

These AuxTarget changes can be split off. I observed a segfault while testing stuff.

clang/lib/Headers/CMakeLists.txt

146

the math_declares above will provide all C math functions so we do not need math.h anymore here.

clang/lib/Headers/__clang_cuda_cmath.h

315

when this is included with OpenMP it is actually following the system <cmath> and these using declarations are neither needed nor helpful (afaict)

clang/lib/Headers/__clang_cuda_math_forward_declares.h

42–43

yes, thx! I make so many errors modifying these files, it is ridicules...

[jdoerfert]

Add and repair tests to show new features, e.g. math macro support

[jdoerfert]

Remove OpenMP from clang/lib/Headers/__clang_cuda_math_forward_declares.h

[jdoerfert]

Rewrite. Wrap math.h, time.h, and cmath. Preload only device functions.

Passes all 185 math c++11 tests from [0] which do not deal with long double.
[0] https://github.com/TApplencourt/OmpVal

[ye-luo]

Good work.
I verified that PR42798 and PR42799 are fixed by this.
Tests are completed on Ubuntu 18.04.
Clang now becomes usable for application developers.

There are still issues on RHEL that openmp_wrappers is not added before the system math searching path. Option -isystem openmp_wrappers can be used as a workaround.
This remaining issue should not be a blocker of accepting this patch. It can be dealt separately later without complicating the review of the current patch.

[ye-luo]

My RHEL issue was caused by a CPLUS_INCLUDE_PATH environment variable. So this is feature not a bug. After removing it, everything works smoothly for me.

[jdoerfert]

Cleanup and rebase

[sammccall]

clang/test/Headers/nvptx_device_math_macro.cpp

12

Hmm, this fails if the test is run from a directory containing "call".

; ModuleID = '/usr/local/google/home/sammccall/src/llvm-project/clang/test/Headers/nvptx_device_math_macro.cpp

I guess I can blame my parents :-)
Is CHECK-NOT: call double correct for both of these?

[jdoerfert]

clang/test/Headers/nvptx_device_math_macro.cpp

12

Oh my, ... sorry.

I guess call{{.*}}@ should do the trick. No globals are referenced in a call

[chapuni]

clang/lib/Headers/CMakeLists.txt

142

It doesn't contain <time.h>, intentional?

[jdoerfert]

clang/lib/Headers/CMakeLists.txt

142

We don't wrap that yet. Do you need it?

clang/lib/Headers/CMakeLists.txt

set(ppc_wrapper_files
   ppc_wrappers/xmmintrin.h
   ppc_wrappers/mm_malloc.h
   ppc_wrappers/emmintrin.h
   ppc_wrappers/pmmintrin.h
   ppc_wrappers/tmmintrin.h
   ppc_wrappers/smmintrin.h
 )
 
 set(openmp_wrapper_files
-  openmp_wrappers/math.h
-  openmp_wrappers/cmath
-  openmp_wrappers/__clang_openmp_math.h
   openmp_wrappers/__clang_openmp_math_declares.h
+  openmp_wrappers/cmath
   openmp_wrappers/new
 )
 
 set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
 set(out_files)
 set(generated_files)
 
 function(copy_header_to_output_dir src_dir file)
   set(src ${src_dir}/${file})
   set(dst ${output_dir}/${file})

clang/lib/Headers/__clang_cuda_cmath.h

 // 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); }
 __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); }
 __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);
 }
 __DEVICE__ int fpclassify(double __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
                               FP_ZERO, __x);
 }
-#endif
 __DEVICE__ float frexp(float __arg, int *__exp) {
   return ::frexpf(__arg, __exp);
 }
 
 // For inscrutable reasons, the CUDA headers define these functions for us on
 // Windows.
 #ifndef _MSC_VER
 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }

 
 template <typename __T>
 __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer,
                                            double>::type
 scalbn(__T __x, int __exp) {
   return std::scalbn((double)__x, __exp);
 }
 
+#ifndef _OPENMP
 // We need to define these overloads in exactly the namespace our standard
 // library uses (including the right inline namespace), otherwise they won't be
 // picked up by other functions in the standard library (e.g. functions in
 // <complex>).  Thus the ugliness below.
 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
 _LIBCPP_BEGIN_NAMESPACE_STD
 #else
 namespace std {

 using ::modff;
 using ::nearbyintf;
 using ::nextafterf;
 using ::powf;
 using ::remainderf;
 using ::remquof;
 using ::rintf;
 using ::roundf;
-// TODO: remove once variant is supported
-#ifndef _OPENMP
 using ::scalblnf;
-#endif
 using ::scalbnf;
 using ::sinf;
 using ::sinhf;
 using ::sqrtf;
 using ::tanf;
 using ::tanhf;
 using ::tgammaf;
 using ::truncf;
 
 #ifdef _LIBCPP_END_NAMESPACE_STD
 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
+#endif
 
 #undef __NOEXCEPT
 #undef __DEVICE__
 
 #endif

clang/lib/Headers/__clang_cuda_device_functions.h

 __DEVICE__ float atan2f(float __a, float __b) { return __nv_atan2f(__a, __b); }
 __DEVICE__ float atanf(float __a) { return __nv_atanf(__a); }
 __DEVICE__ double atanh(double __a) { return __nv_atanh(__a); }
 __DEVICE__ float atanhf(float __a) { return __nv_atanhf(__a); }
 __DEVICE__ double cbrt(double __a) { return __nv_cbrt(__a); }
 __DEVICE__ float cbrtf(float __a) { return __nv_cbrtf(__a); }
 __DEVICE__ double ceil(double __a) { return __nv_ceil(__a); }
 __DEVICE__ float ceilf(float __a) { return __nv_ceilf(__a); }
-#ifndef _OPENMP
 __DEVICE__ int clock() { return __nvvm_read_ptx_sreg_clock(); }
 __DEVICE__ long long clock64() { return __nvvm_read_ptx_sreg_clock64(); }
-#endif
 __DEVICE__ double copysign(double __a, double __b) {
   return __nv_copysign(__a, __b);
 }
 __DEVICE__ float copysignf(float __a, float __b) {
   return __nv_copysignf(__a, __b);
 }
 __DEVICE__ double cos(double __a) { return __nv_cos(__a); }
 __DEVICE__ float cosf(float __a) {
__DEVICE__ float rnormf(int __dim, const float *__t) {
   return __nv_rnormf(__dim, __t);
 }
 __DEVICE__ double round(double __a) { return __nv_round(__a); }
 __DEVICE__ float roundf(float __a) { return __nv_roundf(__a); }
 __DEVICE__ double rsqrt(double __a) { return __nv_rsqrt(__a); }
 __DEVICE__ float rsqrtf(float __a) { return __nv_rsqrtf(__a); }
 __DEVICE__ double scalbn(double __a, int __b) { return __nv_scalbn(__a, __b); }
 __DEVICE__ float scalbnf(float __a, int __b) { return __nv_scalbnf(__a, __b); }
-// TODO: remove once variant is supported
-#ifndef _OPENMP
 __DEVICE__ double scalbln(double __a, long __b) {
   if (__b > INT_MAX)
     return __a > 0 ? HUGE_VAL : -HUGE_VAL;
   if (__b < INT_MIN)
     return __a > 0 ? 0.0 : -0.0;
   return scalbn(__a, (int)__b);
 }
 __DEVICE__ float scalblnf(float __a, long __b) {
   if (__b > INT_MAX)
     return __a > 0 ? HUGE_VALF : -HUGE_VALF;
   if (__b < INT_MIN)
     return __a > 0 ? 0.f : -0.f;
   return scalbnf(__a, (int)__b);
 }
-#endif
 __DEVICE__ double sin(double __a) { return __nv_sin(__a); }
 __DEVICE__ void sincos(double __a, double *__s, double *__c) {
   return __nv_sincos(__a, __s, __c);
 }
 __DEVICE__ void sincosf(float __a, float *__s, float *__c) {
   return __FAST_OR_SLOW(__nv_fast_sincosf, __nv_sincosf)(__a, __s, __c);
 }
 __DEVICE__ void sincospi(double __a, double *__s, double *__c) {

clang/lib/Headers/__clang_cuda_math_forward_declares.h

 // This file forward-declares of some math functions we (or the CUDA headers)
 // will define later.  We need to do this, and do it before cmath is included,
 // because the standard library may have constexpr math functions.  In the
 // absence of a prior __device__ decl, those constexpr functions may become
 // implicitly host+device.  host+device functions can't be overloaded, so that
 // would preclude the use of our own __device__ overloads for these functions.
 
 #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);
 __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__ int abs(int);
 __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__ float asinh(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) __NOEXCEPT;
-__DEVICE__ float fabs(float) __NOEXCEPT;
+__DEVICE__ double fabs(double);
+__DEVICE__ float fabs(float);
 __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) __NOEXCEPT;
+__DEVICE__ long labs(long);
 __DEVICE__ double ldexp(double, int);
 __DEVICE__ float ldexp(float, int);
 __DEVICE__ double lgamma(double);
 __DEVICE__ float lgamma(float);
-__DEVICE__ long long llabs(long long) __NOEXCEPT;
+__DEVICE__ long long llabs(long long);
 __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);
 __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 *);

 _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

clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h

-/*===---- __clang_openmp_math.h - OpenMP target math support ---------------===
- *
- * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
- * See https://llvm.org/LICENSE.txt for license information.
- * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
- *
- *===-----------------------------------------------------------------------===
- */
-
-#if defined(__NVPTX__) && defined(_OPENMP)
-/// TODO:
-/// We are currently reusing the functionality of the Clang-CUDA code path
-/// as an alternative to the host declarations provided by math.h and cmath.
-/// This is suboptimal.
-///
-/// We should instead declare the device functions in a similar way, e.g.,
-/// through OpenMP 5.0 variants, and afterwards populate the module with the
-/// host declarations by unconditionally including the host math.h or cmath,
-/// respectively. This is actually what the Clang-CUDA code path does, using
-/// __device__ instead of variants to avoid redeclarations and get the desired
-/// overload resolution.
-
-#define __CUDA__
-
-#if defined(__cplusplus)
-  #include <__clang_cuda_cmath.h>
-#endif
-
-#undef __CUDA__
-
-/// Magic macro for stopping the math.h/cmath host header from being included.
-#define __CLANG_NO_HOST_MATH__
-
-#endif
-

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

 /*===---- __clang_openmp_math_declares.h - OpenMP math declares ------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_OPENMP_MATH_DECLARES_H__
 #define __CLANG_OPENMP_MATH_DECLARES_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
-#if defined(__NVPTX__) && defined(_OPENMP)
+/**
+ * A positive float constant expression. HUGE_VALF evaluates
+ * to +infinity. Used as an error value returned by the built-in
+ * math functions.
+ */
+#define HUGE_VALF (__builtin_huge_valf())
 
-#define __CUDA__
+/**
+ * A positive double constant expression. HUGE_VAL evaluates
+ * to +infinity. Used as an error value returned by the built-in
+ * math functions.
+ */
+#define HUGE_VAL (__builtin_huge_val())
 
 #if defined(__cplusplus)
-  #include <__clang_cuda_math_forward_declares.h>
+  #include <limits>
+  #include <climits>
+#else
+  #include <limits.h>
 #endif
 
+#pragma omp begin declare variant match(device={arch(nvptx64)})
+#define __CUDA__
+
 /// Include declarations for libdevice functions.
 #include <__clang_cuda_libdevice_declares.h>
+
 /// Provide definitions for these functions.
 #include <__clang_cuda_device_functions.h>
 
 #undef __CUDA__
+// TODO: Hack until we support an extension to the match clause that allows "or".
+#undef __CLANG_CUDA_LIBDEVICE_DECLARES_H__
+#undef __CLANG_CUDA_DEVICE_FUNCTIONS_H__
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch(nvptx)})
+#define __CUDA__
+
+/// Include declarations for libdevice functions.
+#include <__clang_cuda_libdevice_declares.h>
+
+/// Provide definitions for these functions.
+#include <__clang_cuda_device_functions.h>
+
+#undef __CUDA__
+#pragma omp end declare variant
+
+
+#undef HUGE_VAL
+#undef HUGE_VALF
 
-#endif
 #endif

clang/lib/Headers/openmp_wrappers/cmath

-/*===-------------- cmath - Alternative cmath header -----------------------===
+/*===---- __clang_openmp_math_declares.h - OpenMP math declares ------ c++ -===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
-#include <__clang_openmp_math.h>
+#ifndef __CLANG_OPENMP_CMATH_H__
+#define __CLANG_OPENMP_CMATH_H__
 
-#ifndef __CLANG_NO_HOST_MATH__
+#ifndef _OPENMP
+#error "This file is for OpenMP compilation only."
+#endif
+
+#include_next <cstdlib>
 #include_next <cmath>
-#else
-#undef __CLANG_NO_HOST_MATH__
+
+#pragma omp begin declare variant match(device={arch(nvptx64)})
+#define __CUDA__
+
+#include <__clang_cuda_cmath.h>
+
+// TODO: Hack until we support an extension to the match clause that allows "or".
+#undef __CLANG_CUDA_CMATH_H__
+
+#undef __CUDA__
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch(nvptx)})
+#define __CUDA__
+
+#include <__clang_cuda_cmath.h>
+
+#undef __CUDA__
+#pragma omp end declare variant
+
 #endif

clang/lib/Headers/openmp_wrappers/math.h

-/*===------------- math.h - Alternative math.h header ----------------------===
- *
- * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
- * See https://llvm.org/LICENSE.txt for license information.
- * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
- *
- *===-----------------------------------------------------------------------===
- */
-
-#include <__clang_openmp_math.h>
-
-#ifndef __CLANG_NO_HOST_MATH__
-#include_next <math.h>
-#else
-#undef __CLANG_NO_HOST_MATH__
-#endif
-

clang/test/Headers/Inputs/include/climits

+#pragma once
+
+#define INT_MIN  -2147483648
+#define INT_MAX   2147483647

clang/test/Headers/Inputs/include/cmath

 #pragma once
 
-double sqrt(double);
-double pow(double, double);
-double modf(double, double*);
+static double acos(double);
+static float acos(float);
+static double acosh(double);
+static float acosh(float);
+static double asin(double);
+static float asin(float);
+static double asinh(double);
+static float asinh(float);
+static double atan2(double, double);
+static float atan2(float, float);
+static double atan(double);
+static float atan(float);
+static double atanh(double);
+static float atanh(float);
+static double cbrt(double);
+static float cbrt(float);
+static double ceil(double);
+static float ceil(float);
+static double copysign(double, double);
+static float copysign(float, float);
+static double cos(double);
+static float cos(float);
+static double cosh(double);
+static float cosh(float);
+static double erfc(double);
+static float erfc(float);
+static double erf(double);
+static float erf(float);
+static double exp2(double);
+static float exp2(float);
+static double exp(double);
+static float exp(float);
+static double expm1(double);
+static float expm1(float);
+static double fdim(double, double);
+static float fdim(float, float);
+static double floor(double);
+static float floor(float);
+static double fma(double, double, double);
+static float fma(float, float, float);
+static double fmax(double, double);
+static float fmax(float, float);
+static double fmin(double, double);
+static float fmin(float, float);
+static double fmod(double, double);
+static float fmod(float, float);
+static int fpclassify(double);
+static int fpclassify(float);
+static double frexp(double, int *);
+static float frexp(float, int *);
+static double hypot(double, double);
+static float hypot(float, float);
+static int ilogb(double);
+static int ilogb(float);
+static bool isfinite(long double);
+static bool isfinite(double);
+static bool isfinite(float);
+static bool isgreater(double, double);
+static bool isgreaterequal(double, double);
+static bool isgreaterequal(float, float);
+static bool isgreater(float, float);
+static bool isinf(long double);
+static bool isinf(double);
+static bool isinf(float);
+static bool isless(double, double);
+static bool islessequal(double, double);
+static bool islessequal(float, float);
+static bool isless(float, float);
+static bool islessgreater(double, double);
+static bool islessgreater(float, float);
+static bool isnan(long double);
+static bool isnan(double);
+static bool isnan(float);
+static bool isnormal(double);
+static bool isnormal(float);
+static bool isunordered(double, double);
+static bool isunordered(float, float);
+static double ldexp(double, int);
+static float ldexp(float, int);
+static double lgamma(double);
+static float lgamma(float);
+static long long llrint(double);
+static long long llrint(float);
+static double log10(double);
+static float log10(float);
+static double log1p(double);
+static float log1p(float);
+static double log2(double);
+static float log2(float);
+static double logb(double);
+static float logb(float);
+static double log(double);
+static float log(float);
+static long lrint(double);
+static long lrint(float);
+static long lround(double);
+static long lround(float);
+static long long llround(float); // No llround(double).
+static double modf(double, double *);
+static float modf(float, float *);
+static double nan(const char *);
+static float nanf(const char *);
+static double nearbyint(double);
+static float nearbyint(float);
+static double nextafter(double, double);
+static float nextafter(float, float);
+static double pow(double, double);
+static double pow(double, int);
+static float pow(float, float);
+static float pow(float, int);
+static double remainder(double, double);
+static float remainder(float, float);
+static double remquo(double, double, int *);
+static float remquo(float, float, int *);
+static double rint(double);
+static float rint(float);
+static double round(double);
+static float round(float);
+static double scalbln(double, long);
+static float scalbln(float, long);
+static double scalbn(double, int);
+static float scalbn(float, int);
+static bool signbit(double);
+static bool signbit(float);
+static long double sin(long double);
+static double sin(double);
+static float sin(float);
+static double sinh(double);
+static float sinh(float);
+static double sqrt(double);
+static float sqrt(float);
+static double tan(double);
+static float tan(float);
+static double tanh(double);
+static float tanh(float);
+static double tgamma(double);
+static float tgamma(float);
+static double trunc(double);
+static float trunc(float);
+
+namespace std {
+
+using ::acos;
+using ::acosh;
+using ::asin;
+using ::asinh;
+using ::atan;
+using ::atan2;
+using ::atanh;
+using ::cbrt;
+using ::ceil;
+using ::copysign;
+using ::cos;
+using ::cosh;
+using ::erf;
+using ::erfc;
+using ::exp;
+using ::exp2;
+using ::expm1;
+using ::fdim;
+using ::floor;
+using ::fma;
+using ::fmax;
+using ::fmin;
+using ::fmod;
+using ::fpclassify;
+using ::frexp;
+using ::hypot;
+using ::ilogb;
+using ::isfinite;
+using ::isgreater;
+using ::isgreaterequal;
+using ::isinf;
+using ::isless;
+using ::islessequal;
+using ::islessgreater;
+using ::isnan;
+using ::isnormal;
+using ::isunordered;
+using ::ldexp;
+using ::lgamma;
+using ::llrint;
+using ::log;
+using ::log10;
+using ::log1p;
+using ::log2;
+using ::logb;
+using ::lrint;
+using ::lround;
+using ::llround;
+using ::modf;
+using ::nan;
+using ::nanf;
+using ::nearbyint;
+using ::nextafter;
+using ::pow;
+using ::remainder;
+using ::remquo;
+using ::rint;
+using ::round;
+using ::scalbln;
+using ::scalbn;
+using ::signbit;
+using ::sin;
+using ::sinh;
+using ::sqrt;
+using ::tan;
+using ::tanh;
+using ::tgamma;
+using ::trunc;
+
+} // namespace std
+
+#define FP_NAN 0
+#define FP_INFINITE 1
+#define FP_ZERO 2
+#define FP_SUBNORMAL 3
+#define FP_NORMAL 4

clang/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__)) ;
-extern float fabs (float __x) throw() __attribute__ ((__const__)) ;
+static int abs(int __x) throw() __attribute__((__const__)) { return __builtin_abs(__x); }
+static long int labs(long int __x) throw() __attribute__((__const__)) { return __builtin_labs(__x); }
+static float fabs(float __x) throw() __attribute__((__const__)) { return __builtin_fabs(__x); }
 #else
-extern int abs (int __x) __attribute__ ((__const__)) ;
-extern long int labs (long int __x) __attribute__ ((__const__)) ;
-extern float fabs (float __x) __attribute__ ((__const__)) ;
+static int abs(int __x) __attribute__((__const__)) { return __builtin_abs(__x); }
+static long int labs(long int __x) __attribute__((__const__)) { return __builtin_labs(__x); }
+static float fabs(float __x) __attribute__((__const__)) { return __builtin_fabs(__x); }
 #endif
 
-namespace std
-{
+namespace std {
 
 using ::abs;
 
-inline long
-abs(long __i) { return __builtin_labs(__i); }
-
-inline long long
-abs(long long __x) { return __builtin_llabs (__x); }
 }

clang/test/Headers/Inputs/include/math.h

 #pragma once
 
-double sqrt(double);
-double pow(double, double);
-double modf(double, double*);
+static double fabs(double __a);
+static double acos(double __a);
+static float acosf(float __a);
+static double acosh(double __a);
+static float acoshf(float __a);
+static double asin(double __a);
+static float asinf(float __a);
+static double asinh(double __a);
+static float asinhf(float __a);
+static double atan(double __a);
+static double atan2(double __a, double __b);
+static float atan2f(float __a, float __b);
+static float atanf(float __a);
+static double atanh(double __a);
+static float atanhf(float __a);
+static double cbrt(double __a);
+static float cbrtf(float __a);
+static double ceil(double __a);
+static float ceilf(float __a);
+static int clock();
+static long long clock64();
+static double copysign(double __a, double __b);
+static float copysignf(float __a, float __b);
+static double cos(double __a);
+static float cosf(float __a);
+static double cosh(double __a);
+static float coshf(float __a);
+static double cospi(double __a);
+static float cospif(float __a);
+static double cyl_bessel_i0(double __a);
+static float cyl_bessel_i0f(float __a);
+static double cyl_bessel_i1(double __a);
+static float cyl_bessel_i1f(float __a);
+static double erf(double __a);
+static double erfc(double __a);
+static float erfcf(float __a);
+static double erfcinv(double __a);
+static float erfcinvf(float __a);
+static double erfcx(double __a);
+static float erfcxf(float __a);
+static float erff(float __a);
+static double erfinv(double __a);
+static float erfinvf(float __a);
+static double exp(double __a);
+static double exp10(double __a);
+static float exp10f(float __a);
+static double exp2(double __a);
+static float exp2f(float __a);
+static float expf(float __a);
+static double expm1(double __a);
+static float expm1f(float __a);
+static float fabsf(float __a);
+static double fdim(double __a, double __b);
+static float fdimf(float __a, float __b);
+static double fdivide(double __a, double __b);
+static float fdividef(float __a, float __b);
+static double floor(double __f);
+static float floorf(float __f);
+static double fma(double __a, double __b, double __c);
+static float fmaf(float __a, float __b, float __c);
+static double fmax(double __a, double __b);
+static float fmaxf(float __a, float __b);
+static double fmin(double __a, double __b);
+static float fminf(float __a, float __b);
+static double fmod(double __a, double __b);
+static float fmodf(float __a, float __b);
+static double frexp(double __a, int *__b);
+static float frexpf(float __a, int *__b);
+static double hypot(double __a, double __b);
+static float hypotf(float __a, float __b);
+static int ilogb(double __a);
+static int ilogbf(float __a);
+static double j0(double __a);
+static float j0f(float __a);
+static double j1(double __a);
+static float j1f(float __a);
+static double jn(int __n, double __a);
+static float jnf(int __n, float __a);
+static double ldexp(double __a, int __b);
+static float ldexpf(float __a, int __b);
+static double lgamma(double __a);
+static float lgammaf(float __a);
+static long long llmax(long long __a, long long __b);
+static long long llmin(long long __a, long long __b);
+static long long llrint(double __a);
+static long long llrintf(float __a);
+static long long llround(double __a);
+static long long llroundf(float __a);
+static double log(double __a);
+static double log10(double __a);
+static float log10f(float __a);
+static double log1p(double __a);
+static float log1pf(float __a);
+static double log2(double __a);
+static float log2f(float __a);
+static double logb(double __a);
+static float logbf(float __a);
+static float logf(float __a);
+static long lrint(double __a);
+static long lrintf(float __a);
+static long lround(double __a);
+static long lroundf(float __a);
+static int max(int __a, int __b);
+static int min(int __a, int __b);
+static double modf(double __a, double *__b);
+static float modff(float __a, float *__b);
+static double nearbyint(double __a);
+static float nearbyintf(float __a);
+static double nextafter(double __a, double __b);
+static float nextafterf(float __a, float __b);
+static double norm(int __dim, const double *__t);
+static double norm3d(double __a, double __b, double __c);
+static float norm3df(float __a, float __b, float __c);
+static double norm4d(double __a, double __b, double __c, double __d);
+static float norm4df(float __a, float __b, float __c, float __d);
+static double normcdf(double __a);
+static float normcdff(float __a);
+static double normcdfinv(double __a);
+static float normcdfinvf(float __a);
+static float normf(int __dim, const float *__t);
+static double pow(double __a, double __b);
+static float powf(float __a, float __b);
+static double powi(double __a, int __b);
+static float powif(float __a, int __b);
+static double rcbrt(double __a);
+static float rcbrtf(float __a);
+static double remainder(double __a, double __b);
+static float remainderf(float __a, float __b);
+static double remquo(double __a, double __b, int *__c);
+static float remquof(float __a, float __b, int *__c);
+static double rhypot(double __a, double __b);
+static float rhypotf(float __a, float __b);
+static double rint(double __a);
+static float rintf(float __a);
+static double rnorm(int __a, const double *__b);
+static double rnorm3d(double __a, double __b, double __c);
+static float rnorm3df(float __a, float __b, float __c);
+static double rnorm4d(double __a, double __b, double __c, double __d);
+static float rnorm4df(float __a, float __b, float __c, float __d);
+static float rnormf(int __dim, const float *__t);
+static double round(double __a);
+static float roundf(float __a);
+static double rsqrt(double __a);
+static float rsqrtf(float __a);
+static double scalbn(double __a, int __b);
+static float scalbnf(float __a, int __b);
+static double scalbln(double __a, long __b);
+static float scalblnf(float __a, long __b);
+static double sin(double __a);
+static void sincos(double __a, double *__s, double *__c);
+static void sincosf(float __a, float *__s, float *__c);
+static void sincospi(double __a, double *__s, double *__c);
+static void sincospif(float __a, float *__s, float *__c);
+static float sinf(float __a);
+static double sinh(double __a);
+static float sinhf(float __a);
+static double sinpi(double __a);
+static float sinpif(float __a);
+static double sqrt(double __a);
+static float sqrtf(float __a);
+static double tan(double __a);
+static float tanf(float __a);
+static double tanh(double __a);
+static float tanhf(float __a);
+static double tgamma(double __a);
+static float tgammaf(float __a);
+static double trunc(double __a);
+static float truncf(float __a);
+static unsigned long long ullmax(unsigned long long __a,
+                                 unsigned long long __b);
+static unsigned long long ullmin(unsigned long long __a,
+                                 unsigned long long __b);
+static unsigned int umax(unsigned int __a, unsigned int __b);
+static unsigned int umin(unsigned int __a, unsigned int __b);
+static double y0(double __a);
+static float y0f(float __a);
+static double y1(double __a);
+static float y1f(float __a);
+static double yn(int __a, double __b);
+static float ynf(int __a, float __b);

clang/test/Headers/Inputs/include/stdlib.h

 #pragma once
 typedef __SIZE_TYPE__ size_t;
+
+static int abs(int __a);

clang/test/Headers/nvptx_device_cmath_functions.c

-// Test calling of device math functions.
-///==========================================================================///
-
-// REQUIRES: nvptx-registered-target
-
-// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// 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 -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 -o - | FileCheck -check-prefix CHECK-YES %s
-
-#include <cmath>
-
-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);
-    // CHECK-YES: call double @__nv_fabs(double
-    double l4 = fabs(a1);
-    // CHECK-YES: call i32 @__nv_abs(i32
-    double l5 = abs((int)a1);
-  }
-}

clang/test/Headers/nvptx_device_cmath_functions.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
-// 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 -o - | FileCheck -check-prefix CHECK-YES %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/Inputs/include -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 -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

clang/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
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -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 -internal-isystem %S/Inputs/include -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

clang/test/Headers/nvptx_device_math_complex.c

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+// expected-no-diagnostics
+
+// CHECK-DAG: call { float, float } @__divsc3(
+// CHECK-DAG: call { float, float } @__mulsc3(
+void test_scmplx(float _Complex a) {
+#pragma omp target
+  {
+    (void)(a * (a / a));
+  }
+}
+
+
+// CHECK-DAG: call { double, double } @__divdc3(
+// CHECK-DAG: call { double, double } @__muldc3(
+void test_dcmplx(double _Complex a) {
+#pragma omp target
+  {
+    (void)(a * (a / a));
+  }
+}

clang/test/Headers/nvptx_device_math_functions.c

 // Test calling of device math functions.
 ///==========================================================================///
 
 // REQUIRES: nvptx-registered-target
 
-// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// 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 -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 -o - | FileCheck -check-prefix CHECK-YES %s
+// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/Inputs/include -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 -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c++ -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/Inputs/include -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 -o - | FileCheck %s
 
+#include <stdlib.h>
 #include <math.h>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
-    // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK: call double @__nv_sqrt(double
     double l1 = sqrt(a1);
-    // CHECK-YES: call double @__nv_pow(double
+    // CHECK: call double @__nv_pow(double
     double l2 = pow(a1, a1);
-    // CHECK-YES: call double @__nv_modf(double
+    // CHECK: call double @__nv_modf(double
     double l3 = modf(a1 + 3.5, &a1);
-    // CHECK-YES: call double @__nv_fabs(double
+    // CHECK: call double @__nv_fabs(double
     double l4 = fabs(a1);
-    // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK: call i32 @__nv_abs(i32
     double l5 = abs((int)a1);
   }
 }

clang/test/Headers/nvptx_device_math_functions.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
-// 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 -o - | FileCheck -check-prefix CHECK-YES %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - | FileCheck -check-prefix CHECK-YES %s
 
 #include <cstdlib>
-#include <math.h>
+#include <cmath>
 
 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);
     // CHECK-YES: call double @__nv_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
     double l5 = abs((int)a1);
   }
 }

clang/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
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -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 -internal-isystem %S/Inputs/include -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>
+#include <cmath>
 
 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);
     // CHECK-YES: call double @__nv_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
     double l5 = abs((int)a1);
   }
 }

clang/test/Headers/nvptx_device_math_macro.cpp

+// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c++ -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - | FileCheck %s
+// expected-no-diagnostics
+
+#include <cmath>
+
+#pragma omp declare target
+int use_macro() {
+  double a(0);
+// CHECK:  %call = call i32 @_ZL10fpclassifyd(double %0)
+// CHECK:  %cmp = icmp ne i32 %call, 2
+  return (std::fpclassify(a) != FP_ZERO);
+}
+#pragma omp end declare target

clang/test/Headers/nvptx_device_math_sin.c

+// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - | FileCheck %s --check-prefix=SLOW
+// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math
+// RUN: %clang_cc1 -x c -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - -ffast-math | FileCheck %s --check-prefix=FAST
+// expected-no-diagnostics
+
+#include <math.h>
+
+double math(float f, double d, long double ld) {
+  double r = 0;
+// SLOW:  call float @__nv_sinf(float
+// FAST:  call fast float @__nv_fast_sinf(float
+  r += sinf(f);
+// SLOW:  call double @__nv_sin(double
+// FAST:  call fast double @__nv_sin(double
+  r += sin(d);
+  return r;
+}
+
+long double foo(float f, double d, long double ld) {
+  double r = ld;
+  r += math(f, d, ld);
+#pragma omp target map(r)
+  { r += math(f, d, ld); }
+  return r;
+}

clang/test/Headers/nvptx_device_math_sin.cpp

+// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -x c++ -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - | FileCheck %s --check-prefix=SLOW
+// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math
+// RUN: %clang_cc1 -x c++ -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 -o - -ffast-math | FileCheck %s --check-prefix=FAST
+// expected-no-diagnostics
+
+#include <cmath>
+
+double math(float f, double d, long double ld) {
+  double r = 0;
+// SLOW:  call float @__nv_sinf(float
+// FAST:  call fast float @__nv_fast_sinf(float
+  r += sin(f);
+// SLOW:  call double @__nv_sin(double
+// FAST:  call fast double @__nv_sin(double
+  r += sin(d);
+  return r;
+}
+
+long double foo(float f, double d, long double ld) {
+  double r = ld;
+  r += math(f, d, ld);
+#pragma omp target map(r)
+  { r += math(f, d, ld); }
+  return r;
+}