[OpenMP] Replace math headers with OpenMP wrapper calls

[jhuber6]

Summary

This patch replaces the math functions when called on the device with an
__omp_ variant that we can use. This will later be turned back into
the regular math call so we can avoid the problems with including math.h
on the device. This patch introduces the -fopenmp-device-libm flag
which enables the new math wrappers by defining MATH_WRAPPERS in the
header.

[tianshilei1992]

LGTM. Support the header file is actually generated automatically. :-)

[JonChesterfield]

We've got quite a lot of debt in this area and seem at risk of taking on more. Ideally the cuda and hip and openmp headers would be closer to a single header containing:

double acosh(double);
INSTANTIATE(1, acosh, cuda_acosh, amdgpu_acosh, intel_acosh);

clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h

43

This is not totally good. Can we do something with macros or code generators instead, even in the first instance?

Also wonder if what we actually want to do is map these things onto __gpu_acosh or __llvm_gpu_acosh or similar, there's nothing openmp specific about this and it would be really nice to not end up spinning another copy of this for some other gpu language.

[jdoerfert]

We want one copy of this, yes. s/omp/llvm_gpu/, sure, though gpu might also not the best name. Generators, macros, etc. I'm not sold. Even if we want to do this, let's not conflate that.

clang/include/clang/Driver/Options.td

defm openmp_target_new_runtime: BoolFOption<"openmp-target-new-runtime",
   NegFlag<SetFalse>>;
 defm openmp_optimistic_collapse : BoolFOption<"openmp-optimistic-collapse",
   LangOpts<"OpenMPOptimisticCollapse">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option]>, NegFlag<SetFalse>, BothFlags<[NoArgumentUnused, HelpHidden]>>;
 def static_openmp: Flag<["-"], "static-openmp">,
   HelpText<"Use the static host OpenMP runtime while linking.">;
 def fopenmp_new_driver : Flag<["-"], "fopenmp-new-driver">, Flags<[CC1Option]>, Group<Action_Group>,
   HelpText<"Use the new driver for OpenMP offloading.">;
+def fopenmp_device_libm : Flag<["-"], "fopenmp-device-libm">, Flags<[CC1Option]>, Group<Action_Group>,
+  HelpText<"Use the OpenMP math wrappers for offloading.">;
 def fno_optimize_sibling_calls : Flag<["-"], "fno-optimize-sibling-calls">, Group<f_Group>;
 def foptimize_sibling_calls : Flag<["-"], "foptimize-sibling-calls">, Group<f_Group>;
 defm escaping_block_tail_calls : BoolFOption<"escaping-block-tail-calls",
   CodeGenOpts<"NoEscapingBlockTailCalls">, DefaultFalse,
   NegFlag<SetTrue, [CC1Option]>, PosFlag<SetFalse>>;
 def force__cpusubtype__ALL : Flag<["-"], "force_cpusubtype_ALL">;
 def force__flat__namespace : Flag<["-"], "force_flat_namespace">;
 def force__load : Separate<["-"], "force_load">;

clang/lib/Driver/ToolChains/Clang.cpp

if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
     if (!Args.hasArg(options::OPT_nobuiltininc)) {
       // Add openmp_wrappers/* to our system include path.  This lets us wrap
       // standard library headers.
       SmallString<128> P(D.ResourceDir);
       llvm::sys::path::append(P, "include");
       llvm::sys::path::append(P, "openmp_wrappers");
       CmdArgs.push_back("-internal-isystem");
       CmdArgs.push_back(Args.MakeArgString(P));
+
+      // If using the device math library we use math wrapper functions.
+      if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
+          Args.hasArg(options::OPT_fopenmp_device_libm))
+        CmdArgs.push_back("-D__MATH_WRAPPERS__");
     }
 
     CmdArgs.push_back("-include");
     CmdArgs.push_back("__clang_openmp_device_functions.h");
   }
 
   // Add -i* options, and automatically translate to
   // -include-pch/-include-pth for transparent PCH support. It's

clang/lib/Headers/CMakeLists.txt

 )
 
 set(openmp_wrapper_files
   openmp_wrappers/math.h
   openmp_wrappers/cmath
   openmp_wrappers/complex.h
   openmp_wrappers/complex
   openmp_wrappers/__clang_openmp_device_functions.h
+  openmp_wrappers/__clang_openmp_math_forward_declares.h
+  openmp_wrappers/__clang_openmp_math.h
   openmp_wrappers/complex_cmath.h
   openmp_wrappers/new
 )
 
 set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
 set(out_files)
 set(generated_files)
 

clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h

+//===- __clang_math_forward_declares.h - Prototypes of __device__ math fns --===
+//
+// 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_H__
+#define __CLANG__OPENMP_MATH_H__
+
+#if !defined(_OPENMP)
+#error "This file is for OpenMP compilation only."
+#endif
+
+// Forward declares of all the wrappers for the standard math functions.
+#include <__clang_openmp_math_forward_declares.h>
+
+// __DEVICE__ is a helper macro with common set of attributes for the wrappers
+// 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__")
+#if defined(__cplusplus)
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#else
+#define __DEVICE__ static __attribute__((always_inline, nothrow))
+#endif
+
+// Specialized version of __DEVICE__ for functions with void return type. Needed
+// 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 defined(__cplusplus) && __cplusplus < 201402L
+#define __DEVICE_VOID__ static __attribute__((always_inline, nothrow))
+#else
+#define __DEVICE_VOID__ __DEVICE__
+#endif
+
+#if defined(__cplusplus)
+extern "C" {
+#endif
+
+__DEVICE__ int abs(int __a) { return __omp_abs(__a); }
+__DEVICE__ double fabs(double __a) { return __omp_fabs(__a); }
+__DEVICE__ double acos(double __a) { return __omp_acos(__a); }
+__DEVICE__ float acosf(float __a) { return __omp_acosf(__a); }
+__DEVICE__ double acosh(double __a) { return __omp_acosh(__a); }
+__DEVICE__ float acoshf(float __a) { return __omp_acoshf(__a); }
+__DEVICE__ double asin(double __a) { return __omp_asin(__a); }
+__DEVICE__ float asinf(float __a) { return __omp_asinf(__a); }
+__DEVICE__ double asinh(double __a) { return __omp_asinh(__a); }
+__DEVICE__ float asinhf(float __a) { return __omp_asinhf(__a); }
+__DEVICE__ double atan(double __a) { return __omp_atan(__a); }
+__DEVICE__ double atan2(double __a, double __b) {
+  return __omp_atan2(__a, __b);
+}
+__DEVICE__ float atan2f(float __a, float __b) { return __omp_atan2f(__a, __b); }
+__DEVICE__ float atanf(float __a) { return __omp_atanf(__a); }
+__DEVICE__ double atanh(double __a) { return __omp_atanh(__a); }
+__DEVICE__ float atanhf(float __a) { return __omp_atanhf(__a); }
+__DEVICE__ double cbrt(double __a) { return __omp_cbrt(__a); }
+__DEVICE__ float cbrtf(float __a) { return __omp_cbrtf(__a); }
+__DEVICE__ double ceil(double __a) { return __omp_ceil(__a); }
+__DEVICE__ float ceilf(float __a) { return __omp_ceilf(__a); }
+__DEVICE__ double copysign(double __a, double __b) {
+  return __omp_copysign(__a, __b);
+}
+__DEVICE__ float copysignf(float __a, float __b) {
+  return __omp_copysignf(__a, __b);
+}
+__DEVICE__ double cos(double __a) { return __omp_cos(__a); }
+__DEVICE__ float cosf(float __a) { return __omp_cosf(__a); }
+__DEVICE__ double cosh(double __a) { return __omp_cosh(__a); }
+__DEVICE__ float coshf(float __a) { return __omp_coshf(__a); }
+__DEVICE__ double cospi(double __a) { return __omp_cospi(__a); }
+__DEVICE__ float cospif(float __a) { return __omp_cospif(__a); }
+__DEVICE__ double cyl_bessel_i0(double __a) { return __omp_cyl_bessel_i0(__a); }
+__DEVICE__ float cyl_bessel_i0f(float __a) { return __omp_cyl_bessel_i0f(__a); }
+__DEVICE__ double cyl_bessel_i1(double __a) { return __omp_cyl_bessel_i1(__a); }
+__DEVICE__ float cyl_bessel_i1f(float __a) { return __omp_cyl_bessel_i1f(__a); }
+__DEVICE__ double erf(double __a) { return __omp_erf(__a); }
+__DEVICE__ double erfc(double __a) { return __omp_erfc(__a); }
+__DEVICE__ float erfcf(float __a) { return __omp_erfcf(__a); }
+__DEVICE__ double erfcinv(double __a) { return __omp_erfcinv(__a); }
+__DEVICE__ float erfcinvf(float __a) { return __omp_erfcinvf(__a); }
+__DEVICE__ double erfcx(double __a) { return __omp_erfcx(__a); }
+__DEVICE__ float erfcxf(float __a) { return __omp_erfcxf(__a); }
+__DEVICE__ float erff(float __a) { return __omp_erff(__a); }
+__DEVICE__ double erfinv(double __a) { return __omp_erfinv(__a); }
+__DEVICE__ float erfinvf(float __a) { return __omp_erfinvf(__a); }
+__DEVICE__ double exp(double __a) { return __omp_exp(__a); }
+__DEVICE__ double exp10(double __a) { return __omp_exp10(__a); }
+__DEVICE__ float exp10f(float __a) { return __omp_exp10f(__a); }
+__DEVICE__ double exp2(double __a) { return __omp_exp2(__a); }
+__DEVICE__ float exp2f(float __a) { return __omp_exp2f(__a); }
+__DEVICE__ float expf(float __a) { return __omp_expf(__a); }
+__DEVICE__ double expm1(double __a) { return __omp_expm1(__a); }
+__DEVICE__ float expm1f(float __a) { return __omp_expm1f(__a); }
+__DEVICE__ float fabsf(float __a) { return __omp_fabsf(__a); }
+__DEVICE__ double fdim(double __a, double __b) { return __omp_fdim(__a, __b); }
+__DEVICE__ float fdimf(float __a, float __b) { return __omp_fdimf(__a, __b); }
+__DEVICE__ double fdivide(double __a, double __b) { return __a / __b; }
+__DEVICE__ float fdividef(float __a, float __b) { return __a / __b; }
+__DEVICE__ double floor(double __f) { return __omp_floor(__f); }
+__DEVICE__ float floorf(float __f) { return __omp_floorf(__f); }
+__DEVICE__ double fma(double __a, double __b, double __c) {
+  return __omp_fma(__a, __b, __c);
+}
+__DEVICE__ float fmaf(float __a, float __b, float __c) {
+  return __omp_fmaf(__a, __b, __c);
+}
+__DEVICE__ double fmax(double __a, double __b) { return __omp_fmax(__a, __b); }
+__DEVICE__ float fmaxf(float __a, float __b) { return __omp_fmaxf(__a, __b); }
+__DEVICE__ double fmin(double __a, double __b) { return __omp_fmin(__a, __b); }
+__DEVICE__ float fminf(float __a, float __b) { return __omp_fminf(__a, __b); }
+__DEVICE__ double fmod(double __a, double __b) { return __omp_fmod(__a, __b); }
+__DEVICE__ float fmodf(float __a, float __b) { return __omp_fmodf(__a, __b); }
+__DEVICE__ double frexp(double __a, int *__b) { return __omp_frexp(__a, __b); }
+__DEVICE__ float frexpf(float __a, int *__b) { return __omp_frexpf(__a, __b); }
+__DEVICE__ double hypot(double __a, double __b) {
+  return __omp_hypot(__a, __b);
+}
+__DEVICE__ float hypotf(float __a, float __b) { return __omp_hypotf(__a, __b); }
+__DEVICE__ int ilogb(double __a) { return __omp_ilogb(__a); }
+__DEVICE__ int ilogbf(float __a) { return __omp_ilogbf(__a); }
+__DEVICE__ double j0(double __a) { return __omp_j0(__a); }
+__DEVICE__ float j0f(float __a) { return __omp_j0f(__a); }
+__DEVICE__ double j1(double __a) { return __omp_j1(__a); }
+__DEVICE__ float j1f(float __a) { return __omp_j1f(__a); }
+__DEVICE__ double jn(int __n, double __a) { return __omp_jn(__n, __a); }
+__DEVICE__ float jnf(int __n, float __a) { return __omp_jnf(__n, __a); }
+#if defined(__LP64__) || defined(_WIN64)
+__DEVICE__ long labs(long __a) { return __omp_llabs(__a); };
+#else
+__DEVICE__ long labs(long __a) { return __omp_abs(__a); };
+#endif
+__DEVICE__ double ldexp(double __a, int __b) { return __omp_ldexp(__a, __b); }
+__DEVICE__ float ldexpf(float __a, int __b) { return __omp_ldexpf(__a, __b); }
+__DEVICE__ double lgamma(double __a) { return __omp_lgamma(__a); }
+__DEVICE__ float lgammaf(float __a) { return __omp_lgammaf(__a); }
+__DEVICE__ long long llabs(long long __a) { return __omp_llabs(__a); }
+__DEVICE__ long long llmax(long long __a, long long __b) {
+  return __omp_llmax(__a, __b);
+}
+__DEVICE__ long long llmin(long long __a, long long __b) {
+  return __omp_llmin(__a, __b);
+}
+__DEVICE__ long long llrint(double __a) { return __omp_llrint(__a); }
+__DEVICE__ long long llrintf(float __a) { return __omp_llrintf(__a); }
+__DEVICE__ long long llround(double __a) { return __omp_llround(__a); }
+__DEVICE__ long long llroundf(float __a) { return __omp_llroundf(__a); }
+__DEVICE__ double round(double __a) { return __omp_round(__a); }
+__DEVICE__ float roundf(float __a) { return __omp_roundf(__a); }
+__DEVICE__ double log(double __a) { return __omp_log(__a); }
+__DEVICE__ double log10(double __a) { return __omp_log10(__a); }
+__DEVICE__ float log10f(float __a) { return __omp_log10f(__a); }
+__DEVICE__ double log1p(double __a) { return __omp_log1p(__a); }
+__DEVICE__ float log1pf(float __a) { return __omp_log1pf(__a); }
+__DEVICE__ double log2(double __a) { return __omp_log2(__a); }
+__DEVICE__ float log2f(float __a) { return __omp_log2f(__a); }
+__DEVICE__ double logb(double __a) { return __omp_logb(__a); }
+__DEVICE__ float logbf(float __a) { return __omp_logbf(__a); }
+__DEVICE__ float logf(float __a) { return __omp_logf(__a); }
+__DEVICE__ long lrint(double __a) { return __omp_lrint(__a); }
+__DEVICE__ long lrintf(float __a) { return __omp_lrintf(__a); }
+__DEVICE__ long lround(double __a) { return __omp_lround(__a); }
+__DEVICE__ long lroundf(float __a) { return __omp_lroundf(__a); }
+__DEVICE__ int max(int __a, int __b) { return __omp_max(__a, __b); }
+__DEVICE__ int min(int __a, int __b) { return __omp_min(__a, __b); }
+__DEVICE__ double modf(double __a, double *__b) { return __omp_modf(__a, __b); }
+__DEVICE__ float modff(float __a, float *__b) { return __omp_modff(__a, __b); }
+__DEVICE__ double nearbyint(double __a) { return __builtin_nearbyint(__a); }
+__DEVICE__ float nearbyintf(float __a) { return __builtin_nearbyintf(__a); }
+__DEVICE__ double nextafter(double __a, double __b) {
+  return __omp_nextafter(__a, __b);
+}
+__DEVICE__ float nextafterf(float __a, float __b) {
+  return __omp_nextafterf(__a, __b);
+}
+__DEVICE__ double norm(int __dim, const double *__t) {
+  return __omp_norm(__dim, __t);
+}
+__DEVICE__ double norm3d(double __a, double __b, double __c) {
+  return __omp_norm3d(__a, __b, __c);
+}
+__DEVICE__ float norm3df(float __a, float __b, float __c) {
+  return __omp_norm3df(__a, __b, __c);
+}
+__DEVICE__ double norm4d(double __a, double __b, double __c, double __d) {
+  return __omp_norm4d(__a, __b, __c, __d);
+}
+__DEVICE__ float norm4df(float __a, float __b, float __c, float __d) {
+  return __omp_norm4df(__a, __b, __c, __d);
+}
+__DEVICE__ double normcdf(double __a) { return __omp_normcdf(__a); }
+__DEVICE__ float normcdff(float __a) { return __omp_normcdff(__a); }
+__DEVICE__ double normcdfinv(double __a) { return __omp_normcdfinv(__a); }
+__DEVICE__ float normcdfinvf(float __a) { return __omp_normcdfinvf(__a); }
+__DEVICE__ float normf(int __dim, const float *__t) {
+  return __omp_normf(__dim, __t);
+}
+__DEVICE__ double pow(double __a, double __b) { return __omp_pow(__a, __b); }
+__DEVICE__ float powf(float __a, float __b) { return __omp_powf(__a, __b); }
+__DEVICE__ double powi(double __a, int __b) { return __omp_powi(__a, __b); }
+__DEVICE__ float powif(float __a, int __b) { return __omp_powif(__a, __b); }
+__DEVICE__ double rcbrt(double __a) { return __omp_rcbrt(__a); }
+__DEVICE__ float rcbrtf(float __a) { return __omp_rcbrtf(__a); }
+__DEVICE__ double remainder(double __a, double __b) {
+  return __omp_remainder(__a, __b);
+}
+__DEVICE__ float remainderf(float __a, float __b) {
+  return __omp_remainderf(__a, __b);
+}
+__DEVICE__ double remquo(double __a, double __b, int *__c) {
+  return __omp_remquo(__a, __b, __c);
+}
+__DEVICE__ float remquof(float __a, float __b, int *__c) {
+  return __omp_remquof(__a, __b, __c);
+}
+__DEVICE__ double rhypot(double __a, double __b) {
+  return __omp_rhypot(__a, __b);
+}
+__DEVICE__ float rhypotf(float __a, float __b) {
+  return __omp_rhypotf(__a, __b);
+}
+// __omp_rint* in libdevice is buggy and produces incorrect results.
+__DEVICE__ double rint(double __a) { return __builtin_rint(__a); }
+__DEVICE__ float rintf(float __a) { return __builtin_rintf(__a); }
+__DEVICE__ double rnorm(int __a, const double *__b) {
+  return __omp_rnorm(__a, __b);
+}
+__DEVICE__ double rnorm3d(double __a, double __b, double __c) {
+  return __omp_rnorm3d(__a, __b, __c);
+}
+__DEVICE__ float rnorm3df(float __a, float __b, float __c) {
+  return __omp_rnorm3df(__a, __b, __c);
+}
+__DEVICE__ double rnorm4d(double __a, double __b, double __c, double __d) {
+  return __omp_rnorm4d(__a, __b, __c, __d);
+}
+__DEVICE__ float rnorm4df(float __a, float __b, float __c, float __d) {
+  return __omp_rnorm4df(__a, __b, __c, __d);
+}
+__DEVICE__ float rnormf(int __dim, const float *__t) {
+  return __omp_rnormf(__dim, __t);
+}
+__DEVICE__ double rsqrt(double __a) { return __omp_rsqrt(__a); }
+__DEVICE__ float rsqrtf(float __a) { return __omp_rsqrtf(__a); }
+__DEVICE__ double scalbn(double __a, int __b) { return __omp_scalbn(__a, __b); }
+__DEVICE__ float scalbnf(float __a, int __b) { return __omp_scalbnf(__a, __b); }
+__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);
+}
+__DEVICE__ double sin(double __a) { return __omp_sin(__a); }
+__DEVICE_VOID__ void sincos(double __a, double *__s, double *__c) {
+  return __omp_sincos(__a, __s, __c);
+}
+__DEVICE_VOID__ void sincosf(float __a, float *__s, float *__c) {
+  return __omp_sincosf(__a, __s, __c);
+}
+__DEVICE_VOID__ void sincospi(double __a, double *__s, double *__c) {
+  return __omp_sincospi(__a, __s, __c);
+}
+__DEVICE_VOID__ void sincospif(float __a, float *__s, float *__c) {
+  return __omp_sincospif(__a, __s, __c);
+}
+__DEVICE__ float sinf(float __a) { return __omp_sinf(__a); }
+__DEVICE__ double sinh(double __a) { return __omp_sinh(__a); }
+__DEVICE__ float sinhf(float __a) { return __omp_sinhf(__a); }
+__DEVICE__ double sinpi(double __a) { return __omp_sinpi(__a); }
+__DEVICE__ float sinpif(float __a) { return __omp_sinpif(__a); }
+__DEVICE__ double sqrt(double __a) { return __omp_sqrt(__a); }
+__DEVICE__ float sqrtf(float __a) { return __omp_sqrtf(__a); }
+__DEVICE__ double tan(double __a) { return __omp_tan(__a); }
+__DEVICE__ float tanf(float __a) { return __omp_tanf(__a); }
+__DEVICE__ double tanh(double __a) { return __omp_tanh(__a); }
+__DEVICE__ float tanhf(float __a) { return __omp_tanhf(__a); }
+__DEVICE__ double tgamma(double __a) { return __omp_tgamma(__a); }
+__DEVICE__ float tgammaf(float __a) { return __omp_tgammaf(__a); }
+__DEVICE__ double trunc(double __a) { return __omp_trunc(__a); }
+__DEVICE__ float truncf(float __a) { return __omp_truncf(__a); }
+__DEVICE__ unsigned long long ullmax(unsigned long long __a,
+                                     unsigned long long __b) {
+  return __omp_ullmax(__a, __b);
+}
+__DEVICE__ unsigned long long ullmin(unsigned long long __a,
+                                     unsigned long long __b) {
+  return __omp_ullmin(__a, __b);
+}
+__DEVICE__ unsigned int umax(unsigned int __a, unsigned int __b) {
+  return __omp_umax(__a, __b);
+}
+__DEVICE__ unsigned int umin(unsigned int __a, unsigned int __b) {
+  return __omp_umin(__a, __b);
+}
+__DEVICE__ double y0(double __a) { return __omp_y0(__a); }
+__DEVICE__ float y0f(float __a) { return __omp_y0f(__a); }
+__DEVICE__ double y1(double __a) { return __omp_y1(__a); }
+__DEVICE__ float y1f(float __a) { return __omp_y1f(__a); }
+__DEVICE__ double yn(int __a, double __b) { return __omp_yn(__a, __b); }
+__DEVICE__ float ynf(int __a, float __b) { return __omp_ynf(__a, __b); }
+
+#if defined(__cplusplus)
+}
+#endif
+
+#endif

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_forward_declares.h

+//===- glang_math_forward_declares.h - Prototypes of evice__ math fns --===
+//
+// 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 gLANG__OPENMP_MATH_FORWARD_DECLARES_H__
+#define gLANG__OPENMP_MATH_FORWARD_DECLARES_H__
+
+#if !defined(_OPENMP)
+#error "This file is for OpenMP compilation only."
+#endif
+
+#pragma push_macro("__DEVICE__")
+#define __DEVICE__
+
+#if defined(__cplusplus)
+extern "C" {
+#endif
+
+__DEVICE__ int __omp_abs(int);
+__DEVICE__ double __omp_fabs(double);
+__DEVICE__ double __omp_acos(double);
+__DEVICE__ float __omp_acosf(float);
+__DEVICE__ double __omp_acosh(double);
+__DEVICE__ float __omp_acoshf(float);
+__DEVICE__ double __omp_asin(double);
+__DEVICE__ float __omp_asinf(float);
+__DEVICE__ double __omp_asinh(double);
+__DEVICE__ float __omp_asinhf(float);
+__DEVICE__ double __omp_atan(double);
+__DEVICE__ double __omp_atan2(double, double);
+__DEVICE__ float __omp_atan2f(float, float);
+__DEVICE__ float __omp_atanf(float);
+__DEVICE__ double __omp_atanh(double);
+__DEVICE__ float __omp_atanhf(float);
+__DEVICE__ double __omp_cbrt(double);
+__DEVICE__ float __omp_cbrtf(float);
+__DEVICE__ double __omp_ceil(double);
+__DEVICE__ float __omp_ceilf(float);
+__DEVICE__ double __omp_copysign(double, double);
+__DEVICE__ float __omp_copysignf(float, float);
+__DEVICE__ double __omp_cos(double);
+__DEVICE__ float __omp_cosf(float);
+__DEVICE__ double __omp_cosh(double);
+__DEVICE__ float __omp_coshf(float);
+__DEVICE__ double __omp_cospi(double);
+__DEVICE__ float __omp_cospif(float);
+__DEVICE__ double __omp_cyl_bessel_i0(double);
+__DEVICE__ float __omp_cyl_bessel_i0f(float);
+__DEVICE__ double __omp_cyl_bessel_i1(double);
+__DEVICE__ float __omp_cyl_bessel_i1f(float);
+__DEVICE__ double __omp_erf(double);
+__DEVICE__ double __omp_erfc(double);
+__DEVICE__ float __omp_erfcf(float);
+__DEVICE__ double __omp_erfcinv(double);
+__DEVICE__ float __omp_erfcinvf(float);
+__DEVICE__ double __omp_erfcx(double);
+__DEVICE__ float __omp_erfcxf(float);
+__DEVICE__ float __omp_erff(float);
+__DEVICE__ double __omp_erfinv(double);
+__DEVICE__ float __omp_erfinvf(float);
+__DEVICE__ double __omp_exp(double);
+__DEVICE__ double __omp_exp10(double);
+__DEVICE__ float __omp_exp10f(float);
+__DEVICE__ double __omp_exp2(double);
+__DEVICE__ float __omp_exp2f(float);
+__DEVICE__ float __omp_expf(float);
+__DEVICE__ double __omp_expm1(double);
+__DEVICE__ float __omp_expm1f(float);
+__DEVICE__ float __omp_fabsf(float);
+__DEVICE__ double __omp_fdim(double, double);
+__DEVICE__ float __omp_fdimf(float, float);
+__DEVICE__ double __omp_fdivide(double, double);
+__DEVICE__ float __omp_fdividef(float, float);
+__DEVICE__ double __omp_floor(double __f);
+__DEVICE__ float __omp_floorf(float __f);
+__DEVICE__ double __omp_fma(double, double, double);
+__DEVICE__ float __omp_fmaf(float, float, float);
+__DEVICE__ double __omp_fmax(double, double);
+__DEVICE__ float __omp_fmaxf(float, float);
+__DEVICE__ double __omp_fmin(double, double);
+__DEVICE__ float __omp_fminf(float, float);
+__DEVICE__ double __omp_fmod(double, double);
+__DEVICE__ float __omp_fmodf(float, float);
+__DEVICE__ double __omp_frexp(double, int *);
+__DEVICE__ float __omp_frexpf(float, int *);
+__DEVICE__ double __omp_hypot(double, double);
+__DEVICE__ float __omp_hypotf(float, float);
+__DEVICE__ int __omp_ilogb(double);
+__DEVICE__ int __omp_ilogbf(float);
+__DEVICE__ double __omp_j0(double);
+__DEVICE__ float __omp_j0f(float);
+__DEVICE__ double __omp_j1(double);
+__DEVICE__ float __omp_j1f(float);
+__DEVICE__ double __omp_jn(int __n, double);
+__DEVICE__ float __omp_jnf(int __n, float);
+__DEVICE__ long __omp_labs(long);
+__DEVICE__ double __omp_ldexp(double, int);
+__DEVICE__ float __omp_ldexpf(float, int);
+__DEVICE__ double __omp_lgamma(double);
+__DEVICE__ float __omp_lgammaf(float);
+__DEVICE__ long long __omp_llabs(long long);
+__DEVICE__ long long __omp_llmax(long long, long long);
+__DEVICE__ long long __omp_llmin(long long, long long);
+__DEVICE__ long long __omp_llrint(double);
+__DEVICE__ long long __omp_llrintf(float);
+__DEVICE__ long long __omp_llround(double);
+__DEVICE__ long long __omp_llroundf(float);
+__DEVICE__ double __omp_round(double);
+__DEVICE__ float __omp_roundf(float);
+__DEVICE__ double __omp_log(double);
+__DEVICE__ double __omp_log10(double);
+__DEVICE__ float __omp_log10f(float);
+__DEVICE__ double __omp_log1p(double);
+__DEVICE__ float __omp_log1pf(float);
+__DEVICE__ double __omp_log2(double);
+__DEVICE__ float __omp_log2f(float);
+__DEVICE__ double __omp_logb(double);
+__DEVICE__ float __omp_logbf(float);
+__DEVICE__ float __omp_logf(float);
+__DEVICE__ long __omp_lrint(double);
+__DEVICE__ long __omp_lrintf(float);
+__DEVICE__ long __omp_lround(double);
+__DEVICE__ long __omp_lroundf(float);
+__DEVICE__ int __omp_max(int, int);
+__DEVICE__ int __omp_min(int, int);
+__DEVICE__ double __omp_modf(double, double *);
+__DEVICE__ float __omp_modff(float, float *);
+__DEVICE__ double __omp_nearbyint(double);
+__DEVICE__ float __omp_nearbyintf(float);
+__DEVICE__ double __omp_nextafter(double, double);
+__DEVICE__ float __omp_nextafterf(float, float);
+__DEVICE__ double __omp_norm(int im, const double *);
+__DEVICE__ double __omp_norm3d(double, double, double);
+__DEVICE__ float __omp_norm3df(float, float, float);
+__DEVICE__ double __omp_norm4d(double, double, double, double);
+__DEVICE__ float __omp_norm4df(float, float, float, float);
+__DEVICE__ double __omp_normcdf(double);
+__DEVICE__ float __omp_normcdff(float);
+__DEVICE__ double __omp_normcdfinv(double);
+__DEVICE__ float __omp_normcdfinvf(float);
+__DEVICE__ float __omp_normf(int im, const float *);
+__DEVICE__ double __omp_pow(double, double);
+__DEVICE__ float __omp_powf(float, float);
+__DEVICE__ double __omp_powi(double, int);
+__DEVICE__ float __omp_powif(float, int);
+__DEVICE__ double __omp_rcbrt(double);
+__DEVICE__ float __omp_rcbrtf(float);
+__DEVICE__ double __omp_remainder(double, double);
+__DEVICE__ float __omp_remainderf(float, float);
+__DEVICE__ double __omp_remquo(double, double, int *);
+__DEVICE__ float __omp_remquof(float, float, int *);
+__DEVICE__ double __omp_rhypot(double, double);
+__DEVICE__ float __omp_rhypotf(float, float);
+__DEVICE__ double __omp_rint(double);
+__DEVICE__ float __omp_rintf(float);
+__DEVICE__ double __omp_rnorm(int, const double *);
+__DEVICE__ double __omp_rnorm3d(double, double, double);
+__DEVICE__ float __omp_rnorm3df(float, float, float);
+__DEVICE__ double __omp_rnorm4d(double, double, double, double);
+__DEVICE__ float __omp_rnorm4df(float, float, float, float);
+__DEVICE__ float __omp_rnormf(int im, const float *);
+__DEVICE__ double __omp_rsqrt(double);
+__DEVICE__ float __omp_rsqrtf(float);
+__DEVICE__ double __omp_scalbn(double, int);
+__DEVICE__ float __omp_scalbnf(float, int);
+__DEVICE__ double __omp_scalbln(double, long);
+__DEVICE__ float __omp_scalblnf(float, long);
+__DEVICE__ double __omp_sin(double);
+__DEVICE__ void __omp_sincos(double, double *, double *);
+__DEVICE__ void __omp_sincosf(float, float *, float *);
+__DEVICE__ void __omp_sincospi(double, double *, double *);
+__DEVICE__ void __omp_sincospif(float, float *, float *);
+__DEVICE__ float __omp_sinf(float);
+__DEVICE__ double __omp_sinh(double);
+__DEVICE__ float __omp_sinhf(float);
+__DEVICE__ double __omp_sinpi(double);
+__DEVICE__ float __omp_sinpif(float);
+__DEVICE__ double __omp_sqrt(double);
+__DEVICE__ float __omp_sqrtf(float);
+__DEVICE__ double __omp_tan(double);
+__DEVICE__ float __omp_tanf(float);
+__DEVICE__ double __omp_tanh(double);
+__DEVICE__ float __omp_tanhf(float);
+__DEVICE__ double __omp_tgamma(double);
+__DEVICE__ float __omp_tgammaf(float);
+__DEVICE__ double __omp_trunc(double);
+__DEVICE__ float __omp_truncf(float);
+__DEVICE__ unsigned long long __omp_ullmax(unsigned long long,
+                                           unsigned long long);
+__DEVICE__ unsigned long long __omp_ullmin(unsigned long long,
+                                           unsigned long long);
+__DEVICE__ unsigned int __omp_umax(unsigned int, unsigned int);
+__DEVICE__ unsigned int __omp_umin(unsigned int, unsigned int);
+__DEVICE__ double __omp_y0(double);
+__DEVICE__ float __omp_y0f(float);
+__DEVICE__ double __omp_y1(double);
+__DEVICE__ float __omp_y1f(float);
+__DEVICE__ double __omp_yn(int, double);
+__DEVICE__ float __omp_ynf(int, float);
+
+#if defined(__cplusplus)
+}
+#endif
+
+#pragma pop_macro("__DEVICE__")
+
+#endif

clang/lib/Headers/openmp_wrappers/math.h

 // We need limits.h for __clang_cuda_math.h below and because it should not hurt
 // we include it eagerly here.
 #include <limits.h>
 
 // We need stdlib.h because (for now) __clang_cuda_math.h below declares `abs`
 // which should live in stdlib.h.
 #include <stdlib.h>
 
+// Use the OpenMP math wrappers and library to call device math routines.
+#if defined(__MATH_WRAPPERS__)
+// Math routines on the device will call an OpenMP wrapper to be defined later.
+#pragma omp begin declare variant match(                                       \
+    device = {arch(nvptx, nvptx64, amdgcn)},                                   \
+    implementation = {extension(match_any)})
+
+#include <__clang_openmp_math.h>
+
+#pragma omp end declare variant
+#else
 #pragma omp begin declare variant match(                                       \
     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
 
 #ifdef __AMDGCN__
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 #define __OPENMP_AMDGCN__
 #include <__clang_hip_math.h>
 #undef __OPENMP_AMDGCN__
 
 #pragma omp end declare variant
 #endif
+#endif
 
 #endif

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 -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_device_functions.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 -check-prefix CHECK-YES %s
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix CHECK-LIBM %s
 
 #include <stdlib.h>
 #include <math.h>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK-LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK-YES: call double @__nv_pow(double
+    // CHECK-LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK-YES: call double @__nv_modf(double
+    // CHECK-LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK-YES: call double @__nv_fabs(double
+    // CHECK-LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK-LIBM: call i32 @__omp_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 -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_device_functions.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
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix CHECK-LIBM %s
 
 #include <cmath>
 #include <cstdlib>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK-LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK-YES: call double @__nv_pow(double
+    // CHECK-LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK-YES: call double @__nv_modf(double
+    // CHECK-LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK-YES: call double @__nv_fabs(double
+    // CHECK-LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK-LIBM: call i32 @__omp_abs(i32
     double l5 = abs((int)a1);
   }
 }

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 -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_device_functions.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
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix CHECK-LIBM %s
 
 #include <cmath>
 #include <cstdlib>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK-LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK-YES: call double @__nv_pow(double
+    // CHECK-LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK-YES: call double @__nv_modf(double
+    // CHECK-LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK-YES: call double @__nv_fabs(double
+    // CHECK-LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK-LIBM: call i32 @__omp_abs(i32
     double l5 = abs((int)a1);
   }
 }

clang/test/Headers/nvptx_device_math_functions.c

 // Test calling of device math functions.
 ///==========================================================================///
 
 // REQUIRES: nvptx-registered-target
 
 // 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_device_functions.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/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix LIBM %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_device_functions.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/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix LIBM %s
 
 #ifdef __cplusplus
 #include <cstdlib>
 #include <cmath>
 #else
 #include <stdlib.h>
 #include <math.h>
 #endif
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK: call double @__nv_sqrt(double
+    // LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK: call double @__nv_pow(double
+    // LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK: call double @__nv_modf(double
+    // LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK: call double @__nv_fabs(double
+    // LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK: call i32 @__nv_abs(i32
+    // LIBM: call i32 @__omp_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 -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_device_functions.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 <cmath>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK-LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK-YES: call double @__nv_pow(double
+    // CHECK-LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK-YES: call double @__nv_modf(double
+    // CHECK-LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK-YES: call double @__nv_fabs(double
+    // CHECK-LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK-LIBM: call i32 @__omp_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 -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_device_functions.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
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.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 -D__MATH_WRAPPERS__ -o - | FileCheck -check-prefix CHECK-LIBM %s
 
 #include <cstdlib>
 #include <cmath>
 
 void test_sqrt(double a1) {
   #pragma omp target
   {
     // CHECK-YES: call double @__nv_sqrt(double
+    // CHECK-LIBM: call double @__omp_sqrt(double
     double l1 = sqrt(a1);
     // CHECK-YES: call double @__nv_pow(double
+    // CHECK-LIBM: call double @__omp_pow(double
     double l2 = pow(a1, a1);
     // CHECK-YES: call double @__nv_modf(double
+    // CHECK-LIBM: call double @__omp_modf(double
     double l3 = modf(a1 + 3.5, &a1);
     // CHECK-YES: call double @__nv_fabs(double
+    // CHECK-LIBM: call double @__omp_fabs(double
     double l4 = fabs(a1);
     // CHECK-YES: call i32 @__nv_abs(i32
+    // CHECK-LIBM: call i32 @__omp_abs(i32
     double l5 = abs((int)a1);
   }
 }