[clang] [OFFLOAD] Introduce OpenMP cmath wrappers for SPIRV backend (PR #195386)
via cfe-commits
cfe-commits at lists.llvm.org
Wed May 6 22:50:25 PDT 2026
https://github.com/fineg74 updated https://github.com/llvm/llvm-project/pull/195386
>From 221662cbad181eb35111eceb916360da54e0582d Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Fri, 1 May 2026 10:26:07 -0700
Subject: [PATCH 1/4] Introduce cmath wrappers for SPIRV backend
---
clang/lib/Headers/CMakeLists.txt | 1 +
clang/lib/Headers/__clang_spirv_cmath.h | 505 ++++++++++++++++++++++++
clang/lib/Headers/openmp_wrappers/cmath | 52 +++
3 files changed, 558 insertions(+)
create mode 100644 clang/lib/Headers/__clang_spirv_cmath.h
diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
index d60ae2b5961e0..4252f8c4685b6 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -146,6 +146,7 @@ set(spirv_files
__clang_spirv_builtins.h
__clang_spirv_libdevice_declares.h
__clang_spirv_math.h
+ __clang_spirv_cmath.h
)
set(systemz_files
diff --git a/clang/lib/Headers/__clang_spirv_cmath.h b/clang/lib/Headers/__clang_spirv_cmath.h
new file mode 100644
index 0000000000000..b2bdb804b4956
--- /dev/null
+++ b/clang/lib/Headers/__clang_spirv_cmath.h
@@ -0,0 +1,505 @@
+ /*===---- __clang_spirv_cmath.h - SPIRV cmath decls -----------------------===
+ *
+ * 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_SPIRV_CMATH_H__
+#define __CLANG_SPIRV_CMATH_H__
+
+#if !defined(__SPIRV__) && !defined(__OPENMP_SPIRV__)
+#error "This file is for SPIRV OpenMP device compilation only."
+#endif
+
+#if defined(__cplusplus)
+#include <limits>
+#include <type_traits>
+#include <utility>
+#endif
+#include <limits.h>
+#include <stdint.h>
+
+#pragma push_macro("__DEVICE__")
+#ifdef __OPENMP_SPIRV__
+#if defined(__cplusplus)
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#else
+#define __DEVICE__ static __attribute__((always_inline, nothrow))
+#endif
+#else
+#define __DEVICE__ static __device__ __forceinline__
+#endif
+
+__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
+__DEVICE__ float sin(float __x) { return ::sinf(__x); }
+__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
+__DEVICE__ float cos(float __x) { return ::cosf(__x); }
+__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
+__DEVICE__ double abs(double __x) { return ::fabs(__x); }
+__DEVICE__ float abs(float __x) { return ::fabsf(__x); }
+__DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
+__DEVICE__ long abs(long __n) { return ::labs(__n); }
+__DEVICE__ float fma(float __x, float __y, float __z) {
+ return ::fmaf(__x, __y, __z);
+}
+__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);
+}
+__DEVICE__ float frexp(float __arg, int *__exp) {
+ return ::frexpf(__arg, __exp);
+}
+__DEVICE__ float acos(float __x) { return ::acosf(__x); }
+__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
+__DEVICE__ float asin(float __x) { return ::asinf(__x); }
+__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
+__DEVICE__ float atan(float __x) { return ::atanf(__x); }
+__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
+__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
+__DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
+__DEVICE__ float exp(float __x) { return ::expf(__x); }
+__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
+__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
+__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
+__DEVICE__ float floor(float __x) { return ::floorf(__x); }
+__DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
+__DEVICE__ float fmax(float __x, float __y) { return ::fmaxf(__x, __y); }
+__DEVICE__ float fmin(float __x, float __y) { return ::fminf(__x, __y); }
+__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
+
+#if defined(__OPENMP_SPIRV__)
+// For OpenMP we work around some old system headers that have non-conforming
+// `isinf(float)` and `isnan(float)` implementations that return an `int`. We do
+// this by providing two versions of these functions, differing only in the
+// return type. To avoid conflicting definitions we disable implicit base
+// function generation. That means we will end up with two specializations, one
+// per type, but only one has a base function defined by the system header.
+#pragma omp begin declare variant match( \
+ implementation = {extension(disable_implicit_base)})
+
+// FIXME: We lack an extension to customize the mangling of the variants, e.g.,
+// add a suffix. This means we would clash with the names of the variants
+// (note that we do not create implicit base functions here). To avoid
+// this clash we add a new trait to some of them that is always true
+// (this is LLVM after all ;)). It will only influence the mangled name
+// of the variants inside the inner region and avoid the clash.
+#pragma omp begin declare variant match(implementation = {vendor(llvm)})
+
+__DEVICE__ int isinf(float __x) { return ::__isinff(__x); }
+__DEVICE__ int isinf(double __x) { return ::__isinf(__x); }
+__DEVICE__ int isfinite(float __x) { return ::__finitef(__x); }
+__DEVICE__ int isfinite(double __x) { return ::__finite(__x); }
+__DEVICE__ int isnan(float __x) { return ::__isnanf(__x); }
+__DEVICE__ int isnan(double __x) { return ::__isnan(__x); }
+
+#pragma omp end declare variant
+#endif // defined(__OPENMP_SPIRV__)
+
+__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
+__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
+__DEVICE__ bool isfinite(float __x) { return ::__finitef(__x); }
+__DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
+__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
+__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
+
+#if defined(__OPENMP_SPIRV__)
+#pragma omp end declare variant
+#endif // defined(__OPENMP_SPIRV__)
+
+__DEVICE__ bool isgreater(float __x, float __y) {
+ return __builtin_isgreater(__x, __y);
+}
+__DEVICE__ bool isgreater(double __x, double __y) {
+ return __builtin_isgreater(__x, __y);
+}
+__DEVICE__ bool isgreaterequal(float __x, float __y) {
+ return __builtin_isgreaterequal(__x, __y);
+}
+__DEVICE__ bool isgreaterequal(double __x, double __y) {
+ return __builtin_isgreaterequal(__x, __y);
+}
+__DEVICE__ bool isless(float __x, float __y) {
+ return __builtin_isless(__x, __y);
+}
+__DEVICE__ bool isless(double __x, double __y) {
+ return __builtin_isless(__x, __y);
+}
+__DEVICE__ bool islessequal(float __x, float __y) {
+ return __builtin_islessequal(__x, __y);
+}
+__DEVICE__ bool islessequal(double __x, double __y) {
+ return __builtin_islessequal(__x, __y);
+}
+__DEVICE__ bool islessgreater(float __x, float __y) {
+ return __builtin_islessgreater(__x, __y);
+}
+__DEVICE__ bool islessgreater(double __x, double __y) {
+ return __builtin_islessgreater(__x, __y);
+}
+__DEVICE__ bool isnormal(float __x) {
+ return __builtin_isnormal(__x);
+}
+__DEVICE__ bool isnormal(double __x) {
+ return __builtin_isnormal(__x);
+}
+__DEVICE__ bool isunordered(float __x, float __y) {
+ return __builtin_isunordered(__x, __y);
+}
+__DEVICE__ bool isunordered(double __x, double __y) {
+ return __builtin_isunordered(__x, __y);
+}
+__DEVICE__ float modf(float __x, float *__iptr) {
+ return ::modff(__x, __iptr);
+}
+__DEVICE__ float pow(float __base, int __iexp) {
+ return ::powif(__base, __iexp);
+}
+__DEVICE__ double pow(double __base, int __iexp) {
+ return ::powi(__base, __iexp);
+}
+__DEVICE__ float remquo(float __x, float __y, int *__quo) {
+ return ::remquof(__x, __y, __quo);
+}
+__DEVICE__ float scalbln(float __x, long int __n) {
+ return ::scalblnf(__x, __n);
+}
+__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
+__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
+__DEVICE__ float ldexp(float __arg, int __exp) {
+ return ::ldexpf(__arg, __exp);
+}
+__DEVICE__ float log(float __x) { return ::logf(__x); }
+__DEVICE__ float log10(float __x) { return ::log10f(__x); }
+__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
+__DEVICE__ float log2(float __x) { return ::log2f(__x); }
+__DEVICE__ float logb(float __x) { return ::logbf(__x); }
+
+__DEVICE__ float pow(float __base, float __exp) {
+ return ::powf(__base, __exp);
+}
+__DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
+__DEVICE__ float tan(float __x) { return ::tanf(__x); }
+__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
+__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
+__DEVICE__ float copysign(float __a, float __b) { return ::copysignf(__a, __b); }
+__DEVICE__ float erf(float __x) { return ::erff(__x); }
+__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
+__DEVICE__ float fdim(float __a, float __b) { return ::fdimf(__a, __b); }
+__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
+__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
+__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
+__DEVICE__ long long llrint(float __x) { return ::llrintf(__x); }
+__DEVICE__ long long llround(float __x) { return ::llroundf(__x); }
+__DEVICE__ long lrint(float __x) { return ::lrintf(__x); }
+__DEVICE__ long lround(float __x) { return ::lroundf(__x); }
+__DEVICE__ float rint(float __x) { return ::rintf(__x); }
+__DEVICE__ float round(float __x) { return ::roundf(__x); }
+__DEVICE__ float trunc(float __x) { return ::truncf(__x); }
+__DEVICE__ float nearbyint(float __x) { return ::nearbyintf(__x); }
+__DEVICE__ float nextafter(float __a, float __b) { return ::nextafterf(__a, __b); }
+__DEVICE__ float remainder(float __a, float __b) { return ::remainderf(__a, __b); }
+__DEVICE__ float scalbn(float __a, int __b) { return ::scalbnf(__a, __b); }
+
+#ifndef __OPENMP_SPIRV__
+#pragma push_macro("__SPIRV_OVERLOAD1")
+#pragma push_macro("__SPIRV_OVERLOAD2")
+
+// __SPIRV_OVERLOAD1 is used to resolve function calls with integer argument to
+// avoid compilation error due to ambiguity. e.g. floor(5) is resolved with
+// floor(double).
+#define __SPIRV_OVERLOAD1(__retty, __fn) \
+ template <typename __T> \
+ __DEVICE__ \
+ std::enable_if_t<std::numeric_limits<__T>::is_integer, __retty> \
+ __fn(__T __x) { \
+ return ::__fn((double)__x); \
+ }
+
+#define __SPIRV_OVERLOAD2(__retty, __fn) \
+ template <typename __T1, typename __T2> \
+ __DEVICE__ \
+ std::enable_if_t<std::numeric_limits<__T1>::is_specialized && \
+ std::numeric_limits<__T2>::is_specialized, \
+ __retty> \
+ __fn(__T1 __x, __T2 __y) { \
+ return __fn((double)__x, (double)__y); \
+ }
+
+__SPIRV_OVERLOAD1(double, acos)
+__SPIRV_OVERLOAD1(double, acosh)
+__SPIRV_OVERLOAD1(double, asin)
+__SPIRV_OVERLOAD1(double, asinh)
+__SPIRV_OVERLOAD1(double, atan)
+__SPIRV_OVERLOAD2(double, atan2)
+__SPIRV_OVERLOAD1(double, atanh)
+__SPIRV_OVERLOAD1(double, cbrt)
+__SPIRV_OVERLOAD1(double, ceil)
+__SPIRV_OVERLOAD2(double, copysign)
+__SPIRV_OVERLOAD1(double, cos)
+__SPIRV_OVERLOAD1(double, cosh)
+__SPIRV_OVERLOAD1(double, erf)
+__SPIRV_OVERLOAD1(double, erfc)
+__SPIRV_OVERLOAD1(double, exp)
+__SPIRV_OVERLOAD1(double, exp2)
+__SPIRV_OVERLOAD1(double, expm1)
+__SPIRV_OVERLOAD1(double, fabs)
+__SPIRV_OVERLOAD2(double, fdim)
+__SPIRV_OVERLOAD1(double, floor)
+__SPIRV_OVERLOAD2(double, fmax)
+__SPIRV_OVERLOAD2(double, fmin)
+__SPIRV_OVERLOAD2(double, fmod)
+__SPIRV_OVERLOAD1(int, fpclassify)
+__SPIRV_OVERLOAD2(double, hypot)
+__SPIRV_OVERLOAD1(int, ilogb)
+__SPIRV_OVERLOAD1(bool, isfinite)
+__SPIRV_OVERLOAD2(bool, isgreater)
+__SPIRV_OVERLOAD2(bool, isgreaterequal)
+__SPIRV_OVERLOAD1(bool, isinf)
+__SPIRV_OVERLOAD2(bool, isless)
+__SPIRV_OVERLOAD2(bool, islessequal)
+__SPIRV_OVERLOAD2(bool, islessgreater)
+__SPIRV_OVERLOAD1(bool, isnan)
+__SPIRV_OVERLOAD1(bool, isnormal)
+__SPIRV_OVERLOAD2(bool, isunordered)
+__SPIRV_OVERLOAD1(double, lgamma)
+__SPIRV_OVERLOAD1(double, log)
+__SPIRV_OVERLOAD1(double, log10)
+__SPIRV_OVERLOAD1(double, log1p)
+__SPIRV_OVERLOAD1(double, log2)
+__SPIRV_OVERLOAD1(double, logb)
+__SPIRV_OVERLOAD1(long long, llrint)
+__SPIRV_OVERLOAD1(long long, llround)
+__SPIRV_OVERLOAD1(long, lrint)
+__SPIRV_OVERLOAD1(long, lround)
+__SPIRV_OVERLOAD1(double, nearbyint)
+__SPIRV_OVERLOAD2(double, nextafter)
+__SPIRV_OVERLOAD2(double, pow)
+__SPIRV_OVERLOAD2(double, remainder)
+__SPIRV_OVERLOAD1(double, rint)
+__SPIRV_OVERLOAD1(double, round)
+__SPIRV_OVERLOAD1(bool, signbit)
+__SPIRV_OVERLOAD1(double, sin)
+__SPIRV_OVERLOAD1(double, sinh)
+__SPIRV_OVERLOAD1(double, sqrt)
+__SPIRV_OVERLOAD1(double, tan)
+__SPIRV_OVERLOAD1(double, tanh)
+__SPIRV_OVERLOAD1(double, tgamma)
+__SPIRV_OVERLOAD1(double, trunc)
+
+// Overload these but don't add them to std, they are not part of cmath.
+__SPIRV_OVERLOAD2(double, max)
+__SPIRV_OVERLOAD2(double, min)
+
+template <typename __T1, typename __T2, typename __T3>
+__DEVICE__ std::enable_if_t<
+ std::numeric_limits<__T1>::is_specialized &&
+ std::numeric_limits<__T2>::is_specialized &&
+ std::numeric_limits<__T3>::is_specialized,
+ double>
+fma(__T1 __x, __T2 __y, __T3 __z) {
+ return ::fma((double)__x, (double)__y, (double)__z);
+}
+
+
+template <typename __T>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
+ frexp(__T __x, int *__exp) {
+ return ::frexp((double)__x, __exp);
+}
+
+template <typename __T>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
+ ldexp(__T __x, int __exp) {
+ return ::ldexp((double)__x, __exp);
+}
+
+template <typename __T>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
+ modf(__T __x, double *__exp) {
+ return ::modf((double)__x, __exp);
+}
+
+template <typename __T1, typename __T2>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized &&
+ std::numeric_limits<__T2>::is_specialized,
+ double>
+ remquo(__T1 __x, __T2 __y, int *__quo) {
+ return ::remquo((double)__x, (double)__y, __quo);
+}
+
+template <typename __T>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
+ scalbln(__T __x, long int __exp) {
+ return ::scalbln((double)__x, __exp);
+}
+
+template <typename __T>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
+ scalbn(__T __x, int __exp) {
+ return ::scalbn((double)__x, __exp);
+}
+
+#pragma pop_macro("__SPIRV_OVERLOAD1")
+#pragma pop_macro("__SPIRV_OVERLOAD2")
+
+// Define these overloads inside the namespace our standard library uses.
+
+#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
+_LIBCPP_BEGIN_NAMESPACE_STD
+#else
+namespace std {
+#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
+_GLIBCXX_BEGIN_NAMESPACE_VERSION
+#endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
+#endif // _LIBCPP_BEGIN_NAMESPACE_STD
+
+// Pull the new overloads we defined above into namespace std.
+// using ::abs; - This may be considered for C++.
+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 ::fabs;
+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 ::isless;
+using ::islessequal;
+using ::islessgreater;
+using ::isnormal;
+using ::isunordered;
+using ::ldexp;
+using ::lgamma;
+using ::llrint;
+using ::llround;
+using ::log;
+using ::log10;
+using ::log1p;
+using ::log2;
+using ::logb;
+using ::lrint;
+using ::lround;
+using ::modf;
+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;
+
+// Well this is fun: We need to pull these symbols in for libc++, but we can't
+// pull them in with libstdc++, because its ::isinf and ::isnan are different
+// than its std::isinf and std::isnan.
+#ifndef __GLIBCXX__
+using ::isinf;
+using ::isnan;
+#endif
+
+// Finally, pull the "foobarf" functions that HIP defines into std.
+using ::acosf;
+using ::acoshf;
+using ::asinf;
+using ::asinhf;
+using ::atan2f;
+using ::atanf;
+using ::atanhf;
+using ::cbrtf;
+using ::ceilf;
+using ::copysignf;
+using ::cosf;
+using ::coshf;
+using ::erfcf;
+using ::erff;
+using ::exp2f;
+using ::expf;
+using ::expm1f;
+using ::fabsf;
+using ::fdimf;
+using ::floorf;
+using ::fmaf;
+using ::fmaxf;
+using ::fminf;
+using ::fmodf;
+using ::frexpf;
+using ::hypotf;
+using ::ilogbf;
+using ::ldexpf;
+using ::lgammaf;
+using ::llrintf;
+using ::llroundf;hfgh fghdggf h
+using ::log10f;
+using ::log1pf;
+using ::log2f;
+using ::logbf;
+using ::logf;
+using ::lrintf;
+using ::lroundf;
+using ::modff;
+using ::nearbyintf;
+using ::nextafterf;
+using ::powf;
+using ::remainderf;
+using ::remquof;
+using ::rintf;
+using ::roundf;
+using ::scalblnf;
+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 // _GLIBCXX_BEGIN_NAMESPACE_VERSION
+} // namespace std
+#endif // _LIBCPP_END_NAMESPACE_STD
+#endif // ifndef __OPENMP_SPIRV__
+#endif // __CLANG_SPIRV_CMATH_H__
\ No newline at end of file
diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath
index e1b71516e72c2..e6b887ff70507 100644
--- a/clang/lib/Headers/openmp_wrappers/cmath
+++ b/clang/lib/Headers/openmp_wrappers/cmath
@@ -129,4 +129,56 @@ __DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
#pragma omp end declare variant
#endif // __AMDGCN__
+#ifdef __SPIRV__
+#pragma omp begin declare variant match(device = {arch(spirv64)})
+
+#define __OPENMP_SPIRV__
+
+#include <__clang_spirv_cmath.h>
+
+
+#undef __OPENMP_SPIRV__
+
+// Define overloads otherwise which are absent
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+
+__DEVICE__ float acos(float __x) { return ::acosf(__x); }
+__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
+__DEVICE__ float asin(float __x) { return ::asinf(__x); }
+__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
+__DEVICE__ float atan(float __x) { return ::atanf(__x); }
+__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
+__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
+__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
+__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
+__DEVICE__ float erf(float __x) { return ::erff(__x); }
+__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
+__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
+__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
+__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); }
+__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
+__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
+__DEVICE__ float ldexp(float __arg, int __exp) {
+ return ::ldexpf(__arg, __exp);
+}
+__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
+__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
+__DEVICE__ float logb(float __x) { return ::logbf(__x); }
+__DEVICE__ float nextafter(float __x, float __y) {
+ return ::nextafterf(__x, __y);
+}
+__DEVICE__ float remainder(float __x, float __y) {
+ return ::remainderf(__x, __y);
+}
+__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
+__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
+__DEVICE__ float tan(float __x) { return ::tanf(__x); }
+__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
+__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
+
+#undef __DEVICE__
+
+#pragma omp end declare variant
+#endif // __SPIRV__
+
#endif
>From 894a16648df0fa07bbf24086331af1e1706923c3 Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Fri, 1 May 2026 14:51:59 -0700
Subject: [PATCH 2/4] Fix several minor issues
---
clang/lib/Headers/__clang_spirv_cmath.h | 71 ++++++++++++-------------
clang/lib/Headers/openmp_wrappers/cmath | 43 +--------------
2 files changed, 35 insertions(+), 79 deletions(-)
diff --git a/clang/lib/Headers/__clang_spirv_cmath.h b/clang/lib/Headers/__clang_spirv_cmath.h
index b2bdb804b4956..0e2bc899fc1c5 100644
--- a/clang/lib/Headers/__clang_spirv_cmath.h
+++ b/clang/lib/Headers/__clang_spirv_cmath.h
@@ -1,4 +1,4 @@
- /*===---- __clang_spirv_cmath.h - SPIRV cmath decls -----------------------===
+/*===---- __clang_spirv_cmath.h - SPIRV cmath decls -----------------------===
*
* Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
* See https://llvm.org/LICENSE.txt for license information.
@@ -33,7 +33,6 @@
#define __DEVICE__ static __device__ __forceinline__
#endif
-__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
__DEVICE__ float sin(float __x) { return ::sinf(__x); }
__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
__DEVICE__ float cos(float __x) { return ::cosf(__x); }
@@ -82,7 +81,7 @@ __DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
// function generation. That means we will end up with two specializations, one
// per type, but only one has a base function defined by the system header.
#pragma omp begin declare variant match( \
- implementation = {extension(disable_implicit_base)})
+ implementation = {extension(disable_implicit_base)})
// FIXME: We lack an extension to customize the mangling of the variants, e.g.,
// add a suffix. This means we would clash with the names of the variants
@@ -143,21 +142,15 @@ __DEVICE__ bool islessgreater(float __x, float __y) {
__DEVICE__ bool islessgreater(double __x, double __y) {
return __builtin_islessgreater(__x, __y);
}
-__DEVICE__ bool isnormal(float __x) {
- return __builtin_isnormal(__x);
-}
-__DEVICE__ bool isnormal(double __x) {
- return __builtin_isnormal(__x);
-}
+__DEVICE__ bool isnormal(float __x) { return __builtin_isnormal(__x); }
+__DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
__DEVICE__ bool isunordered(float __x, float __y) {
return __builtin_isunordered(__x, __y);
}
__DEVICE__ bool isunordered(double __x, double __y) {
return __builtin_isunordered(__x, __y);
}
-__DEVICE__ float modf(float __x, float *__iptr) {
- return ::modff(__x, __iptr);
-}
+__DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
__DEVICE__ float pow(float __base, int __iexp) {
return ::powif(__base, __iexp);
}
@@ -171,7 +164,7 @@ __DEVICE__ float scalbln(float __x, long int __n) {
return ::scalblnf(__x, __n);
}
__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
-__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
+__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
__DEVICE__ float ldexp(float __arg, int __exp) {
return ::ldexpf(__arg, __exp);
}
@@ -188,7 +181,9 @@ __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
__DEVICE__ float tan(float __x) { return ::tanf(__x); }
__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
-__DEVICE__ float copysign(float __a, float __b) { return ::copysignf(__a, __b); }
+__DEVICE__ float copysign(float __a, float __b) {
+ return ::copysignf(__a, __b);
+}
__DEVICE__ float erf(float __x) { return ::erff(__x); }
__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
__DEVICE__ float fdim(float __a, float __b) { return ::fdimf(__a, __b); }
@@ -203,8 +198,12 @@ __DEVICE__ float rint(float __x) { return ::rintf(__x); }
__DEVICE__ float round(float __x) { return ::roundf(__x); }
__DEVICE__ float trunc(float __x) { return ::truncf(__x); }
__DEVICE__ float nearbyint(float __x) { return ::nearbyintf(__x); }
-__DEVICE__ float nextafter(float __a, float __b) { return ::nextafterf(__a, __b); }
-__DEVICE__ float remainder(float __a, float __b) { return ::remainderf(__a, __b); }
+__DEVICE__ float nextafter(float __a, float __b) {
+ return ::nextafterf(__a, __b);
+}
+__DEVICE__ float remainder(float __a, float __b) {
+ return ::remainderf(__a, __b);
+}
__DEVICE__ float scalbn(float __a, int __b) { return ::scalbnf(__a, __b); }
#ifndef __OPENMP_SPIRV__
@@ -216,19 +215,17 @@ __DEVICE__ float scalbn(float __a, int __b) { return ::scalbnf(__a, __b); }
// floor(double).
#define __SPIRV_OVERLOAD1(__retty, __fn) \
template <typename __T> \
- __DEVICE__ \
- std::enable_if_t<std::numeric_limits<__T>::is_integer, __retty> \
- __fn(__T __x) { \
+ __DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, __retty> \
+ __fn(__T __x) { \
return ::__fn((double)__x); \
}
#define __SPIRV_OVERLOAD2(__retty, __fn) \
template <typename __T1, typename __T2> \
- __DEVICE__ \
- std::enable_if_t<std::numeric_limits<__T1>::is_specialized && \
- std::numeric_limits<__T2>::is_specialized, \
- __retty> \
- __fn(__T1 __x, __T2 __y) { \
+ __DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized && \
+ std::numeric_limits<__T2>::is_specialized, \
+ __retty> \
+ __fn(__T1 __x, __T2 __y) { \
return __fn((double)__x, (double)__y); \
}
@@ -298,51 +295,49 @@ __SPIRV_OVERLOAD2(double, max)
__SPIRV_OVERLOAD2(double, min)
template <typename __T1, typename __T2, typename __T3>
-__DEVICE__ std::enable_if_t<
- std::numeric_limits<__T1>::is_specialized &&
- std::numeric_limits<__T2>::is_specialized &&
- std::numeric_limits<__T3>::is_specialized,
- double>
+__DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized &&
+ std::numeric_limits<__T2>::is_specialized &&
+ std::numeric_limits<__T3>::is_specialized,
+ double>
fma(__T1 __x, __T2 __y, __T3 __z) {
return ::fma((double)__x, (double)__y, (double)__z);
}
-
template <typename __T>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
- frexp(__T __x, int *__exp) {
+frexp(__T __x, int *__exp) {
return ::frexp((double)__x, __exp);
}
template <typename __T>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
- ldexp(__T __x, int __exp) {
+ldexp(__T __x, int __exp) {
return ::ldexp((double)__x, __exp);
}
template <typename __T>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
- modf(__T __x, double *__exp) {
+modf(__T __x, double *__exp) {
return ::modf((double)__x, __exp);
}
template <typename __T1, typename __T2>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized &&
- std::numeric_limits<__T2>::is_specialized,
+ std::numeric_limits<__T2>::is_specialized,
double>
- remquo(__T1 __x, __T2 __y, int *__quo) {
+remquo(__T1 __x, __T2 __y, int *__quo) {
return ::remquo((double)__x, (double)__y, __quo);
}
template <typename __T>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
- scalbln(__T __x, long int __exp) {
+scalbln(__T __x, long int __exp) {
return ::scalbln((double)__x, __exp);
}
template <typename __T>
__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
- scalbn(__T __x, int __exp) {
+scalbn(__T __x, int __exp) {
return ::scalbn((double)__x, __exp);
}
@@ -467,7 +462,7 @@ using ::ilogbf;
using ::ldexpf;
using ::lgammaf;
using ::llrintf;
-using ::llroundf;hfgh fghdggf h
+using ::llroundf;
using ::log10f;
using ::log1pf;
using ::log2f;
diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath
index e6b887ff70507..a277126304d37 100644
--- a/clang/lib/Headers/openmp_wrappers/cmath
+++ b/clang/lib/Headers/openmp_wrappers/cmath
@@ -28,7 +28,8 @@
#include <limits>
#pragma omp begin declare variant match( \
- device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any, allow_templates)})
+ device = {arch(nvptx, nvptx64)}, \
+ implementation = {extension(match_any, allow_templates)})
#define __CUDA__
#define __OPENMP_NVPTX__
@@ -136,48 +137,8 @@ __DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
#include <__clang_spirv_cmath.h>
-
#undef __OPENMP_SPIRV__
-// Define overloads otherwise which are absent
-#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
-
-__DEVICE__ float acos(float __x) { return ::acosf(__x); }
-__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
-__DEVICE__ float asin(float __x) { return ::asinf(__x); }
-__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
-__DEVICE__ float atan(float __x) { return ::atanf(__x); }
-__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
-__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
-__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
-__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
-__DEVICE__ float erf(float __x) { return ::erff(__x); }
-__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
-__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
-__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
-__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); }
-__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
-__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
-__DEVICE__ float ldexp(float __arg, int __exp) {
- return ::ldexpf(__arg, __exp);
-}
-__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
-__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
-__DEVICE__ float logb(float __x) { return ::logbf(__x); }
-__DEVICE__ float nextafter(float __x, float __y) {
- return ::nextafterf(__x, __y);
-}
-__DEVICE__ float remainder(float __x, float __y) {
- return ::remainderf(__x, __y);
-}
-__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
-__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
-__DEVICE__ float tan(float __x) { return ::tanf(__x); }
-__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
-__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
-
-#undef __DEVICE__
-
#pragma omp end declare variant
#endif // __SPIRV__
>From 45c3ccd4c5eb04589a8cc6e637e35db943983d5c Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Fri, 1 May 2026 18:09:02 -0700
Subject: [PATCH 3/4] Fix test failure
---
clang/test/Headers/openmp_device_math_isnan.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/test/Headers/openmp_device_math_isnan.cpp b/clang/test/Headers/openmp_device_math_isnan.cpp
index 3fd98813f2480..7c8dfd4672227 100644
--- a/clang/test/Headers/openmp_device_math_isnan.cpp
+++ b/clang/test/Headers/openmp_device_math_isnan.cpp
@@ -30,18 +30,18 @@ double math(float f, double d) {
// INT_RETURN: call noundef i32 @__nv_isnanf(float
// AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
// AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
- // SPIRV_INT_RETURN: call spir_func noundef i32 @_Z5isnanf(float
+ // SPIRV_INT_RETURN: call spir_func zeroext i1 @_Z13__spirv_IsNanf(float
// BOOL_RETURN: call noundef i32 @__nv_isnanf(float
- // SPIRV_BOOL_RETURN: call spir_func noundef zeroext i1 @_Z5isnanf(float
+ // SPIRV_BOOL_RETURN: call spir_func zeroext i1 @_Z13__spirv_IsNanf(float
// AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
// AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
r += std::isnan(f);
// INT_RETURN: call noundef i32 @__nv_isnand(double
- // SPIRV_INT_RETURN: call spir_func noundef i32 @_Z5isnand(double
+ // SPIRV_INT_RETURN: call spir_func zeroext i1 @_Z13__spirv_IsNand(double
// AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
// AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
// BOOL_RETURN: call noundef i32 @__nv_isnand(double
- // SPIRV_BOOL_RETURN: call spir_func noundef zeroext i1 @_Z5isnand(double
+ // SPIRV_BOOL_RETURN: call spir_func zeroext i1 @_Z13__spirv_IsNand(double
// AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
// AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
r += std::isnan(d);
>From c719a0da977f19aaacf31e635e61ad06f287c808 Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Wed, 6 May 2026 22:50:13 -0700
Subject: [PATCH 4/4] Unify cmath headers for OpenMP wrappers
---
clang/lib/Headers/__clang_spirv_cmath.h | 327 ++----------------------
clang/lib/Headers/openmp_wrappers/cmath | 88 +------
2 files changed, 22 insertions(+), 393 deletions(-)
diff --git a/clang/lib/Headers/__clang_spirv_cmath.h b/clang/lib/Headers/__clang_spirv_cmath.h
index 0e2bc899fc1c5..38565b93dc7b5 100644
--- a/clang/lib/Headers/__clang_spirv_cmath.h
+++ b/clang/lib/Headers/__clang_spirv_cmath.h
@@ -10,10 +10,10 @@
#ifndef __CLANG_SPIRV_CMATH_H__
#define __CLANG_SPIRV_CMATH_H__
-#if !defined(__SPIRV__) && !defined(__OPENMP_SPIRV__)
-#error "This file is for SPIRV OpenMP device compilation only."
+#if !defined(__OPENMP_SPIRV__) && !defined(__OPENMP_AMDGCN__) && \
+ !defined(__OPENMP_NVPTX__)
+#error "This file is for SPIRV/HIP/CUDA OpenMP device compilation only."
#endif
-
#if defined(__cplusplus)
#include <limits>
#include <type_traits>
@@ -21,18 +21,8 @@
#endif
#include <limits.h>
#include <stdint.h>
-
-#pragma push_macro("__DEVICE__")
-#ifdef __OPENMP_SPIRV__
-#if defined(__cplusplus)
#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
-#else
-#define __DEVICE__ static __attribute__((always_inline, nothrow))
-#endif
-#else
-#define __DEVICE__ static __device__ __forceinline__
-#endif
-
+#if defined(__cplusplus)
__DEVICE__ float sin(float __x) { return ::sinf(__x); }
__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
__DEVICE__ float cos(float __x) { return ::cosf(__x); }
@@ -69,11 +59,8 @@ __DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
__DEVICE__ float floor(float __x) { return ::floorf(__x); }
__DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
-__DEVICE__ float fmax(float __x, float __y) { return ::fmaxf(__x, __y); }
-__DEVICE__ float fmin(float __x, float __y) { return ::fminf(__x, __y); }
__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
-#if defined(__OPENMP_SPIRV__)
// For OpenMP we work around some old system headers that have non-conforming
// `isinf(float)` and `isnan(float)` implementations that return an `int`. We do
// this by providing two versions of these functions, differing only in the
@@ -99,7 +86,6 @@ __DEVICE__ int isnan(float __x) { return ::__isnanf(__x); }
__DEVICE__ int isnan(double __x) { return ::__isnan(__x); }
#pragma omp end declare variant
-#endif // defined(__OPENMP_SPIRV__)
__DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }
__DEVICE__ bool isinf(double __x) { return ::__isinf(__x); }
@@ -108,9 +94,7 @@ __DEVICE__ bool isfinite(double __x) { return ::__finite(__x); }
__DEVICE__ bool isnan(float __x) { return ::__isnanf(__x); }
__DEVICE__ bool isnan(double __x) { return ::__isnan(__x); }
-#if defined(__OPENMP_SPIRV__)
#pragma omp end declare variant
-#endif // defined(__OPENMP_SPIRV__)
__DEVICE__ bool isgreater(float __x, float __y) {
return __builtin_isgreater(__x, __y);
@@ -164,7 +148,13 @@ __DEVICE__ float scalbln(float __x, long int __n) {
return ::scalblnf(__x, __n);
}
__DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
-__DEVICE__ bool signbit(double __x) { return ::__signbit(__x); }
+__DEVICE__ bool signbit(double __x) {
+#if defined(__OPENMP_NVPTX__)
+ return ::__signbitd(__x);
+#else
+ return ::__signbit(__x);
+#endif
+}
__DEVICE__ float ldexp(float __arg, int __exp) {
return ::ldexpf(__arg, __exp);
}
@@ -206,295 +196,14 @@ __DEVICE__ float remainder(float __a, float __b) {
}
__DEVICE__ float scalbn(float __a, int __b) { return ::scalbnf(__a, __b); }
-#ifndef __OPENMP_SPIRV__
-#pragma push_macro("__SPIRV_OVERLOAD1")
-#pragma push_macro("__SPIRV_OVERLOAD2")
-
-// __SPIRV_OVERLOAD1 is used to resolve function calls with integer argument to
-// avoid compilation error due to ambiguity. e.g. floor(5) is resolved with
-// floor(double).
-#define __SPIRV_OVERLOAD1(__retty, __fn) \
- template <typename __T> \
- __DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, __retty> \
- __fn(__T __x) { \
- return ::__fn((double)__x); \
- }
-
-#define __SPIRV_OVERLOAD2(__retty, __fn) \
- template <typename __T1, typename __T2> \
- __DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized && \
- std::numeric_limits<__T2>::is_specialized, \
- __retty> \
- __fn(__T1 __x, __T2 __y) { \
- return __fn((double)__x, (double)__y); \
- }
-
-__SPIRV_OVERLOAD1(double, acos)
-__SPIRV_OVERLOAD1(double, acosh)
-__SPIRV_OVERLOAD1(double, asin)
-__SPIRV_OVERLOAD1(double, asinh)
-__SPIRV_OVERLOAD1(double, atan)
-__SPIRV_OVERLOAD2(double, atan2)
-__SPIRV_OVERLOAD1(double, atanh)
-__SPIRV_OVERLOAD1(double, cbrt)
-__SPIRV_OVERLOAD1(double, ceil)
-__SPIRV_OVERLOAD2(double, copysign)
-__SPIRV_OVERLOAD1(double, cos)
-__SPIRV_OVERLOAD1(double, cosh)
-__SPIRV_OVERLOAD1(double, erf)
-__SPIRV_OVERLOAD1(double, erfc)
-__SPIRV_OVERLOAD1(double, exp)
-__SPIRV_OVERLOAD1(double, exp2)
-__SPIRV_OVERLOAD1(double, expm1)
-__SPIRV_OVERLOAD1(double, fabs)
-__SPIRV_OVERLOAD2(double, fdim)
-__SPIRV_OVERLOAD1(double, floor)
-__SPIRV_OVERLOAD2(double, fmax)
-__SPIRV_OVERLOAD2(double, fmin)
-__SPIRV_OVERLOAD2(double, fmod)
-__SPIRV_OVERLOAD1(int, fpclassify)
-__SPIRV_OVERLOAD2(double, hypot)
-__SPIRV_OVERLOAD1(int, ilogb)
-__SPIRV_OVERLOAD1(bool, isfinite)
-__SPIRV_OVERLOAD2(bool, isgreater)
-__SPIRV_OVERLOAD2(bool, isgreaterequal)
-__SPIRV_OVERLOAD1(bool, isinf)
-__SPIRV_OVERLOAD2(bool, isless)
-__SPIRV_OVERLOAD2(bool, islessequal)
-__SPIRV_OVERLOAD2(bool, islessgreater)
-__SPIRV_OVERLOAD1(bool, isnan)
-__SPIRV_OVERLOAD1(bool, isnormal)
-__SPIRV_OVERLOAD2(bool, isunordered)
-__SPIRV_OVERLOAD1(double, lgamma)
-__SPIRV_OVERLOAD1(double, log)
-__SPIRV_OVERLOAD1(double, log10)
-__SPIRV_OVERLOAD1(double, log1p)
-__SPIRV_OVERLOAD1(double, log2)
-__SPIRV_OVERLOAD1(double, logb)
-__SPIRV_OVERLOAD1(long long, llrint)
-__SPIRV_OVERLOAD1(long long, llround)
-__SPIRV_OVERLOAD1(long, lrint)
-__SPIRV_OVERLOAD1(long, lround)
-__SPIRV_OVERLOAD1(double, nearbyint)
-__SPIRV_OVERLOAD2(double, nextafter)
-__SPIRV_OVERLOAD2(double, pow)
-__SPIRV_OVERLOAD2(double, remainder)
-__SPIRV_OVERLOAD1(double, rint)
-__SPIRV_OVERLOAD1(double, round)
-__SPIRV_OVERLOAD1(bool, signbit)
-__SPIRV_OVERLOAD1(double, sin)
-__SPIRV_OVERLOAD1(double, sinh)
-__SPIRV_OVERLOAD1(double, sqrt)
-__SPIRV_OVERLOAD1(double, tan)
-__SPIRV_OVERLOAD1(double, tanh)
-__SPIRV_OVERLOAD1(double, tgamma)
-__SPIRV_OVERLOAD1(double, trunc)
-
-// Overload these but don't add them to std, they are not part of cmath.
-__SPIRV_OVERLOAD2(double, max)
-__SPIRV_OVERLOAD2(double, min)
-
-template <typename __T1, typename __T2, typename __T3>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized &&
- std::numeric_limits<__T2>::is_specialized &&
- std::numeric_limits<__T3>::is_specialized,
- double>
-fma(__T1 __x, __T2 __y, __T3 __z) {
- return ::fma((double)__x, (double)__y, (double)__z);
-}
-
-template <typename __T>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
-frexp(__T __x, int *__exp) {
- return ::frexp((double)__x, __exp);
+#if defined(__OPENMP_AMDGCN__)
+__DEVICE__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) {
+ return __builtin_fmaf16(__x, __y, __z);
}
-
-template <typename __T>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
-ldexp(__T __x, int __exp) {
- return ::ldexp((double)__x, __exp);
-}
-
-template <typename __T>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
-modf(__T __x, double *__exp) {
- return ::modf((double)__x, __exp);
-}
-
-template <typename __T1, typename __T2>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T1>::is_specialized &&
- std::numeric_limits<__T2>::is_specialized,
- double>
-remquo(__T1 __x, __T2 __y, int *__quo) {
- return ::remquo((double)__x, (double)__y, __quo);
+__DEVICE__ _Float16 pow(_Float16 __base, int __iexp) {
+ return __ocml_pown_f16(__base, __iexp);
}
-
-template <typename __T>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
-scalbln(__T __x, long int __exp) {
- return ::scalbln((double)__x, __exp);
-}
-
-template <typename __T>
-__DEVICE__ std::enable_if_t<std::numeric_limits<__T>::is_integer, double>
-scalbn(__T __x, int __exp) {
- return ::scalbn((double)__x, __exp);
-}
-
-#pragma pop_macro("__SPIRV_OVERLOAD1")
-#pragma pop_macro("__SPIRV_OVERLOAD2")
-
-// Define these overloads inside the namespace our standard library uses.
-
-#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
-_LIBCPP_BEGIN_NAMESPACE_STD
-#else
-namespace std {
-#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
-_GLIBCXX_BEGIN_NAMESPACE_VERSION
-#endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
-#endif // _LIBCPP_BEGIN_NAMESPACE_STD
-
-// Pull the new overloads we defined above into namespace std.
-// using ::abs; - This may be considered for C++.
-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 ::fabs;
-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 ::isless;
-using ::islessequal;
-using ::islessgreater;
-using ::isnormal;
-using ::isunordered;
-using ::ldexp;
-using ::lgamma;
-using ::llrint;
-using ::llround;
-using ::log;
-using ::log10;
-using ::log1p;
-using ::log2;
-using ::logb;
-using ::lrint;
-using ::lround;
-using ::modf;
-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;
-
-// Well this is fun: We need to pull these symbols in for libc++, but we can't
-// pull them in with libstdc++, because its ::isinf and ::isnan are different
-// than its std::isinf and std::isnan.
-#ifndef __GLIBCXX__
-using ::isinf;
-using ::isnan;
#endif
-
-// Finally, pull the "foobarf" functions that HIP defines into std.
-using ::acosf;
-using ::acoshf;
-using ::asinf;
-using ::asinhf;
-using ::atan2f;
-using ::atanf;
-using ::atanhf;
-using ::cbrtf;
-using ::ceilf;
-using ::copysignf;
-using ::cosf;
-using ::coshf;
-using ::erfcf;
-using ::erff;
-using ::exp2f;
-using ::expf;
-using ::expm1f;
-using ::fabsf;
-using ::fdimf;
-using ::floorf;
-using ::fmaf;
-using ::fmaxf;
-using ::fminf;
-using ::fmodf;
-using ::frexpf;
-using ::hypotf;
-using ::ilogbf;
-using ::ldexpf;
-using ::lgammaf;
-using ::llrintf;
-using ::llroundf;
-using ::log10f;
-using ::log1pf;
-using ::log2f;
-using ::logbf;
-using ::logf;
-using ::lrintf;
-using ::lroundf;
-using ::modff;
-using ::nearbyintf;
-using ::nextafterf;
-using ::powf;
-using ::remainderf;
-using ::remquof;
-using ::rintf;
-using ::roundf;
-using ::scalblnf;
-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 // _GLIBCXX_BEGIN_NAMESPACE_VERSION
-} // namespace std
-#endif // _LIBCPP_END_NAMESPACE_STD
-#endif // ifndef __OPENMP_SPIRV__
+#endif
+#undef __DEVICE__
#endif // __CLANG_SPIRV_CMATH_H__
\ No newline at end of file
diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath
index a277126304d37..393e28c5fc863 100644
--- a/clang/lib/Headers/openmp_wrappers/cmath
+++ b/clang/lib/Headers/openmp_wrappers/cmath
@@ -26,107 +26,27 @@
// We need limits because __clang_cuda_cmath.h below uses `std::numeric_limit`.
#include <limits>
-
+#ifdef __NVPTX__
#pragma omp begin declare variant match( \
device = {arch(nvptx, nvptx64)}, \
implementation = {extension(match_any, allow_templates)})
#define __CUDA__
#define __OPENMP_NVPTX__
-#include <__clang_cuda_cmath.h>
+#include <__clang_spirv_cmath.h>
#undef __OPENMP_NVPTX__
#undef __CUDA__
-// Overloads not provided by the CUDA wrappers but by the CUDA system headers.
-// Since we do not include the latter we define them ourselves.
-#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
-
-__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
-__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
-__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
-__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
-__DEVICE__ float erf(float __x) { return ::erff(__x); }
-__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
-__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
-__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
-__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); }
-__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
-__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
-__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
-__DEVICE__ long long int llrint(float __x) { return ::llrintf(__x); }
-__DEVICE__ long long int llround(float __x) { return ::llroundf(__x); }
-__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
-__DEVICE__ float log2(float __x) { return ::log2f(__x); }
-__DEVICE__ float logb(float __x) { return ::logbf(__x); }
-__DEVICE__ long int lrint(float __x) { return ::lrintf(__x); }
-__DEVICE__ long int lround(float __x) { return ::lroundf(__x); }
-__DEVICE__ float nextafter(float __x, float __y) {
- return ::nextafterf(__x, __y);
-}
-__DEVICE__ float remainder(float __x, float __y) {
- return ::remainderf(__x, __y);
-}
-__DEVICE__ float scalbln(float __x, long int __y) {
- return ::scalblnf(__x, __y);
-}
-__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
-__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
-
-#undef __DEVICE__
-
#pragma omp end declare variant
+#endif // __NVPTX__
#ifdef __AMDGCN__
#pragma omp begin declare variant match(device = {arch(amdgcn)})
-#pragma push_macro("__constant__")
-#define __constant__ __attribute__((constant))
#define __OPENMP_AMDGCN__
-
-#include <__clang_hip_cmath.h>
-
-#pragma pop_macro("__constant__")
+#include <__clang_spirv_cmath.h>
#undef __OPENMP_AMDGCN__
-// Define overloads otherwise which are absent
-#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
-
-__DEVICE__ float acos(float __x) { return ::acosf(__x); }
-__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
-__DEVICE__ float asin(float __x) { return ::asinf(__x); }
-__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
-__DEVICE__ float atan(float __x) { return ::atanf(__x); }
-__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
-__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
-__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
-__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
-__DEVICE__ float erf(float __x) { return ::erff(__x); }
-__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
-__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
-__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
-__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); }
-__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
-__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
-__DEVICE__ float ldexp(float __arg, int __exp) {
- return ::ldexpf(__arg, __exp);
-}
-__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
-__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
-__DEVICE__ float logb(float __x) { return ::logbf(__x); }
-__DEVICE__ float nextafter(float __x, float __y) {
- return ::nextafterf(__x, __y);
-}
-__DEVICE__ float remainder(float __x, float __y) {
- return ::remainderf(__x, __y);
-}
-__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
-__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
-__DEVICE__ float tan(float __x) { return ::tanf(__x); }
-__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
-__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
-
-#undef __DEVICE__
-
#pragma omp end declare variant
#endif // __AMDGCN__
More information about the cfe-commits
mailing list