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" <[email protected]> Date: Fri, 1 May 2026 10:26:07 -0700 Subject: [PATCH 1/5] 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" <[email protected]> Date: Fri, 1 May 2026 14:51:59 -0700 Subject: [PATCH 2/5] 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" <[email protected]> Date: Fri, 1 May 2026 18:09:02 -0700 Subject: [PATCH 3/5] 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" <[email protected]> Date: Wed, 6 May 2026 22:50:13 -0700 Subject: [PATCH 4/5] 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__ >From 5c8ab80268e5fba4fd7e6eec27708d90cf1e856b Mon Sep 17 00:00:00 2001 From: "Fine, Gregory" <[email protected]> Date: Thu, 7 May 2026 10:36:59 -0700 Subject: [PATCH 5/5] Rename files --- clang/lib/Headers/CMakeLists.txt | 2 +- .../lib/Headers/{__clang_spirv_cmath.h => __clang_cmath.h} | 6 +++--- clang/lib/Headers/openmp_wrappers/cmath | 6 +++--- 3 files changed, 7 insertions(+), 7 deletions(-) rename clang/lib/Headers/{__clang_spirv_cmath.h => __clang_cmath.h} (98%) diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index 4252f8c4685b6..e57e64a524ae2 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -42,6 +42,7 @@ set(core_files tgmath.h unwind.h varargs.h + __clang_cmath.h ) set(arm_common_files @@ -146,7 +147,6 @@ 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_cmath.h similarity index 98% rename from clang/lib/Headers/__clang_spirv_cmath.h rename to clang/lib/Headers/__clang_cmath.h index 38565b93dc7b5..28c6e85707dee 100644 --- a/clang/lib/Headers/__clang_spirv_cmath.h +++ b/clang/lib/Headers/__clang_cmath.h @@ -7,8 +7,8 @@ *===-----------------------------------------------------------------------=== */ -#ifndef __CLANG_SPIRV_CMATH_H__ -#define __CLANG_SPIRV_CMATH_H__ +#ifndef __CLANG_CMATH_H__ +#define __CLANG_CMATH_H__ #if !defined(__OPENMP_SPIRV__) && !defined(__OPENMP_AMDGCN__) && \ !defined(__OPENMP_NVPTX__) @@ -206,4 +206,4 @@ __DEVICE__ _Float16 pow(_Float16 __base, int __iexp) { #endif #endif #undef __DEVICE__ -#endif // __CLANG_SPIRV_CMATH_H__ \ No newline at end of file +#endif // __CLANG_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 393e28c5fc863..d49e51b06adbb 100644 --- a/clang/lib/Headers/openmp_wrappers/cmath +++ b/clang/lib/Headers/openmp_wrappers/cmath @@ -33,7 +33,7 @@ #define __CUDA__ #define __OPENMP_NVPTX__ -#include <__clang_spirv_cmath.h> +#include <__clang_cmath.h> #undef __OPENMP_NVPTX__ #undef __CUDA__ @@ -44,7 +44,7 @@ #pragma omp begin declare variant match(device = {arch(amdgcn)}) #define __OPENMP_AMDGCN__ -#include <__clang_spirv_cmath.h> +#include <__clang_cmath.h> #undef __OPENMP_AMDGCN__ #pragma omp end declare variant @@ -55,7 +55,7 @@ #define __OPENMP_SPIRV__ -#include <__clang_spirv_cmath.h> +#include <__clang_cmath.h> #undef __OPENMP_SPIRV__ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
