From 0c9aa18d3816f756f1694e0ebd6ccf232062ccf2 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:09:37 -0700 Subject: [PATCH 1/6] Revert "[SYCL] Make SYCL math functions overloads instead of templates (#9753)" This reverts commit 826d868af1c33b9630881c615737343c01746293. --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 8 +- sycl/include/sycl/builtins.hpp | 1208 +++++++---------- .../ext/intel/esimd/detail/math_intrin.hpp | 18 +- .../ext/oneapi/experimental/sycl_complex.hpp | 2 +- sycl/test-e2e/Basic/half_builtins.cpp | 11 +- .../Basic/sycl_2020_images/common.hpp | 2 +- .../discard_events_check_images.cpp | 3 +- sycl/test-e2e/ESIMD/kmeans/kmeans.cpp | 8 +- 8 files changed, 550 insertions(+), 710 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 1185dd35c4a91..0c29357f7672f 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -53,10 +53,10 @@ static const char *LegalSYCLFunctions[] = { "^sycl::_V1::sub_group<.+>::.+", "^sycl::_V1::range<.+>::.+", "^sycl::_V1::kernel_handler::.+", - "^sycl::_V1::cos", - "^sycl::_V1::sin", - "^sycl::_V1::log", - "^sycl::_V1::exp", + "^sycl::_V1::cos<.+>", + "^sycl::_V1::sin<.+>", + "^sycl::_V1::log<.+>", + "^sycl::_V1::exp<.+>", "^sycl::_V1::bit_cast<.+>", "^sycl::_V1::operator.+<.+>", "^sycl::_V1::ext::oneapi::sub_group::.+", diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index f4391fd44a711..71d24e532a3c4 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -486,9 +486,12 @@ namespace __sycl_std = __host_std; // TODO: Replace with overloads. #ifdef __FAST_MATH__ +#define __FAST_MATH_GENFLOAT(T) \ + (detail::is_svgenfloatd::value || detail::is_svgenfloath::value) #define __FAST_MATH_SGENFLOAT(T) \ (std::is_same_v || std::is_same_v) #else +#define __FAST_MATH_GENFLOAT(T) (detail::is_svgenfloat::value) #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat::value) #endif @@ -514,7 +517,6 @@ namespace __sycl_std = __host_std; /* ----------------- 4.13.3 Math functions. ---------------------------------*/ -// TODO: Replace with overloads. // These macros for marray math function implementations use vectorizations of // size two as a simple general optimization. A more complex implementation // using larger vectorizations for large marray sizes is possible; however more @@ -773,261 +775,207 @@ __SYCL_MATH_FUNCTION_3_OVERLOAD(mad) __SYCL_MATH_FUNCTION_3_OVERLOAD(mix) #undef __SYCL_MATH_FUNCTION_3_OVERLOAD -// genfloat acos (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE acos(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_acos(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF + // svgenfloat acos (svgenfloat x) + template + std::enable_if_t::value, T> acos(T x) __NOEXC { + return __sycl_std::__invoke_acos(x); +} -// genfloat acosh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE acosh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_acosh(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat acosh (svgenfloat x) +template +std::enable_if_t::value, T> acosh(T x) __NOEXC { + return __sycl_std::__invoke_acosh(x); +} -// genfloat acospi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE acospi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_acospi(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat acospi (svgenfloat x) +template +std::enable_if_t::value, T> acospi(T x) __NOEXC { + return __sycl_std::__invoke_acospi(x); +} -// genfloat asin (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE asin(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_asin(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat asin (svgenfloat x) +template +std::enable_if_t::value, T> asin(T x) __NOEXC { + return __sycl_std::__invoke_asin(x); +} -// genfloat asinh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE asinh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_asinh(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat asinh (svgenfloat x) +template +std::enable_if_t::value, T> asinh(T x) __NOEXC { + return __sycl_std::__invoke_asinh(x); +} -// genfloat asinpi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE asinpi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_asinpi(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat asinpi (svgenfloat x) +template +std::enable_if_t::value, T> asinpi(T x) __NOEXC { + return __sycl_std::__invoke_asinpi(x); +} -// genfloat atan (genfloat y_over_x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE atan(TYPE y_over_x) __NOEXC { \ - return __sycl_std::__invoke_atan(y_over_x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat atan (svgenfloat y_over_x) +template +std::enable_if_t::value, T> atan(T y_over_x) __NOEXC { + return __sycl_std::__invoke_atan(y_over_x); +} -// genfloat atan2 (genfloat y, genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE atan2(TYPE y, TYPE x) __NOEXC { \ - return __sycl_std::__invoke_atan2(y, x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat atan2 (svgenfloat y, svgenfloat x) +template +std::enable_if_t::value, T> atan2(T y, T x) __NOEXC { + return __sycl_std::__invoke_atan2(y, x); +} -// genfloat atanh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE atanh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_atanh(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat atanh (svgenfloat x) +template +std::enable_if_t::value, T> atanh(T x) __NOEXC { + return __sycl_std::__invoke_atanh(x); +} -// genfloat atanpi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE atanpi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_atanpi(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat atanpi (svgenfloat x) +template +std::enable_if_t::value, T> atanpi(T x) __NOEXC { + return __sycl_std::__invoke_atanpi(x); +} -// genfloat atan2pi (genfloat y, genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE atan2pi(TYPE y, TYPE x) __NOEXC { \ - return __sycl_std::__invoke_atan2pi(y, x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat atan2pi (svgenfloat y, svgenfloat x) +template +std::enable_if_t::value, T> atan2pi(T y, T x) __NOEXC { + return __sycl_std::__invoke_atan2pi(y, x); +} -// genfloat cbrt (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cbrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_cbrt(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat cbrt (svgenfloat x) +template +std::enable_if_t::value, T> cbrt(T x) __NOEXC { + return __sycl_std::__invoke_cbrt(x); +} -// genfloat ceil (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE ceil(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_ceil(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat ceil (svgenfloat x) +template +std::enable_if_t::value, T> ceil(T x) __NOEXC { + return __sycl_std::__invoke_ceil(x); +} -// genfloat copysign (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE copysign(TYPE y, TYPE x) __NOEXC { \ - return __sycl_std::__invoke_copysign(y, x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat copysign (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> copysign(T x, + T y) __NOEXC { + return __sycl_std::__invoke_copysign(x, y); +} -// genfloat cos (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cos(TYPE x) __NOEXC { return __sycl_std::__invoke_cos(x); } - __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat cos (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> cos(T x) __NOEXC { + return __sycl_std::__invoke_cos(x); +} -// genfloat cosh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cosh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_cosh(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat cosh (svgenfloat x) +template +std::enable_if_t::value, T> cosh(T x) __NOEXC { + return __sycl_std::__invoke_cosh(x); +} -// genfloat cospi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cospi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_cospi(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat cospi (svgenfloat x) +template +std::enable_if_t::value, T> cospi(T x) __NOEXC { + return __sycl_std::__invoke_cospi(x); +} -// genfloat erfc (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE erfc(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_erfc(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat erfc (svgenfloat x) +template +std::enable_if_t::value, T> erfc(T x) __NOEXC { + return __sycl_std::__invoke_erfc(x); +} -// genfloat erf (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE erf(TYPE x) __NOEXC { return __sycl_std::__invoke_erf(x); } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat erf (svgenfloat x) +template +std::enable_if_t::value, T> erf(T x) __NOEXC { + return __sycl_std::__invoke_erf(x); +} -// genfloat exp (genfloat x ) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp(TYPE x) __NOEXC { return __sycl_std::__invoke_exp(x); } - __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat exp (svgenfloat x ) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp(T x) __NOEXC { + return __sycl_std::__invoke_exp(x); +} -// genfloat exp2 (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_exp2(x); \ - } - __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat exp2 (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp2(T x) __NOEXC { + return __sycl_std::__invoke_exp2(x); +} -// genfloat exp10 (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_exp10(x); \ - } - __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat exp10 (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp10(T x) __NOEXC { + return __sycl_std::__invoke_exp10(x); +} -// genfloat expm1 (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE expm1(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_expm1(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat expm1 (svgenfloat x) +template +std::enable_if_t::value, T> expm1(T x) __NOEXC { + return __sycl_std::__invoke_expm1(x); +} -// genfloat fabs (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fabs(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_fabs(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fabs (svgenfloat x) +template +std::enable_if_t::value, T> fabs(T x) __NOEXC { + return __sycl_std::__invoke_fabs(x); +} -// genfloat fdim (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fdim(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_fdim(x, y); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fdim (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> fdim(T x, T y) __NOEXC { + return __sycl_std::__invoke_fdim(x, y); +} -// genfloat floor (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE floor(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_floor(x); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat floor (svgenfloat x) +template +std::enable_if_t::value, T> floor(T x) __NOEXC { + return __sycl_std::__invoke_floor(x); +} -// genfloat fma (genfloat a, genfloat b, genfloat c) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fma(TYPE a, TYPE b, TYPE c) __NOEXC { \ - return __sycl_std::__invoke_fma(a, b, c); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fma (svgenfloat a, svgenfloat b, svgenfloat c) +template +std::enable_if_t::value, T> fma(T a, T b, + T c) __NOEXC { + return __sycl_std::__invoke_fma(a, b, c); +} -// genfloat fmax (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fmax(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_fmax(x, y); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fmax (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> fmax(T x, T y) __NOEXC { + return __sycl_std::__invoke_fmax(x, y); +} -// genfloat fmax (genfloat x, sgenfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fmax(TYPE x, TYPE::element_type y) __NOEXC { \ - return __sycl_std::__invoke_fmax(x, TYPE(y)); \ - } - __SYCL_DEF_BUILTIN_VGENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fmax (svgenfloat x, sgenfloat y) +template +std::enable_if_t::value, T> +fmax(T x, typename T::element_type y) __NOEXC { + return __sycl_std::__invoke_fmax(x, T(y)); +} -// genfloat fmin (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fmin(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_fmin(x, y); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fmin (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> fmin(T x, T y) __NOEXC { + return __sycl_std::__invoke_fmin(x, y); +} -// genfloat fmin (genfloat x, sgenfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fmin(TYPE x, TYPE::element_type y) __NOEXC { \ - return __sycl_std::__invoke_fmin(x, TYPE(y)); \ - } - __SYCL_DEF_BUILTIN_VGENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fmin (svgenfloat x, sgenfloat y) +template +std::enable_if_t::value, T> +fmin(T x, typename T::element_type y) __NOEXC { + return __sycl_std::__invoke_fmin(x, T(y)); +} -// genfloat fmod (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fmod(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_fmod(x, y); \ - } - __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat fmod (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> fmod(T x, T y) __NOEXC { + return __sycl_std::__invoke_fmod(x, y); +} - // svgenfloat fract (svgenfloat x, genfloatptr iptr) - template - std::enable_if_t::value && - detail::is_genfloatptr::value, - T> fract(T x, T2 iptr) __NOEXC { +// svgenfloat fract (svgenfloat x, genfloatptr iptr) +template +std::enable_if_t< + detail::is_svgenfloat::value && detail::is_genfloatptr::value, T> +fract(T x, T2 iptr) __NOEXC { detail::check_vector_size(); return __sycl_std::__invoke_fract(x, iptr); } @@ -1041,40 +989,32 @@ frexp(T x, T2 exp) __NOEXC { return __sycl_std::__invoke_frexp(x, exp); } -// genfloat hypot (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE hypot(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_hypot(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat hypot (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> hypot(T x, T y) __NOEXC { + return __sycl_std::__invoke_hypot(x, y); +} -// genint ilogb (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::change_base_type_t ilogb(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_ilogb>( \ - x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// genint ilogb (svgenfloat x) +template ::value, T>> +detail::change_base_type_t ilogb(T x) __NOEXC { + return __sycl_std::__invoke_ilogb>(x); +} // float ldexp (float x, int k) // double ldexp (double x, int k) // half ldexp (half x, int k) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE ldexp(TYPE x, int k) __NOEXC { \ - return __sycl_std::__invoke_ldexp(x, k); \ - } -__SYCL_DEF_BUILTIN_SGENFLOAT -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> ldexp(T x, int k) __NOEXC { + return __sycl_std::__invoke_ldexp(x, k); +} // vgenfloat ldexp (vgenfloat x, int k) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE ldexp(TYPE x, int k) __NOEXC { \ - return __sycl_std::__invoke_ldexp(x, vec(k)); \ - } -__SYCL_DEF_BUILTIN_VGENFLOAT -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> ldexp(T x, int k) __NOEXC { + return __sycl_std::__invoke_ldexp(x, vec(k)); +} // vgenfloat ldexp (vgenfloat x, genint k) template @@ -1085,13 +1025,11 @@ ldexp(T x, T2 k) __NOEXC { return __sycl_std::__invoke_ldexp(x, k); } -// genfloat lgamma (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE lgamma(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_lgamma(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat lgamma (svgenfloat x) +template +std::enable_if_t::value, T> lgamma(T x) __NOEXC { + return __sycl_std::__invoke_lgamma(x); +} // svgenfloat lgamma_r (svgenfloat x, genintptr signp) template @@ -1102,67 +1040,54 @@ lgamma_r(T x, T2 signp) __NOEXC { return __sycl_std::__invoke_lgamma_r(x, signp); } -// genfloat log (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log(TYPE x) __NOEXC { return __sycl_std::__invoke_log(x); } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat log (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log(T x) __NOEXC { + return __sycl_std::__invoke_log(x); +} -// genfloat log2 (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_log2(x); \ - } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat log2 (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log2(T x) __NOEXC { + return __sycl_std::__invoke_log2(x); +} -// genfloat log10 (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_log10(x); \ - } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat log10 (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log10(T x) __NOEXC { + return __sycl_std::__invoke_log10(x); +} -// genfloat log1p (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log1p(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_log1p(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat log1p (svgenfloat x) +template +std::enable_if_t::value, T> log1p(T x) __NOEXC { + return __sycl_std::__invoke_log1p(x); +} -// genfloat logb (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE logb(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_logb(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat logb (svgenfloat x) +template +std::enable_if_t::value, T> logb(T x) __NOEXC { + return __sycl_std::__invoke_logb(x); +} -// genfloat mad (genfloat a, genfloat b, genfloat c) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE mad(TYPE a, TYPE b, TYPE c) __NOEXC { \ - return __sycl_std::__invoke_mad(a, b, c); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat mad (svgenfloat a, svgenfloat b, svgenfloat c) +template +std::enable_if_t::value, T> mad(T a, T b, + T c) __NOEXC { + return __sycl_std::__invoke_mad(a, b, c); +} -// genfloat maxmag (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE maxmag(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_maxmag(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat maxmag (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> maxmag(T x, T y) __NOEXC { + return __sycl_std::__invoke_maxmag(x, y); +} -// genfloat minmag (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE minmag(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_minmag(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat minmag (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> minmag(T x, T y) __NOEXC { + return __sycl_std::__invoke_minmag(x, y); +} // svgenfloat modf (svgenfloat x, genfloatptr iptr) template @@ -1180,21 +1105,18 @@ detail::nan_return_t nan(T nancode) __NOEXC { detail::convert_data_type>()(nancode)); } -// genfloat nextafter (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE nextafter(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_nextafter(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat nextafter (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> nextafter(T x, + T y) __NOEXC { + return __sycl_std::__invoke_nextafter(x, y); +} -// genfloat pow (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE pow(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_pow(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat pow (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> pow(T x, T y) __NOEXC { + return __sycl_std::__invoke_pow(x, y); +} // svgenfloat pown (svgenfloat x, genint y) template @@ -1205,21 +1127,18 @@ pown(T x, T2 y) __NOEXC { return __sycl_std::__invoke_pown(x, y); } -// genfloat powr (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE powr(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_powr(x, y); \ - } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat powr (svgenfloat x, svgenfloat y) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> powr(T x, T y) __NOEXC { + return __sycl_std::__invoke_powr(x, y); +} -// genfloat remainder (genfloat x, genfloat y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE remainder(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_remainder(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat remainder (svgenfloat x, svgenfloat y) +template +std::enable_if_t::value, T> remainder(T x, + T y) __NOEXC { + return __sycl_std::__invoke_remainder(x, y); +} // svgenfloat remquo (svgenfloat x, svgenfloat y, genintptr quo) template @@ -1230,13 +1149,11 @@ remquo(T x, T y, T2 quo) __NOEXC { return __sycl_std::__invoke_remquo(x, y, quo); } -// genfloat rint (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE rint(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_rint(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat rint (svgenfloat x) +template +std::enable_if_t::value, T> rint(T x) __NOEXC { + return __sycl_std::__invoke_rint(x); +} // svgenfloat rootn (svgenfloat x, genint y) template @@ -1247,27 +1164,23 @@ rootn(T x, T2 y) __NOEXC { return __sycl_std::__invoke_rootn(x, y); } -// genfloat round (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE round(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_round(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF - -// genfloat rsqrt (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE rsqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_rsqrt(x); \ - } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat round (svgenfloat x) +template +std::enable_if_t::value, T> round(T x) __NOEXC { + return __sycl_std::__invoke_round(x); +} -// genfloat sin (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sin(TYPE x) __NOEXC { return __sycl_std::__invoke_sin(x); } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat rsqrt (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> rsqrt(T x) __NOEXC { + return __sycl_std::__invoke_rsqrt(x); +} + +// svgenfloat sin (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sin(T x) __NOEXC { + return __sycl_std::__invoke_sin(x); +} // svgenfloat sincos (svgenfloat x, genfloatptr cosval) template @@ -1278,67 +1191,53 @@ sincos(T x, T2 cosval) __NOEXC { return __sycl_std::__invoke_sincos(x, cosval); } -// genfloat sinh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sinh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_sinh(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat sinh (svgenfloat x) +template +std::enable_if_t::value, T> sinh(T x) __NOEXC { + return __sycl_std::__invoke_sinh(x); +} -// genfloat sinpi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sinpi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_sinpi(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat sinpi (svgenfloat x) +template +std::enable_if_t::value, T> sinpi(T x) __NOEXC { + return __sycl_std::__invoke_sinpi(x); +} -// genfloat sqrt (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_sqrt(x); \ - } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat sqrt (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sqrt(T x) __NOEXC { + return __sycl_std::__invoke_sqrt(x); +} -// genfloat tan (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tan(TYPE x) __NOEXC { return __sycl_std::__invoke_tan(x); } -__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat tan (svgenfloat x) +template +std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> tan(T x) __NOEXC { + return __sycl_std::__invoke_tan(x); +} -// genfloat tanh (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tanh(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_tanh(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat tanh (svgenfloat x) +template +std::enable_if_t::value, T> tanh(T x) __NOEXC { + return __sycl_std::__invoke_tanh(x); +} -// genfloat tanpi (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tanpi(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_tanpi(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat tanpi (svgenfloat x) +template +std::enable_if_t::value, T> tanpi(T x) __NOEXC { + return __sycl_std::__invoke_tanpi(x); +} -// genfloat tgamma (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tgamma(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_tgamma(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat tgamma (svgenfloat x) +template +std::enable_if_t::value, T> tgamma(T x) __NOEXC { + return __sycl_std::__invoke_tgamma(x); +} -// genfloat trunc (genfloat x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE trunc(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_trunc(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +// svgenfloat trunc (svgenfloat x) +template +std::enable_if_t::value, T> trunc(T x) __NOEXC { + return __sycl_std::__invoke_trunc(x); +} // other marray math functions @@ -2752,117 +2651,89 @@ __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(powr) #undef __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD -// genfloatf cos (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cos(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_cos(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf cos (svgenfloatf x) +template +std::enable_if_t::value, T> cos(T x) __NOEXC { + return __sycl_std::__invoke_native_cos(x); +} -// genfloatf divide (genfloatf x, genfloatf y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE divide(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_native_divide(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf divide (svgenfloatf x, svgenfloatf y) +template +std::enable_if_t::value, T> divide(T x, T y) __NOEXC { + return __sycl_std::__invoke_native_divide(x, y); +} -// genfloatf exp (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_exp(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp (svgenfloatf x) +template +std::enable_if_t::value, T> exp(T x) __NOEXC { + return __sycl_std::__invoke_native_exp(x); +} -// genfloatf exp2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_exp2(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp2 (svgenfloatf x) +template +std::enable_if_t::value, T> exp2(T x) __NOEXC { + return __sycl_std::__invoke_native_exp2(x); +} -// genfloatf exp10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_exp10(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp10 (svgenfloatf x) +template +std::enable_if_t::value, T> exp10(T x) __NOEXC { + return __sycl_std::__invoke_native_exp10(x); +} -// genfloatf log (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_log(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log (svgenfloatf x) +template +std::enable_if_t::value, T> log(T x) __NOEXC { + return __sycl_std::__invoke_native_log(x); +} -// genfloatf log2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_log2(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log2 (svgenfloatf x) +template +std::enable_if_t::value, T> log2(T x) __NOEXC { + return __sycl_std::__invoke_native_log2(x); +} -// genfloatf log10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_log10(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log10 (svgenfloatf x) +template +std::enable_if_t::value, T> log10(T x) __NOEXC { + return __sycl_std::__invoke_native_log10(x); +} -// genfloatf powr (genfloatf x, genfloatf y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE powr(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_native_powr(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf powr (svgenfloatf x, svgenfloatf y) +template +std::enable_if_t::value, T> powr(T x, T y) __NOEXC { + return __sycl_std::__invoke_native_powr(x, y); +} -// genfloatf recip (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE recip(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_recip(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf recip (svgenfloatf x) +template +std::enable_if_t::value, T> recip(T x) __NOEXC { + return __sycl_std::__invoke_native_recip(x); +} -// genfloatf rsqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE rsqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_rsqrt(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf rsqrt (svgenfloatf x) +template +std::enable_if_t::value, T> rsqrt(T x) __NOEXC { + return __sycl_std::__invoke_native_rsqrt(x); +} -// genfloatf sin (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sin(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_sin(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sin (svgenfloatf x) +template +std::enable_if_t::value, T> sin(T x) __NOEXC { + return __sycl_std::__invoke_native_sin(x); +} -// genfloatf sqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_sqrt(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sqrt (svgenfloatf x) +template +std::enable_if_t::value, T> sqrt(T x) __NOEXC { + return __sycl_std::__invoke_native_sqrt(x); +} -// genfloatf tan (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tan(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_native_tan(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf tan (svgenfloatf x) +template +std::enable_if_t::value, T> tan(T x) __NOEXC { + return __sycl_std::__invoke_native_tan(x); +} } // namespace native namespace half_precision { @@ -2920,117 +2791,89 @@ __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(powr) #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD -// genfloatf cos (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cos(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_cos(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf cos (svgenfloatf x) +template +std::enable_if_t::value, T> cos(T x) __NOEXC { + return __sycl_std::__invoke_half_cos(x); +} -// genfloatf divide (genfloatf x, genfloatf y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE divide(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_half_divide(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf divide (svgenfloatf x, svgenfloatf y) +template +std::enable_if_t::value, T> divide(T x, T y) __NOEXC { + return __sycl_std::__invoke_half_divide(x, y); +} -// genfloatf exp (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_exp(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp (svgenfloatf x) +template +std::enable_if_t::value, T> exp(T x) __NOEXC { + return __sycl_std::__invoke_half_exp(x); +} -// genfloatf exp2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_exp2(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp2 (svgenfloatf x) +template +std::enable_if_t::value, T> exp2(T x) __NOEXC { + return __sycl_std::__invoke_half_exp2(x); +} -// genfloatf exp10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_exp10(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp10 (svgenfloatf x) +template +std::enable_if_t::value, T> exp10(T x) __NOEXC { + return __sycl_std::__invoke_half_exp10(x); +} -// genfloatf log (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_log(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log (svgenfloatf x) +template +std::enable_if_t::value, T> log(T x) __NOEXC { + return __sycl_std::__invoke_half_log(x); +} -// genfloatf log2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log2(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_log2(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log2 (svgenfloatf x) +template +std::enable_if_t::value, T> log2(T x) __NOEXC { + return __sycl_std::__invoke_half_log2(x); +} -// genfloatf log10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log10(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_log10(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log10 (svgenfloatf x) +template +std::enable_if_t::value, T> log10(T x) __NOEXC { + return __sycl_std::__invoke_half_log10(x); +} -// genfloatf powr (genfloatf x, genfloatf y) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE powr(TYPE x, TYPE y) __NOEXC { \ - return __sycl_std::__invoke_half_powr(x, y); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf powr (svgenfloatf x, svgenfloatf y) +template +std::enable_if_t::value, T> powr(T x, T y) __NOEXC { + return __sycl_std::__invoke_half_powr(x, y); +} -// genfloatf recip (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE recip(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_recip(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf recip (svgenfloatf x) +template +std::enable_if_t::value, T> recip(T x) __NOEXC { + return __sycl_std::__invoke_half_recip(x); +} -// genfloatf rsqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE rsqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_rsqrt(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf rsqrt (svgenfloatf x) +template +std::enable_if_t::value, T> rsqrt(T x) __NOEXC { + return __sycl_std::__invoke_half_rsqrt(x); +} -// genfloatf sin (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sin(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_sin(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sin (svgenfloatf x) +template +std::enable_if_t::value, T> sin(T x) __NOEXC { + return __sycl_std::__invoke_half_sin(x); +} -// genfloatf sqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sqrt(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_sqrt(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sqrt (svgenfloatf x) +template +std::enable_if_t::value, T> sqrt(T x) __NOEXC { + return __sycl_std::__invoke_half_sqrt(x); +} -// genfloatf tan (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tan(TYPE x) __NOEXC { \ - return __sycl_std::__invoke_half_tan(x); \ - } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf tan (svgenfloatf x) +template +std::enable_if_t::value, T> tan(T x) __NOEXC { + return __sycl_std::__invoke_half_tan(x); +} } // namespace half_precision @@ -3058,85 +2901,84 @@ __SYCL_MATH_FUNCTION_OVERLOAD_FM(sqrt) __SYCL_MATH_FUNCTION_OVERLOAD_FM(rsqrt) #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM -// genfloatf cos (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cos(TYPE x) __NOEXC { return native::cos(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t, marray> + powr(marray x, marray y) __NOEXC { + return native::powr(x, y); +} -// genfloatf exp (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp(TYPE x) __NOEXC { return native::exp(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf cos (svgenfloatf x) +template +std::enable_if_t::value, T> cos(T x) __NOEXC { + return native::cos(x); +} -// genfloatf exp2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp2(TYPE x) __NOEXC { return native::exp2(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp (svgenfloatf x) +template +std::enable_if_t::value, T> exp(T x) __NOEXC { + return native::exp(x); +} -// genfloatf exp10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE exp10(TYPE x) __NOEXC { return native::exp10(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp2 (svgenfloatf x) +template +std::enable_if_t::value, T> exp2(T x) __NOEXC { + return native::exp2(x); +} -// genfloatf log(genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log(TYPE x) __NOEXC { return native::log(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf exp10 (svgenfloatf x) +template +std::enable_if_t::value, T> exp10(T x) __NOEXC { + return native::exp10(x); +} -// genfloatf log2 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log2(TYPE x) __NOEXC { return native::log2(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log(svgenfloatf x) +template +std::enable_if_t::value, T> log(T x) __NOEXC { + return native::log(x); +} -// genfloatf log10 (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE log10(TYPE x) __NOEXC { return native::log10(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf log2 (svgenfloatf x) +template +std::enable_if_t::value, T> log2(T x) __NOEXC { + return native::log2(x); +} -// genfloatf powr (genfloatf x, genfloatf y) -// TODO: remove when __SYCL_DEF_BUILTIN_MARRAY is defined -template -inline __SYCL_ALWAYS_INLINE - std::enable_if_t, marray> - powr(marray x, marray y) __NOEXC { - return native::powr(x, y); +// svgenfloatf log10 (svgenfloatf x) +template +std::enable_if_t::value, T> log10(T x) __NOEXC { + return native::log10(x); } -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE powr(TYPE x, TYPE y) __NOEXC { return native::powr(x, y); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf powr (svgenfloatf x) +template +std::enable_if_t::value, T> powr(T x, T y) __NOEXC { + return native::powr(x, y); +} -// genfloatf rsqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE rsqrt(TYPE x) __NOEXC { return native::rsqrt(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf rsqrt (svgenfloatf x) +template +std::enable_if_t::value, T> rsqrt(T x) __NOEXC { + return native::rsqrt(x); +} -// genfloatf sin (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sin(TYPE x) __NOEXC { return native::sin(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sin (svgenfloatf x) +template +std::enable_if_t::value, T> sin(T x) __NOEXC { + return native::sin(x); +} -// genfloatf sqrt (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE sqrt(TYPE x) __NOEXC { return native::sqrt(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf sqrt (svgenfloatf x) +template +std::enable_if_t::value, T> sqrt(T x) __NOEXC { + return native::sqrt(x); +} -// genfloatf tan (genfloatf x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE tan(TYPE x) __NOEXC { return native::tan(x); } -__SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_BUILTIN_DEF +// svgenfloatf tan (svgenfloatf x) +template +std::enable_if_t::value, T> tan(T x) __NOEXC { + return native::tan(x); +} #endif // __FAST_MATH__ diff --git a/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp b/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp index db0af8d212c1a..602baa6634037 100644 --- a/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp +++ b/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp @@ -572,11 +572,11 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, SZ) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(inv, 1.f / src_cpp[i]) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(log, logf(src_cpp[i]) / logf(2.f)) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(exp, powf(2.f, src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sqrt, sycl::sqrt(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sqrt, sqrt(src_cpp[i])) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(ieee_sqrt, sqrt(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(rsqrt, 1.f / sycl::sqrt(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sin, sycl::sin(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(cos, sycl::cos(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(rsqrt, 1.f / sqrt(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sin, sin(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(cos, cos(src_cpp[i])) #undef __ESIMD_UNARY_EXT_MATH_HOST_INTRIN @@ -627,7 +627,7 @@ __esimd_rndd(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - retv[i] = sycl::floor(src0[i]); + retv[i] = floor(src0[i]); } return retv; } @@ -640,13 +640,13 @@ __esimd_rndu(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - if (src0[i] - sycl::floor(src0[i]) > 0.0f) { + if (src0[i] - floor(src0[i]) > 0.0f) { increment = 1; } else { increment = 0; } - retv[i] = sycl::floor(src0[i]) + increment; + retv[i] = floor(src0[i]) + increment; } return retv; @@ -682,12 +682,12 @@ __esimd_rndz(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - if (sycl::fabs(src0[i]) < sycl::fabs(sycl::floor(src0[i]))) { + if (fabs(src0[i]) < fabs(floor(src0[i]))) { increment = 1; } else { increment = 0; } - retv[i] = sycl::floor(src0[i]) + increment; + retv[i] = floor(src0[i]) + increment; } return retv; diff --git a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp index 792e5ea61ddbd..cc2e83d754287 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp @@ -550,7 +550,7 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typedef typename cplx::detail::__libcpp_complex_overload_traits<_Tp>::_ValueType _ValueType; - return sycl::atan2(static_cast<_ValueType>(0), __re); + return sycl::atan2<_ValueType>(0, __re); } // norm diff --git a/sycl/test-e2e/Basic/half_builtins.cpp b/sycl/test-e2e/Basic/half_builtins.cpp index 9637742cb67bb..f92983045fa01 100644 --- a/sycl/test-e2e/Basic/half_builtins.cpp +++ b/sycl/test-e2e/Basic/half_builtins.cpp @@ -12,8 +12,7 @@ using namespace sycl; constexpr int SZ_max = 16; bool check(float a, float b) { - return sycl::fabs(2 * (a - b) / (a + b)) < - std::numeric_limits::epsilon() || + return fabs(2 * (a - b) / (a + b)) < std::numeric_limits::epsilon() || a < std::numeric_limits::min(); } @@ -181,10 +180,10 @@ int main() { auto err = err_buf.get_access(cgh); cgh.parallel_for(SZ_max, [=](item<1> index) { size_t i = index.get_id(0); - TEST_BUILTIN_1(sycl::fabs); - TEST_BUILTIN_2(sycl::fmin); - TEST_BUILTIN_2(sycl::fmax); - TEST_BUILTIN_3(sycl::fma); + TEST_BUILTIN_1(fabs); + TEST_BUILTIN_2(fmin); + TEST_BUILTIN_2(fmax); + TEST_BUILTIN_3(fma); }); }); } diff --git a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp index 875a3199ca15c..576290abfb936 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp @@ -393,7 +393,7 @@ float4 CalcLinearRead(typename FormatTraits::rep_elem_type *RefData, CoordT AdjCoord = Coord; if constexpr (AddrMode == addressing_mode::repeat) { assert(Normalized); - AdjCoord -= sycl::floor(AdjCoord); + AdjCoord -= floor(AdjCoord); AdjCoord *= RangeToCoord(ImageRange); } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) { assert(Normalized); diff --git a/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp b/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp index 4f79df9a88d73..bbaf8194fc501 100644 --- a/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp +++ b/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp @@ -35,8 +35,7 @@ void TestHelper(sycl::queue Q, const sycl::image_channel_type ChanType = sycl::image_channel_type::signed_int32; - const sycl::range<2> ImgSize(sycl::sqrt(static_cast(BUFFER_SIZE)), - sycl::sqrt(static_cast(BUFFER_SIZE))); + const sycl::range<2> ImgSize(sqrt(BUFFER_SIZE), sqrt(BUFFER_SIZE)); std::vector ImgHostData( ImgSize.size(), {InitialVal, InitialVal, InitialVal, InitialVal}); sycl::image<2> Img(ImgHostData.data(), ChanOrder, ChanType, ImgSize); diff --git a/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp b/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp index 180f2db54a214..974edb754e1d5 100644 --- a/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp +++ b/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp @@ -75,10 +75,10 @@ bool verify_result(Centroid4 *centroids4, // gpu centroids result int k = 0; int j = 0; for (auto i = 0; i < NUM_CENTROIDS_ACTUAL; i++) { - float errX = std::fabs(centroids4[j].x[k] - centroids[i].x) / - max(std::fabs(centroids4[j].x[k]), std::fabs(centroids[i].x)); - float errY = std::fabs(centroids4[j].y[k] - centroids[i].y) / - max(std::fabs(centroids4[j].y[k]), std::fabs(centroids[i].y)); + float errX = fabs(centroids4[j].x[k] - centroids[i].x) / + max(fabs(centroids4[j].x[k]), fabs(centroids[i].x)); + float errY = fabs(centroids4[j].y[k] - centroids[i].y) / + max(fabs(centroids4[j].y[k]), fabs(centroids[i].y)); float errSize = abs(centroids4[j].num_points[k] - centroids[i].num_points) / max(abs(centroids4[j].num_points[k]), abs(centroids[i].num_points)); From 34cd8e2627955ee9b976e4918a5d381cdad7be86 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:09:43 -0700 Subject: [PATCH 2/6] Revert "[SYCL] Make relational builtins overloads (#10050)" This reverts commit a2aec49fcebfc573d3de07abdd81648370e87e82. --- sycl/include/sycl/builtins.hpp | 534 +++++++++-------------- sycl/include/sycl/detail/type_traits.hpp | 22 +- 2 files changed, 217 insertions(+), 339 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 71d24e532a3c4..6598f9dea1b48 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -40,125 +40,6 @@ template marray to_marray(vec x) { marray[i] = x[i]; return marray; } - -// Vectors need fixed-size integers instead of fundamental types. -template -struct same_fixed_size_int; -template -struct same_fixed_size_int> { - using type = int8_t; -}; -template -struct same_fixed_size_int> { - using type = int16_t; -}; -template -struct same_fixed_size_int> { - using type = int32_t; -}; -template -struct same_fixed_size_int> { - using type = int64_t; -}; -template -struct same_fixed_size_int> { - using type = uint8_t; -}; -template -struct same_fixed_size_int> { - using type = uint16_t; -}; -template -struct same_fixed_size_int> { - using type = uint32_t; -}; -template -struct same_fixed_size_int> { - using type = uint64_t; -}; - -// Trait for getting an integer type of the same size as T. This propagates -// through vec and marray. -template struct same_size_int; -template -struct same_size_int< - T, true, std::enable_if_t && sizeof(T) == 1>> { - using type = signed char; -}; -template -struct same_size_int< - T, true, std::enable_if_t && sizeof(T) == 2>> { - using type = signed short; -}; -template -struct same_size_int< - T, true, std::enable_if_t && sizeof(T) == 4>> { - using type = signed int; -}; -template -struct same_size_int< - T, true, std::enable_if_t && sizeof(T) == 8>> { - using type = signed long long; -}; -template -struct same_size_int< - T, false, std::enable_if_t && sizeof(T) == 1>> { - using type = unsigned char; -}; -template -struct same_size_int< - T, false, std::enable_if_t && sizeof(T) == 2>> { - using type = unsigned short; -}; -template -struct same_size_int< - T, false, std::enable_if_t && sizeof(T) == 4>> { - using type = unsigned int; -}; -template -struct same_size_int< - T, false, std::enable_if_t && sizeof(T) == 8>> { - using type = unsigned long long; -}; -template -struct same_size_int, Signed> { - // Use the fixed-size integer types. - using type = vec::type, N>; -}; -template -struct same_size_int, Signed> { - using type = marray::type, N>; -}; - -// Trait for getting a floating point type of the same size as T. This -// propagates through vec and marray. -template struct same_size_float; -template -struct same_size_float< - T, std::enable_if_t && sizeof(T) == 2>> { - using type = half; -}; -template -struct same_size_float< - T, std::enable_if_t && sizeof(T) == 4>> { - using type = float; -}; -template -struct same_size_float< - T, std::enable_if_t && sizeof(T) == 8>> { - using type = double; -}; -template struct same_size_float> { - using type = vec::type, N>; -}; -template struct same_size_float> { - using type = marray::type, N>; -}; - -template -using same_size_int_t = typename same_size_int::type; -template -using same_size_float_t = typename same_size_float::type; } // namespace detail #ifdef __SYCL_DEVICE_ONLY__ @@ -370,7 +251,6 @@ namespace __sycl_std = __host_std; // longlongn and long{n} have the same types, so we only include one here. #define __SYCL_DEF_BUILTIN_VGENINTEGER \ __SYCL_DEF_BUILTIN_CHAR_VEC \ - __SYCL_DEF_BUILTIN_UCHAR_VEC \ __SYCL_DEF_BUILTIN_SHORT_VEC \ __SYCL_DEF_BUILTIN_USHORT_VEC \ __SYCL_DEF_BUILTIN_INT_VEC \ @@ -507,10 +387,6 @@ namespace __sycl_std = __host_std; __SYCL_DEF_BUILTIN_SGENINTEGER \ __SYCL_DEF_BUILTIN_SGENFLOAT -#define __SYCL_DEF_BUILTIN_VGENTYPE \ - __SYCL_DEF_BUILTIN_VGENINTEGER \ - __SYCL_DEF_BUILTIN_VGENFLOAT - #define __SYCL_DEF_BUILTIN_GENTYPE \ __SYCL_DEF_BUILTIN_GENINTEGER \ __SYCL_DEF_BUILTIN_GENFLOAT @@ -2297,127 +2173,106 @@ __SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ /* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/ -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isequal(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdEqual>(x, \ - y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isequal(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdEqual>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isnotequal(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FUnordNotEqual>( \ - x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isnotequal(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FUnordNotEqual>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isgreater(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdGreaterThan< \ - detail::internal_rel_ret_t>(x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isgreater(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdGreaterThan>(x, + y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isgreaterequal(TYPE x, TYPE y) \ - __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdGreaterThanEqual< \ - detail::internal_rel_ret_t>(x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isgreaterequal(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdGreaterThanEqual>( + x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isless(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdLessThan>( \ - x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isless(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdLessThan>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t islessequal(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdLessThanEqual< \ - detail::internal_rel_ret_t>(x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t islessequal(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdLessThanEqual>(x, + y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t islessgreater(TYPE x, TYPE y) \ - __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_FOrdNotEqual>( \ - x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_FOrdNotEqual>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isfinite(TYPE x) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_IsFinite>(x)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isfinite(T x) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_IsFinite>(x)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isinf(TYPE x) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_IsInf>(x)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isinf(T x) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_IsInf>(x)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isnan(TYPE x) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_IsNan>(x)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isnan(T x) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_IsNan>(x)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isnormal(TYPE x) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_IsNormal>(x)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isnormal(T x) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_IsNormal>(x)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isordered(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_Ordered>(x, y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isordered(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_Ordered>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t isunordered(TYPE x, TYPE y) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_Unordered>(x, \ - y)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t isunordered(T x, T y) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_Unordered>(x, y)); +} -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::common_rel_ret_t signbit(TYPE x) __NOEXC { \ - return detail::RelConverter::apply( \ - __sycl_std::__invoke_SignBitSet>(x)); \ - } -__SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +detail::common_rel_ret_t signbit(T x) __NOEXC { + return detail::RelConverter::apply( + __sycl_std::__invoke_SignBitSet>(x)); +} // marray relational functions @@ -2459,107 +2314,143 @@ __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(isunordered) __SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(signbit) // bool any (sigeninteger x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline bool any(TYPE x) __NOEXC { \ - return detail::Boolean<1>(int(detail::msbIsSet(x))); \ - } -__SYCL_DEF_BUILTIN_SIGENINTEGER -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, bool> any(T x) __NOEXC { + return detail::Boolean<1>(int(detail::msbIsSet(x))); +} // int any (vigeninteger x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline int any(TYPE x) __NOEXC { \ - return detail::rel_sign_bit_test_ret_t( \ - __sycl_std::__invoke_Any>( \ - detail::rel_sign_bit_test_arg_t(x))); \ - } -__SYCL_DEF_BUILTIN_VIGENINTEGER -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, int> any(T x) __NOEXC { + return detail::rel_sign_bit_test_ret_t( + __sycl_std::__invoke_Any>( + detail::rel_sign_bit_test_arg_t(x))); +} // bool all (sigeninteger x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline bool all(TYPE x) __NOEXC { \ - return detail::Boolean<1>(int(detail::msbIsSet(x))); \ - } -__SYCL_DEF_BUILTIN_SIGENINTEGER -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, bool> all(T x) __NOEXC { + return detail::Boolean<1>(int(detail::msbIsSet(x))); +} // int all (vigeninteger x) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline int all(TYPE x) __NOEXC { \ - return detail::rel_sign_bit_test_ret_t( \ - __sycl_std::__invoke_All>( \ - detail::rel_sign_bit_test_arg_t(x))); \ - } -__SYCL_DEF_BUILTIN_VIGENINTEGER -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, int> all(T x) __NOEXC { + return detail::rel_sign_bit_test_ret_t( + __sycl_std::__invoke_All>( + detail::rel_sign_bit_test_arg_t(x))); +} // gentype bitselect (gentype a, gentype b, gentype c) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE bitselect(TYPE a, TYPE b, TYPE c) __NOEXC { \ - return __sycl_std::__invoke_bitselect(a, b, c); \ - } -__SYCL_DEF_BUILTIN_GENTYPE -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> bitselect(T a, T b, + T c) __NOEXC { + return __sycl_std::__invoke_bitselect(a, b, c); +} // sgentype select (sgentype a, sgentype b, bool c) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE select(TYPE a, TYPE b, bool c) __NOEXC { \ - /* sycl::select(sgentype a, sgentype b, bool c) calls OpenCL built-in \ - select(sgentype a, sgentype b, igentype c). This type trait makes the \ - proper conversion for argument c from bool to igentype, based on sgentype \ - == T.*/ \ - using get_select_opencl_builtin_c_arg_type = \ - detail::same_size_int_t>; \ - \ - return __sycl_std::__invoke_select( \ - a, b, static_cast(c)); \ - } -__SYCL_DEF_BUILTIN_SGENTYPE -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> select(T a, T b, + bool c) __NOEXC { + constexpr size_t SizeT = sizeof(T); + + // sycl::select(sgentype a, sgentype b, bool c) calls OpenCL built-in + // select(sgentype a, sgentype b, igentype c). This type trait makes the + // proper conversion for argument c from bool to igentype, based on sgentype + // == T. + using get_select_opencl_builtin_c_arg_type = typename std::conditional_t< + SizeT == 1, char, + std::conditional_t< + SizeT == 2, short, + std::conditional_t< + (detail::is_contained< + T, detail::type_list>::value && + (SizeT == 4 || SizeT == 8)), + long, // long and ulong are 32-bit on + // Windows and 64-bit on Linux + std::conditional_t< + SizeT == 4, int, + std::conditional_t>>>>; + + return __sycl_std::__invoke_select( + a, b, static_cast(c)); +} + +// geninteger select (geninteger a, geninteger b, igeninteger c) +template +std::enable_if_t< + detail::is_geninteger::value && detail::is_igeninteger::value, T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} -// vgentype select(vgentype a, vgentype b, vigeninteger c) -// vgentype select(vgentype a, vgentype b, vugeninteger c) -// Non-standard: -// sgentype select(sgentype a, sgentype b, sigeninteger c) -// sgentype select(sgentype a, sgentype b, sugeninteger c) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE select(TYPE a, TYPE b, detail::same_size_int_t c) \ - __NOEXC { \ - return __sycl_std::__invoke_select(a, b, c); \ - } \ - inline TYPE select(TYPE a, TYPE b, detail::same_size_int_t c) \ - __NOEXC { \ - return __sycl_std::__invoke_select(a, b, c); \ - } -__SYCL_DEF_BUILTIN_VGENTYPE -__SYCL_DEF_BUILTIN_SGENTYPE -#undef __SYCL_BUILTIN_DEF +// geninteger select (geninteger a, geninteger b, ugeninteger c) +template +std::enable_if_t< + detail::is_geninteger::value && detail::is_ugeninteger::value, T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} -// Since same_size_int_t uses long long for 64-bit as it is guaranteed to have -// the appropriate size, we need special cases for long. -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline detail::same_size_int_t select( \ - detail::same_size_int_t a, \ - detail::same_size_int_t b, TYPE c) __NOEXC { \ - return __sycl_std::__invoke_select>( \ - a, b, c); \ - } \ - inline detail::same_size_int_t select( \ - detail::same_size_int_t a, \ - detail::same_size_int_t b, TYPE c) __NOEXC { \ - return __sycl_std::__invoke_select>( \ - a, b, c); \ - } \ - inline detail::same_size_float_t select( \ - detail::same_size_float_t a, detail::same_size_float_t b, \ - TYPE c) __NOEXC { \ - return __sycl_std::__invoke_select>(a, b, \ - c); \ - } -__SYCL_DEF_BUILTIN_LONG_SCALAR -__SYCL_DEF_BUILTIN_ULONG_SCALAR -#undef __SYCL_BUILTIN_DEF +// svgenfloatf select (svgenfloatf a, svgenfloatf b, genint c) +template +std::enable_if_t< + detail::is_svgenfloatf::value && detail::is_genint::value, T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} + +// svgenfloatf select (svgenfloatf a, svgenfloatf b, ugenint c) +template +std::enable_if_t< + detail::is_svgenfloatf::value && detail::is_ugenint::value, T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} + +// svgenfloatd select (svgenfloatd a, svgenfloatd b, igeninteger64 c) +template +std::enable_if_t::value && + detail::is_igeninteger64bit::value, + T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} + +// svgenfloatd select (svgenfloatd a, svgenfloatd b, ugeninteger64 c) +template +std::enable_if_t::value && + detail::is_ugeninteger64bit::value, + T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} + +// svgenfloath select (svgenfloath a, svgenfloath b, igeninteger16 c) +template +std::enable_if_t::value && + detail::is_igeninteger16bit::value, + T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} + +// svgenfloath select (svgenfloath a, svgenfloath b, ugeninteger16 c) +template +std::enable_if_t::value && + detail::is_ugeninteger16bit::value, + T> +select(T a, T b, T2 c) __NOEXC { + detail::check_vector_size(); + return __sycl_std::__invoke_select(a, b, c); +} // other marray relational functions @@ -3077,7 +2968,6 @@ std::enable_if_t::value, T> tan(T x) __NOEXC { #undef __SYCL_DEF_BUILTIN_GENGEOFLOAT #undef __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT #undef __SYCL_DEF_BUILTIN_SGENTYPE -#undef __SYCL_DEF_BUILTIN_VGENTYPE #undef __SYCL_DEF_BUILTIN_GENTYPE #undef __SYCL_COMMA } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp index a46fffc97a96a..be072531a7a14 100644 --- a/sycl/include/sycl/detail/type_traits.hpp +++ b/sycl/include/sycl/detail/type_traits.hpp @@ -261,29 +261,17 @@ using is_gen_based_on_type_sizeof = std::bool_constant::value && (sizeof(vector_element_t) == N)>; template struct is_vec : std::false_type {}; -template struct is_vec> : std::true_type {}; - -template constexpr bool is_vec_v = is_vec::value; +template +struct is_vec> : std::true_type {}; template struct get_vec_size { - static constexpr int size = 1; + static constexpr std::size_t size = 1; }; -template struct get_vec_size> { - static constexpr int size = N; +template struct get_vec_size> { + static constexpr std::size_t size = N; }; -// is_marray -template struct is_marray : std::false_type {}; -template -struct is_marray> : std::true_type {}; - -template constexpr bool is_marray_v = is_marray::value; - -// is_marray_or_vec_v -template -constexpr bool is_marray_or_vec_v = is_marray_v || is_vec_v; - // is_integral template struct is_integral : std::is_integral> {}; From 7246ee00cba22b7f10095d3ea1781d6c58c353a1 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:09:45 -0700 Subject: [PATCH 3/6] Revert "[SYCL] Make SYCL geometrical marray functions overloads instead of templates (#10048)" This reverts commit 3c95fc49222be7e7427e468827ff1db7524d0823. --- sycl/include/sycl/builtins.hpp | 148 ++++++++++++--------------------- 1 file changed, 51 insertions(+), 97 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 6598f9dea1b48..fb94b77e71649 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -48,8 +48,6 @@ template marray to_marray(vec x) { namespace __sycl_std = __host_std; #endif -#define __SYCL_COMMA , - #define __SYCL_DEF_BUILTIN_VEC(TYPE) \ __SYCL_BUILTIN_DEF(TYPE##2) \ __SYCL_BUILTIN_DEF(TYPE##3) \ @@ -66,15 +64,6 @@ namespace __sycl_std = __host_std; __SYCL_BUILTIN_DEF(TYPE##3) \ __SYCL_BUILTIN_DEF(TYPE##4) -#define __SYCL_DEF_BUILTIN_GEOMARRAY(TYPE) \ - __SYCL_BUILTIN_DEF(marray) \ - __SYCL_BUILTIN_DEF(marray) \ - __SYCL_BUILTIN_DEF(marray) - -#define __SYCL_DEF_BUILTIN_GEOCROSSMARRAY(TYPE) \ - __SYCL_BUILTIN_DEF(marray) \ - __SYCL_BUILTIN_DEF(marray) - #define __SYCL_DEF_BUILTIN_MARRAY(TYPE) #define __SYCL_DEF_BUILTIN_CHAR_SCALAR __SYCL_BUILTIN_DEF(char) @@ -270,9 +259,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_FLOAT_SCALAR __SYCL_BUILTIN_DEF(float) #define __SYCL_DEF_BUILTIN_FLOAT_VEC __SYCL_DEF_BUILTIN_VEC(float) #define __SYCL_DEF_BUILTIN_FLOAT_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(float) -#define __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_GEOCROSSMARRAY(float) -#define __SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY __SYCL_DEF_BUILTIN_GEOMARRAY(float) #define __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSVEC \ __SYCL_DEF_BUILTIN_GEOCROSSVEC(float) #define __SYCL_DEF_BUILTIN_FLOAT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(float) @@ -289,9 +275,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_DOUBLE_SCALAR __SYCL_BUILTIN_DEF(double) #define __SYCL_DEF_BUILTIN_DOUBLE_VEC __SYCL_DEF_BUILTIN_VEC(double) #define __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_GEOCROSSMARRAY(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_GEOMARRAY __SYCL_DEF_BUILTIN_GEOMARRAY(double) #define __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSVEC \ __SYCL_DEF_BUILTIN_GEOCROSSVEC(double) #define __SYCL_DEF_BUILTIN_DOUBLE_MARRAY __SYCL_DEF_BUILTIN_MARRAY(double) @@ -308,9 +291,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_HALF_SCALAR __SYCL_BUILTIN_DEF(half) #define __SYCL_DEF_BUILTIN_HALF_VEC __SYCL_DEF_BUILTIN_VEC(half) #define __SYCL_DEF_BUILTIN_HALF_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(half) -#define __SYCL_DEF_BUILTIN_HALF_GEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_GEOCROSSMARRAY(half) -#define __SYCL_DEF_BUILTIN_HALF_GEOMARRAY __SYCL_DEF_BUILTIN_GEOMARRAY(half) #define __SYCL_DEF_BUILTIN_HALF_GEOCROSSVEC __SYCL_DEF_BUILTIN_GEOCROSSVEC(half) #define __SYCL_DEF_BUILTIN_HALF_MARRAY __SYCL_DEF_BUILTIN_MARRAY(half) #define __SYCL_DEF_BUILTIN_HALFN \ @@ -343,17 +323,6 @@ namespace __sycl_std = __host_std; __SYCL_DEF_BUILTIN_GENGEOFLOATD \ __SYCL_DEF_BUILTIN_GENGEOFLOATH -#define __SYCL_DEF_BUILTIN_GENGEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSMARRAY \ - __SYCL_DEF_BUILTIN_HALF_GEOCROSSMARRAY - -#define __SYCL_DEF_BUILTIN_GENGEOMARRAY \ - __SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY \ - __SYCL_DEF_BUILTIN_DOUBLE_GEOMARRAY \ - __SYCL_DEF_BUILTIN_HALF_GEOMARRAY - -// TODO: Replace with overloads. #define __SYCL_DEF_BUILTIN_VGENGEOCROSSFLOAT \ __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSVEC \ __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSVEC \ @@ -2094,81 +2063,67 @@ __SYCL_DEF_BUILTIN_GENGEOFLOATF // marray geometric functions -// cross -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cross(TYPE p0, TYPE p1) __NOEXC { \ - return detail::to_marray(cross(detail::to_vec(p0), detail::to_vec(p1))); \ - } -__SYCL_DEF_BUILTIN_GENGEOCROSSMARRAY -#undef __SYCL_BUILTIN_DEF +#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME, ...) \ + vec, T::size()> result_v; \ + result_v = NAME(__VA_ARGS__); \ + return detail::to_marray(result_v); -#undef __SYCL_DEF_BUILTIN_GENGEOCROSSMARRAY -#undef __SYCL_DEF_BUILTIN_HALF_GEOCROSSMARRAY -#undef __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSMARRAY -#undef __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSMARRAY -#undef __SYCL_DEF_BUILTIN_GEOCROSSMARRAY +template +std::enable_if_t::value, T> cross(T p0, + T p1) __NOEXC { + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(cross, detail::to_vec(p0), + detail::to_vec(p1)) +} -// dot -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::value_type dot(TYPE p0, TYPE p1) __NOEXC { \ - return dot(detail::to_vec(p0), detail::to_vec(p1)); \ - } -__SYCL_DEF_BUILTIN_GENGEOMARRAY -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> normalize(T p) __NOEXC { + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(normalize, detail::to_vec(p)) +} -// distance -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::value_type distance(TYPE p0, TYPE p1) __NOEXC { \ - return distance(detail::to_vec(p0), detail::to_vec(p1)); \ - } -__SYCL_DEF_BUILTIN_GENGEOMARRAY -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> +fast_normalize(T p) __NOEXC { + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(fast_normalize, + detail::to_vec(p)) +} -// length -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::value_type length(TYPE p) __NOEXC { \ - return length(detail::to_vec(p)); \ - } -__SYCL_DEF_BUILTIN_GENGEOMARRAY -#undef __SYCL_BUILTIN_DEF +#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL -// normalize -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE normalize(TYPE p) __NOEXC { \ - return detail::to_marray(normalize(detail::to_vec(p))); \ +#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME) \ + template \ + std::enable_if_t::value, \ + detail::marray_element_t> \ + NAME(T p0, T p1) __NOEXC { \ + return NAME(detail::to_vec(p0), detail::to_vec(p1)); \ } -__SYCL_DEF_BUILTIN_GENGEOMARRAY -#undef __SYCL_BUILTIN_DEF -// fast_distance -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline float fast_distance(TYPE p0, TYPE p1) __NOEXC { \ - return fast_distance(detail::to_vec(p0), detail::to_vec(p1)); \ - } -__SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY -#undef __SYCL_BUILTIN_DEF +// clang-format off +__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(dot) +__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(distance) +// clang-format on -// fast_normalize -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fast_normalize(TYPE p) __NOEXC { \ - return detail::to_marray(fast_normalize(detail::to_vec(p))); \ - } -__SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY -#undef __SYCL_BUILTIN_DEF +#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD -// fast_length -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline float fast_length(TYPE p) __NOEXC { \ - return fast_length(detail::to_vec(p)); \ - } -__SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, detail::marray_element_t> +length(T p) __NOEXC { + return __sycl_std::__invoke_length>( + detail::to_vec(p)); +} -#undef __SYCL_DEF_BUILTIN_GENGEOMARRAY -#undef __SYCL_DEF_BUILTIN_HALF_GEOMARRAY -#undef __SYCL_DEF_BUILTIN_DOUBLE_GEOMARRAY -#undef __SYCL_DEF_BUILTIN_FLOAT_GEOMARRAY -#undef __SYCL_DEF_BUILTIN_GEOMARRAY +template +std::enable_if_t::value, + detail::marray_element_t> +fast_distance(T p0, T p1) __NOEXC { + return fast_distance(detail::to_vec(p0), detail::to_vec(p1)); +} + +template +std::enable_if_t::value, + detail::marray_element_t> +fast_length(T p) __NOEXC { + return fast_length(detail::to_vec(p)); +} /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ /* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/ @@ -2969,7 +2924,6 @@ std::enable_if_t::value, T> tan(T x) __NOEXC { #undef __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT #undef __SYCL_DEF_BUILTIN_SGENTYPE #undef __SYCL_DEF_BUILTIN_GENTYPE -#undef __SYCL_COMMA } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl From 799b71a91aa37dbf0c68e375cad1805aaf9e7f85 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:09:48 -0700 Subject: [PATCH 4/6] Revert "[SYCL] Make SYCL geometrical functions overloads instead of templates (#10014)" This reverts commit 49bbcbf52f2f553b740aa6c46650061f05fa99ff. --- sycl/include/sycl/builtins.hpp | 221 ++++++++++++++++++--------------- 1 file changed, 119 insertions(+), 102 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index fb94b77e71649..96a8aa6a560e9 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -60,10 +60,6 @@ namespace __sycl_std = __host_std; __SYCL_BUILTIN_DEF(TYPE##3) \ __SYCL_BUILTIN_DEF(TYPE##4) -#define __SYCL_DEF_BUILTIN_GEOCROSSVEC(TYPE) \ - __SYCL_BUILTIN_DEF(TYPE##3) \ - __SYCL_BUILTIN_DEF(TYPE##4) - #define __SYCL_DEF_BUILTIN_MARRAY(TYPE) #define __SYCL_DEF_BUILTIN_CHAR_SCALAR __SYCL_BUILTIN_DEF(char) @@ -259,8 +255,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_FLOAT_SCALAR __SYCL_BUILTIN_DEF(float) #define __SYCL_DEF_BUILTIN_FLOAT_VEC __SYCL_DEF_BUILTIN_VEC(float) #define __SYCL_DEF_BUILTIN_FLOAT_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(float) -#define __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSVEC \ - __SYCL_DEF_BUILTIN_GEOCROSSVEC(float) #define __SYCL_DEF_BUILTIN_FLOAT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(float) #define __SYCL_DEF_BUILTIN_FLOATN \ __SYCL_DEF_BUILTIN_FLOAT_VEC \ @@ -275,8 +269,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_DOUBLE_SCALAR __SYCL_BUILTIN_DEF(double) #define __SYCL_DEF_BUILTIN_DOUBLE_VEC __SYCL_DEF_BUILTIN_VEC(double) #define __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSVEC \ - __SYCL_DEF_BUILTIN_GEOCROSSVEC(double) #define __SYCL_DEF_BUILTIN_DOUBLE_MARRAY __SYCL_DEF_BUILTIN_MARRAY(double) #define __SYCL_DEF_BUILTIN_DOUBLEN \ __SYCL_DEF_BUILTIN_DOUBLE_VEC \ @@ -291,7 +283,6 @@ namespace __sycl_std = __host_std; #define __SYCL_DEF_BUILTIN_HALF_SCALAR __SYCL_BUILTIN_DEF(half) #define __SYCL_DEF_BUILTIN_HALF_VEC __SYCL_DEF_BUILTIN_VEC(half) #define __SYCL_DEF_BUILTIN_HALF_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(half) -#define __SYCL_DEF_BUILTIN_HALF_GEOCROSSVEC __SYCL_DEF_BUILTIN_GEOCROSSVEC(half) #define __SYCL_DEF_BUILTIN_HALF_MARRAY __SYCL_DEF_BUILTIN_MARRAY(half) #define __SYCL_DEF_BUILTIN_HALFN \ __SYCL_DEF_BUILTIN_HALF_VEC \ @@ -323,17 +314,8 @@ namespace __sycl_std = __host_std; __SYCL_DEF_BUILTIN_GENGEOFLOATD \ __SYCL_DEF_BUILTIN_GENGEOFLOATH -#define __SYCL_DEF_BUILTIN_VGENGEOCROSSFLOAT \ - __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSVEC \ - __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSVEC \ - __SYCL_DEF_BUILTIN_HALF_GEOCROSSVEC - -#define __SYCL_DEF_BUILTIN_VGENGEOFLOAT \ - __SYCL_DEF_BUILTIN_FLOAT_GEOVEC \ - __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC \ - __SYCL_DEF_BUILTIN_HALF_GEOVEC - // TODO: Replace with overloads. + #ifdef __FAST_MATH__ #define __FAST_MATH_GENFLOAT(T) \ (detail::is_svgenfloatd::value || detail::is_svgenfloath::value) @@ -1960,106 +1942,141 @@ __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(upsample, 32bit) // double4 cross (double4 p0, double4 p1) // half3 cross (half3 p0, half3 p1) // half4 cross (half4 p0, half4 p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE cross(TYPE p0, TYPE p1) __NOEXC { \ - return __sycl_std::__invoke_cross(p0, p1); \ - } -__SYCL_DEF_BUILTIN_VGENGEOCROSSFLOAT -#undef __SYCL_BUILTIN_DEF -#undef __SYCL_DEF_BUILTIN_VGENGEOCROSSFLOAT -#undef __SYCL_DEF_BUILTIN_HALF_GEOCROSSVEC -#undef __SYCL_DEF_BUILTIN_DOUBLE_GEOCROSSVEC -#undef __SYCL_DEF_BUILTIN_FLOAT_GEOCROSSVEC -#undef __SYCL_DEF_BUILTIN_GEOCROSSVEC +template +std::enable_if_t::value, T> cross(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_cross(p0, p1); +} // float dot (float p0, float p1) // double dot (double p0, double p1) // half dot (half p0, half p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE dot(TYPE p0, TYPE p1) __NOEXC { return p0 * p1; } -__SYCL_DEF_BUILTIN_SGENFLOAT -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> dot(T p0, T p1) __NOEXC { + return p0 * p1; +} + // float dot (vgengeofloat p0, vgengeofloat p1) +template +std::enable_if_t::value, float> dot(T p0, + T p1) __NOEXC { + return __sycl_std::__invoke_Dot(p0, p1); +} + // double dot (vgengeodouble p0, vgengeodouble p1) +template +std::enable_if_t::value, double> dot(T p0, + T p1) __NOEXC { + return __sycl_std::__invoke_Dot(p0, p1); +} + // half dot (vgengeohalf p0, vgengeohalf p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::element_type dot(TYPE p0, TYPE p1) __NOEXC { \ - return __sycl_std::__invoke_Dot(p0, p1); \ - } -__SYCL_DEF_BUILTIN_VGENGEOFLOAT -#undef __SYCL_BUILTIN_DEF - -// float distance (float p0, float p1) -// double distance (double p0, double p1) -// half distance (half p0, half p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE distance(TYPE p0, TYPE p1) __NOEXC { \ - return __sycl_std::__invoke_distance(p0, p1); \ - } -__SYCL_DEF_BUILTIN_SGENFLOAT -#undef __SYCL_BUILTIN_DEF -// float distance (vgengeofloat p0, vgengeofloat p1) -// double distance (vgengeodouble p0, vgengeodouble p1) -// half distance (vgengeohalf p0, vgengeohalf p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::element_type distance(TYPE p0, TYPE p1) __NOEXC { \ - return __sycl_std::__invoke_distance(p0, p1); \ - } -__SYCL_DEF_BUILTIN_VGENGEOFLOAT -#undef __SYCL_BUILTIN_DEF - -// float length (float p0, float p1) -// double length (double p0, double p1) -// half length (half p0, half p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE length(TYPE p) __NOEXC { \ - return __sycl_std::__invoke_length(p); \ - } -__SYCL_DEF_BUILTIN_SGENFLOAT -#undef __SYCL_BUILTIN_DEF -// float length (vgengeofloat p0, vgengeofloat p1) -// double length (vgengeodouble p0, vgengeodouble p1) -// half length (vgengeohalf p0, vgengeohalf p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE::element_type length(TYPE p) __NOEXC { \ - return __sycl_std::__invoke_length(p); \ - } -__SYCL_DEF_BUILTIN_VGENGEOFLOAT -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, half> dot(T p0, + T p1) __NOEXC { + return __sycl_std::__invoke_Dot(p0, p1); +} + +// float distance (gengeofloat p0, gengeofloat p1) +template ::value, T>> +float distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); +} + +// double distance (gengeodouble p0, gengeodouble p1) +template ::value, T>> +double distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); +} + +// half distance (gengeohalf p0, gengeohalf p1) +template ::value, T>> +half distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); +} + +// float length (gengeofloat p) +template ::value, T>> +float length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); +} + +// double length (gengeodouble p) +template ::value, T>> +double length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); +} + +// half length (gengeohalf p) +template ::value, T>> +half length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); +} // gengeofloat normalize (gengeofloat p) +template +std::enable_if_t::value, T> normalize(T p) __NOEXC { + return __sycl_std::__invoke_normalize(p); +} + // gengeodouble normalize (gengeodouble p) +template +std::enable_if_t::value, T> normalize(T p) __NOEXC { + return __sycl_std::__invoke_normalize(p); +} + // gengeohalf normalize (gengeohalf p) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE normalize(TYPE p) __NOEXC { \ - return __sycl_std::__invoke_normalize(p); \ - } -__SYCL_DEF_BUILTIN_GENGEOFLOAT -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> normalize(T p) __NOEXC { + return __sycl_std::__invoke_normalize(p); +} // float fast_distance (gengeofloat p0, gengeofloat p1) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline float fast_distance(TYPE p0, TYPE p1) __NOEXC { \ - return __sycl_std::__invoke_fast_distance(p0, p1); \ - } -__SYCL_DEF_BUILTIN_GENGEOFLOATF -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +float fast_distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_fast_distance(p0, p1); +} + +// double fast_distance (gengeodouble p0, gengeodouble p1) +template ::value, T>> +double fast_distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_fast_distance(p0, p1); +} // float fast_length (gengeofloat p) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline float fast_length(TYPE p) __NOEXC { \ - return __sycl_std::__invoke_fast_length(p); \ - } -__SYCL_DEF_BUILTIN_GENGEOFLOATF -#undef __SYCL_BUILTIN_DEF +template ::value, T>> +float fast_length(T p) __NOEXC { + return __sycl_std::__invoke_fast_length(p); +} + +// double fast_length (gengeodouble p) +template ::value, T>> +double fast_length(T p) __NOEXC { + return __sycl_std::__invoke_fast_length(p); +} // gengeofloat fast_normalize (gengeofloat p) -#define __SYCL_BUILTIN_DEF(TYPE) \ - inline TYPE fast_normalize(TYPE p) __NOEXC { \ - return __sycl_std::__invoke_fast_normalize(p); \ - } -__SYCL_DEF_BUILTIN_GENGEOFLOATF -#undef __SYCL_BUILTIN_DEF +template +std::enable_if_t::value, T> +fast_normalize(T p) __NOEXC { + return __sycl_std::__invoke_fast_normalize(p); +} + +// gengeodouble fast_normalize (gengeodouble p) +template +std::enable_if_t::value, T> +fast_normalize(T p) __NOEXC { + return __sycl_std::__invoke_fast_normalize(p); +} // marray geometric functions From 8105019ca284ccc093cc7bcf2cb2c78b781ad394 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:09:51 -0700 Subject: [PATCH 5/6] Revert "[SYCL] Add defines for built-in transition (#10026)" This reverts commit fe61b999a1a73a8890b600ba66460f92ce836cb3. --- sycl/include/sycl/builtins.hpp | 381 --------------------------------- 1 file changed, 381 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 96a8aa6a560e9..c56d985d25f07 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -48,274 +48,6 @@ template marray to_marray(vec x) { namespace __sycl_std = __host_std; #endif -#define __SYCL_DEF_BUILTIN_VEC(TYPE) \ - __SYCL_BUILTIN_DEF(TYPE##2) \ - __SYCL_BUILTIN_DEF(TYPE##3) \ - __SYCL_BUILTIN_DEF(TYPE##4) \ - __SYCL_BUILTIN_DEF(TYPE##8) \ - __SYCL_BUILTIN_DEF(TYPE##16) - -#define __SYCL_DEF_BUILTIN_GEOVEC(TYPE) \ - __SYCL_BUILTIN_DEF(TYPE##2) \ - __SYCL_BUILTIN_DEF(TYPE##3) \ - __SYCL_BUILTIN_DEF(TYPE##4) - -#define __SYCL_DEF_BUILTIN_MARRAY(TYPE) - -#define __SYCL_DEF_BUILTIN_CHAR_SCALAR __SYCL_BUILTIN_DEF(char) -#define __SYCL_DEF_BUILTIN_CHAR_VEC __SYCL_DEF_BUILTIN_VEC(char) -#define __SYCL_DEF_BUILTIN_CHAR_MARRAY __SYCL_DEF_BUILTIN_MARRAY(char) -#define __SYCL_DEF_BUILTIN_CHARN \ - __SYCL_DEF_BUILTIN_CHAR_VEC \ - __SYCL_DEF_BUILTIN_CHAR_MARRAY -#define __SYCL_DEF_BUILTIN_SCHAR_SCALAR __SYCL_BUILTIN_DEF(signed char) -#define __SYCL_DEF_BUILTIN_SCHAR_VEC __SYCL_DEF_BUILTIN_VEC(schar) -#define __SYCL_DEF_BUILTIN_SCHAR_MARRAY __SYCL_DEF_BUILTIN_MARRAY(signed char) -#define __SYCL_DEF_BUILTIN_SCHARN \ - __SYCL_DEF_BUILTIN_SCHAR_VEC \ - __SYCL_DEF_BUILTIN_SCHAR_MARRAY -#define __SYCL_DEF_BUILTIN_IGENCHAR \ - __SYCL_DEF_BUILTIN_SCHAR_SCALAR \ - __SYCL_DEF_BUILTIN_SCHARN -#define __SYCL_DEF_BUILTIN_UCHAR_SCALAR __SYCL_BUILTIN_DEF(unsigned char) -#define __SYCL_DEF_BUILTIN_UCHAR_VEC __SYCL_DEF_BUILTIN_VEC(uchar) -#define __SYCL_DEF_BUILTIN_UCHAR_MARRAY __SYCL_DEF_BUILTIN_MARRAY(unsigned char) -#define __SYCL_DEF_BUILTIN_UCHARN \ - __SYCL_DEF_BUILTIN_UCHAR_VEC \ - __SYCL_DEF_BUILTIN_UCHAR_MARRAY -#define __SYCL_DEF_BUILTIN_UGENCHAR \ - __SYCL_DEF_BUILTIN_UCHAR_SCALAR \ - __SYCL_DEF_BUILTIN_UCHARN -// schar{n} and char{n} have the same type, so we skip the char{n} variants. -#define __SYCL_DEF_BUILTIN_GENCHAR \ - __SYCL_DEF_BUILTIN_CHAR_SCALAR \ - __SYCL_DEF_BUILTIN_CHAR_MARRAY \ - __SYCL_DEF_BUILTIN_IGENCHAR \ - __SYCL_DEF_BUILTIN_UGENCHAR - -#define __SYCL_DEF_BUILTIN_SHORT_SCALAR __SYCL_BUILTIN_DEF(short) -#define __SYCL_DEF_BUILTIN_SHORT_VEC __SYCL_DEF_BUILTIN_VEC(short) -#define __SYCL_DEF_BUILTIN_SHORT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(short) -#define __SYCL_DEF_BUILTIN_SHORTN \ - __SYCL_DEF_BUILTIN_SHORT_VEC \ - __SYCL_DEF_BUILTIN_SHORT_MARRAY -#define __SYCL_DEF_BUILTIN_GENSHORT \ - __SYCL_DEF_BUILTIN_SHORT_SCALAR \ - __SYCL_DEF_BUILTIN_SHORTN -#define __SYCL_DEF_BUILTIN_USHORT_SCALAR __SYCL_BUILTIN_DEF(unsigned short) -#define __SYCL_DEF_BUILTIN_USHORT_VEC __SYCL_DEF_BUILTIN_VEC(ushort) -#define __SYCL_DEF_BUILTIN_USHORT_MARRAY \ - __SYCL_DEF_BUILTIN_MARRAY(unsigned short) -#define __SYCL_DEF_BUILTIN_USHORTN \ - __SYCL_DEF_BUILTIN_USHORT_VEC \ - __SYCL_DEF_BUILTIN_USHORT_MARRAY -#define __SYCL_DEF_BUILTIN_UGENSHORT \ - __SYCL_DEF_BUILTIN_USHORT_SCALAR \ - __SYCL_DEF_BUILTIN_USHORTN - -#define __SYCL_DEF_BUILTIN_INT_SCALAR __SYCL_BUILTIN_DEF(int) -#define __SYCL_DEF_BUILTIN_INT_VEC __SYCL_DEF_BUILTIN_VEC(int) -#define __SYCL_DEF_BUILTIN_INT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(int) -#define __SYCL_DEF_BUILTIN_INTN \ - __SYCL_DEF_BUILTIN_INT_VEC \ - __SYCL_DEF_BUILTIN_INT_MARRAY -#define __SYCL_DEF_BUILTIN_GENINT \ - __SYCL_DEF_BUILTIN_INT_SCALAR \ - __SYCL_DEF_BUILTIN_INTN -#define __SYCL_DEF_BUILTIN_UINT_SCALAR __SYCL_BUILTIN_DEF(unsigned int) -#define __SYCL_DEF_BUILTIN_UINT_VEC __SYCL_DEF_BUILTIN_VEC(uint) -#define __SYCL_DEF_BUILTIN_UINT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(unsigned int) -#define __SYCL_DEF_BUILTIN_UINTN \ - __SYCL_DEF_BUILTIN_UINT_VEC \ - __SYCL_DEF_BUILTIN_UINT_MARRAY -#define __SYCL_DEF_BUILTIN_UGENINT \ - __SYCL_DEF_BUILTIN_UINT_SCALAR \ - __SYCL_DEF_BUILTIN_UINTN - -#define __SYCL_DEF_BUILTIN_LONG_SCALAR __SYCL_BUILTIN_DEF(long) -#define __SYCL_DEF_BUILTIN_LONG_VEC __SYCL_DEF_BUILTIN_VEC(long) -#define __SYCL_DEF_BUILTIN_LONG_MARRAY __SYCL_DEF_BUILTIN_MARRAY(long) -#define __SYCL_DEF_BUILTIN_LONGN \ - __SYCL_DEF_BUILTIN_LONG_VEC \ - __SYCL_DEF_BUILTIN_LONG_MARRAY -#define __SYCL_DEF_BUILTIN_GENLONG \ - __SYCL_DEF_BUILTIN_LONG_SCALAR \ - __SYCL_DEF_BUILTIN_LONGN -#define __SYCL_DEF_BUILTIN_ULONG_SCALAR __SYCL_BUILTIN_DEF(unsigned long) -#define __SYCL_DEF_BUILTIN_ULONG_VEC __SYCL_DEF_BUILTIN_VEC(ulong) -#define __SYCL_DEF_BUILTIN_ULONG_MARRAY __SYCL_DEF_BUILTIN_MARRAY(unsigned long) -#define __SYCL_DEF_BUILTIN_ULONGN \ - __SYCL_DEF_BUILTIN_ULONG_VEC \ - __SYCL_DEF_BUILTIN_ULONG_MARRAY -#define __SYCL_DEF_BUILTIN_UGENLONG \ - __SYCL_DEF_BUILTIN_ULONG_SCALAR \ - __SYCL_DEF_BUILTIN_ULONGN - -#define __SYCL_DEF_BUILTIN_LONGLONG_SCALAR __SYCL_BUILTIN_DEF(long long) -#define __SYCL_DEF_BUILTIN_LONGLONG_VEC __SYCL_DEF_BUILTIN_VEC(longlong) -#define __SYCL_DEF_BUILTIN_LONGLONG_MARRAY __SYCL_DEF_BUILTIN_MARRAY(long long) -#define __SYCL_DEF_BUILTIN_LONGLONGN \ - __SYCL_DEF_BUILTIN_LONGLONG_VEC \ - __SYCL_DEF_BUILTIN_LONGLONG_MARRAY -#define __SYCL_DEF_BUILTIN_GENLONGLONG \ - __SYCL_DEF_BUILTIN_LONGLONG_SCALAR \ - __SYCL_DEF_BUILTIN_LONGLONGN -#define __SYCL_DEF_BUILTIN_ULONGLONG_SCALAR \ - __SYCL_BUILTIN_DEF(unsigned long long) -#define __SYCL_DEF_BUILTIN_ULONGLONG_VEC __SYCL_DEF_BUILTIN_VEC(ulonglong) -#define __SYCL_DEF_BUILTIN_ULONGLONG_MARRAY \ - __SYCL_DEF_BUILTIN_MARRAY(unsigned long long) -#define __SYCL_DEF_BUILTIN_ULONGLONGN \ - __SYCL_DEF_BUILTIN_ULONGLONG_VEC \ - __SYCL_DEF_BUILTIN_ULONGLONG_MARRAY -#define __SYCL_DEF_BUILTIN_UGENLONGLONG \ - __SYCL_DEF_BUILTIN_ULONGLONG_SCALAR \ - __SYCL_DEF_BUILTIN_ULONGLONGN - -// longlongn and long{n} have the same types, so we only include one here. -#define __SYCL_DEF_BUILTIN_IGENLONGINTEGER \ - __SYCL_DEF_BUILTIN_LONG_SCALAR \ - __SYCL_DEF_BUILTIN_LONG_MARRAY \ - __SYCL_DEF_BUILTIN_LONGLONG_SCALAR \ - __SYCL_DEF_BUILTIN_LONGLONG_MARRAY \ - __SYCL_DEF_BUILTIN_LONG_VEC - -// longlong{n} and long{n} have the same types, so we only include one here. -#define __SYCL_DEF_BUILTIN_UGENLONGINTEGER \ - __SYCL_DEF_BUILTIN_ULONG_SCALAR \ - __SYCL_DEF_BUILTIN_ULONG_MARRAY \ - __SYCL_DEF_BUILTIN_ULONGLONG_SCALAR \ - __SYCL_DEF_BUILTIN_ULONGLONG_MARRAY \ - __SYCL_DEF_BUILTIN_ULONG_VEC - -#define __SYCL_DEF_BUILTIN_SIGENINTEGER \ - __SYCL_DEF_BUILTIN_SCHAR_SCALAR \ - __SYCL_DEF_BUILTIN_SHORT_SCALAR \ - __SYCL_DEF_BUILTIN_INT_SCALAR \ - __SYCL_DEF_BUILTIN_LONG_SCALAR \ - __SYCL_DEF_BUILTIN_LONGLONG_SCALAR - -// longlongn and longn have the same types, so we only include one here. -#define __SYCL_DEF_BUILTIN_VIGENINTEGER \ - __SYCL_DEF_BUILTIN_CHAR_VEC \ - __SYCL_DEF_BUILTIN_SHORT_VEC \ - __SYCL_DEF_BUILTIN_INT_VEC \ - __SYCL_DEF_BUILTIN_LONG_VEC - -#define __SYCL_DEF_BUILTIN_IGENINTEGER \ - __SYCL_DEF_BUILTIN_IGENCHAR \ - __SYCL_DEF_BUILTIN_GENSHORT \ - __SYCL_DEF_BUILTIN_GENINT \ - __SYCL_DEF_BUILTIN_IGENLONGINTEGER - -#define __SYCL_DEF_BUILTIN_SUGENINTEGER \ - __SYCL_DEF_BUILTIN_UCHAR_SCALAR \ - __SYCL_DEF_BUILTIN_USHORT_SCALAR \ - __SYCL_DEF_BUILTIN_UINT_SCALAR \ - __SYCL_DEF_BUILTIN_ULONG_SCALAR \ - __SYCL_DEF_BUILTIN_ULONGLONG_SCALAR - -// longlongn and longn have the same types, so we only include one here. -#define __SYCL_DEF_BUILTIN_VUGENINTEGER \ - __SYCL_DEF_BUILTIN_UCHAR_VEC \ - __SYCL_DEF_BUILTIN_USHORT_VEC \ - __SYCL_DEF_BUILTIN_UINT_VEC \ - __SYCL_DEF_BUILTIN_ULONG_VEC - -#define __SYCL_DEF_BUILTIN_UGENINTEGER \ - __SYCL_DEF_BUILTIN_UGENCHAR \ - __SYCL_DEF_BUILTIN_UGENSHORT \ - __SYCL_DEF_BUILTIN_UGENINT \ - __SYCL_DEF_BUILTIN_UGENLONGINTEGER - -#define __SYCL_DEF_BUILTIN_SGENINTEGER \ - __SYCL_DEF_BUILTIN_CHAR_SCALAR \ - __SYCL_DEF_BUILTIN_SIGENINTEGER \ - __SYCL_DEF_BUILTIN_SUGENINTEGER - -// longlongn and long{n} have the same types, so we only include one here. -#define __SYCL_DEF_BUILTIN_VGENINTEGER \ - __SYCL_DEF_BUILTIN_CHAR_VEC \ - __SYCL_DEF_BUILTIN_SHORT_VEC \ - __SYCL_DEF_BUILTIN_USHORT_VEC \ - __SYCL_DEF_BUILTIN_INT_VEC \ - __SYCL_DEF_BUILTIN_UINT_VEC \ - __SYCL_DEF_BUILTIN_LONG_VEC \ - __SYCL_DEF_BUILTIN_ULONG_VEC - -#define __SYCL_DEF_BUILTIN_GENINTEGER \ - __SYCL_DEF_BUILTIN_GENCHAR \ - __SYCL_DEF_BUILTIN_GENSHORT \ - __SYCL_DEF_BUILTIN_UGENSHORT \ - __SYCL_DEF_BUILTIN_GENINT \ - __SYCL_DEF_BUILTIN_UGENINT \ - __SYCL_DEF_BUILTIN_UGENLONGINTEGER \ - __SYCL_DEF_BUILTIN_IGENLONGINTEGER - -#define __SYCL_DEF_BUILTIN_FLOAT_SCALAR __SYCL_BUILTIN_DEF(float) -#define __SYCL_DEF_BUILTIN_FLOAT_VEC __SYCL_DEF_BUILTIN_VEC(float) -#define __SYCL_DEF_BUILTIN_FLOAT_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(float) -#define __SYCL_DEF_BUILTIN_FLOAT_MARRAY __SYCL_DEF_BUILTIN_MARRAY(float) -#define __SYCL_DEF_BUILTIN_FLOATN \ - __SYCL_DEF_BUILTIN_FLOAT_VEC \ - __SYCL_DEF_BUILTIN_FLOAT_MARRAY -#define __SYCL_DEF_BUILTIN_GENFLOATF \ - __SYCL_DEF_BUILTIN_FLOAT_SCALAR \ - __SYCL_DEF_BUILTIN_FLOATN -#define __SYCL_DEF_BUILTIN_GENGEOFLOATF \ - __SYCL_DEF_BUILTIN_FLOAT_SCALAR \ - __SYCL_DEF_BUILTIN_FLOAT_GEOVEC - -#define __SYCL_DEF_BUILTIN_DOUBLE_SCALAR __SYCL_BUILTIN_DEF(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_VEC __SYCL_DEF_BUILTIN_VEC(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(double) -#define __SYCL_DEF_BUILTIN_DOUBLE_MARRAY __SYCL_DEF_BUILTIN_MARRAY(double) -#define __SYCL_DEF_BUILTIN_DOUBLEN \ - __SYCL_DEF_BUILTIN_DOUBLE_VEC \ - __SYCL_DEF_BUILTIN_DOUBLE_MARRAY -#define __SYCL_DEF_BUILTIN_GENFLOATD \ - __SYCL_DEF_BUILTIN_DOUBLE_SCALAR \ - __SYCL_DEF_BUILTIN_DOUBLEN -#define __SYCL_DEF_BUILTIN_GENGEOFLOATD \ - __SYCL_DEF_BUILTIN_DOUBLE_SCALAR \ - __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC - -#define __SYCL_DEF_BUILTIN_HALF_SCALAR __SYCL_BUILTIN_DEF(half) -#define __SYCL_DEF_BUILTIN_HALF_VEC __SYCL_DEF_BUILTIN_VEC(half) -#define __SYCL_DEF_BUILTIN_HALF_GEOVEC __SYCL_DEF_BUILTIN_GEOVEC(half) -#define __SYCL_DEF_BUILTIN_HALF_MARRAY __SYCL_DEF_BUILTIN_MARRAY(half) -#define __SYCL_DEF_BUILTIN_HALFN \ - __SYCL_DEF_BUILTIN_HALF_VEC \ - __SYCL_DEF_BUILTIN_HALF_MARRAY -#define __SYCL_DEF_BUILTIN_GENFLOATH \ - __SYCL_DEF_BUILTIN_HALF_SCALAR \ - __SYCL_DEF_BUILTIN_HALFN -#define __SYCL_DEF_BUILTIN_GENGEOFLOATH \ - __SYCL_DEF_BUILTIN_HALF_SCALAR \ - __SYCL_DEF_BUILTIN_HALF_GEOVEC - -#define __SYCL_DEF_BUILTIN_SGENFLOAT \ - __SYCL_DEF_BUILTIN_FLOAT_SCALAR \ - __SYCL_DEF_BUILTIN_DOUBLE_SCALAR \ - __SYCL_DEF_BUILTIN_HALF_SCALAR - -#define __SYCL_DEF_BUILTIN_VGENFLOAT \ - __SYCL_DEF_BUILTIN_FLOAT_VEC \ - __SYCL_DEF_BUILTIN_DOUBLE_VEC \ - __SYCL_DEF_BUILTIN_HALF_VEC - -#define __SYCL_DEF_BUILTIN_GENFLOAT \ - __SYCL_DEF_BUILTIN_GENFLOATF \ - __SYCL_DEF_BUILTIN_GENFLOATD \ - __SYCL_DEF_BUILTIN_GENFLOATH - -#define __SYCL_DEF_BUILTIN_GENGEOFLOAT \ - __SYCL_DEF_BUILTIN_GENGEOFLOATF \ - __SYCL_DEF_BUILTIN_GENGEOFLOATD \ - __SYCL_DEF_BUILTIN_GENGEOFLOATH - -// TODO: Replace with overloads. - #ifdef __FAST_MATH__ #define __FAST_MATH_GENFLOAT(T) \ (detail::is_svgenfloatd::value || detail::is_svgenfloath::value) @@ -326,22 +58,6 @@ namespace __sycl_std = __host_std; #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat::value) #endif -#ifdef __FAST_MATH__ -#define __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT \ - __SYCL_DEF_BUILTIN_GENFLOATD \ - __SYCL_DEF_BUILTIN_GENFLOATH -#else -#define __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT __SYCL_DEF_BUILTIN_GENFLOAT -#endif - -#define __SYCL_DEF_BUILTIN_SGENTYPE \ - __SYCL_DEF_BUILTIN_SGENINTEGER \ - __SYCL_DEF_BUILTIN_SGENFLOAT - -#define __SYCL_DEF_BUILTIN_GENTYPE \ - __SYCL_DEF_BUILTIN_GENINTEGER \ - __SYCL_DEF_BUILTIN_GENFLOAT - /* ----------------- 4.13.3 Math functions. ---------------------------------*/ // These macros for marray math function implementations use vectorizations of @@ -2844,103 +2560,6 @@ std::enable_if_t::value, T> tan(T x) __NOEXC { } #endif // __FAST_MATH__ - -#undef __SYCL_DEF_BUILTIN_VEC -#undef __SYCL_DEF_BUILTIN_GEOVEC -#undef __SYCL_DEF_BUILTIN_MARRAY -#undef __SYCL_DEF_BUILTIN_CHAR_SCALAR -#undef __SYCL_DEF_BUILTIN_CHAR_VEC -#undef __SYCL_DEF_BUILTIN_CHAR_MARRAY -#undef __SYCL_DEF_BUILTIN_CHARN -#undef __SYCL_DEF_BUILTIN_SCHAR_SCALAR -#undef __SYCL_DEF_BUILTIN_SCHAR_VEC -#undef __SYCL_DEF_BUILTIN_SCHAR_MARRAY -#undef __SYCL_DEF_BUILTIN_SCHARN -#undef __SYCL_DEF_BUILTIN_IGENCHAR -#undef __SYCL_DEF_BUILTIN_UCHAR_SCALAR -#undef __SYCL_DEF_BUILTIN_UCHAR_VEC -#undef __SYCL_DEF_BUILTIN_UCHAR_MARRAY -#undef __SYCL_DEF_BUILTIN_UCHARN -#undef __SYCL_DEF_BUILTIN_UGENCHAR -#undef __SYCL_DEF_BUILTIN_GENCHAR -#undef __SYCL_DEF_BUILTIN_SHORT_SCALAR -#undef __SYCL_DEF_BUILTIN_SHORT_VEC -#undef __SYCL_DEF_BUILTIN_SHORT_MARRAY -#undef __SYCL_DEF_BUILTIN_SHORTN -#undef __SYCL_DEF_BUILTIN_GENSHORT -#undef __SYCL_DEF_BUILTIN_USHORT_SCALAR -#undef __SYCL_DEF_BUILTIN_USHORT_MARRAY -#undef __SYCL_DEF_BUILTIN_USHORTN -#undef __SYCL_DEF_BUILTIN_UGENSHORT -#undef __SYCL_DEF_BUILTIN_INT_SCALAR -#undef __SYCL_DEF_BUILTIN_INT_VEC -#undef __SYCL_DEF_BUILTIN_INT_MARRAY -#undef __SYCL_DEF_BUILTIN_INTN -#undef __SYCL_DEF_BUILTIN_GENINT -#undef __SYCL_DEF_BUILTIN_UINT_SCALAR -#undef __SYCL_DEF_BUILTIN_UINT_VEC -#undef __SYCL_DEF_BUILTIN_UINT_MARRAY -#undef __SYCL_DEF_BUILTIN_UINTN -#undef __SYCL_DEF_BUILTIN_UGENINT -#undef __SYCL_DEF_BUILTIN_LONG_SCALAR -#undef __SYCL_DEF_BUILTIN_LONG_VEC -#undef __SYCL_DEF_BUILTIN_LONG_MARRAY -#undef __SYCL_DEF_BUILTIN_LONGN -#undef __SYCL_DEF_BUILTIN_GENLONG -#undef __SYCL_DEF_BUILTIN_ULONG_SCALAR -#undef __SYCL_DEF_BUILTIN_ULONG_VEC -#undef __SYCL_DEF_BUILTIN_ULONG_MARRAY -#undef __SYCL_DEF_BUILTIN_ULONGN -#undef __SYCL_DEF_BUILTIN_UGENLONG -#undef __SYCL_DEF_BUILTIN_LONGLONG_SCALAR -#undef __SYCL_DEF_BUILTIN_LONGLONG_VEC -#undef __SYCL_DEF_BUILTIN_LONGLONG_MARRAY -#undef __SYCL_DEF_BUILTIN_LONGLONGN -#undef __SYCL_DEF_BUILTIN_GENLONGLONG -#undef __SYCL_DEF_BUILTIN_ULONGLONG_SCALAR -#undef __SYCL_DEF_BUILTIN_ULONGLONG_VEC -#undef __SYCL_DEF_BUILTIN_ULONGLONG_MARRAY -#undef __SYCL_DEF_BUILTIN_ULONGLONGN -#undef __SYCL_DEF_BUILTIN_UGENLONGLONG -#undef __SYCL_DEF_BUILTIN_IGENLONGINTEGER -#undef __SYCL_DEF_BUILTIN_UGENLONGINTEGER -#undef __SYCL_DEF_BUILTIN_SIGENINTEGER -#undef __SYCL_DEF_BUILTIN_VIGENINTEGER -#undef __SYCL_DEF_BUILTIN_IGENINTEGER -#undef __SYCL_DEF_BUILTIN_SUGENINTEGER -#undef __SYCL_DEF_BUILTIN_VUGENINTEGER -#undef __SYCL_DEF_BUILTIN_UGENINTEGER -#undef __SYCL_DEF_BUILTIN_SGENINTEGER -#undef __SYCL_DEF_BUILTIN_VGENINTEGER -#undef __SYCL_DEF_BUILTIN_GENINTEGER -#undef __SYCL_DEF_BUILTIN_FLOAT_SCALAR -#undef __SYCL_DEF_BUILTIN_FLOAT_VEC -#undef __SYCL_DEF_BUILTIN_FLOAT_GEOVEC -#undef __SYCL_DEF_BUILTIN_FLOAT_MARRAY -#undef __SYCL_DEF_BUILTIN_FLOATN -#undef __SYCL_DEF_BUILTIN_GENFLOATF -#undef __SYCL_DEF_BUILTIN_GENGEOFLOATF -#undef __SYCL_DEF_BUILTIN_DOUBLE_SCALAR -#undef __SYCL_DEF_BUILTIN_DOUBLE_VEC -#undef __SYCL_DEF_BUILTIN_DOUBLE_GEOVEC -#undef __SYCL_DEF_BUILTIN_DOUBLE_MARRAY -#undef __SYCL_DEF_BUILTIN_DOUBLEN -#undef __SYCL_DEF_BUILTIN_GENFLOATD -#undef __SYCL_DEF_BUILTIN_GENGEOFLOATD -#undef __SYCL_DEF_BUILTIN_HALF_SCALAR -#undef __SYCL_DEF_BUILTIN_HALF_VEC -#undef __SYCL_DEF_BUILTIN_HALF_GEOVEC -#undef __SYCL_DEF_BUILTIN_HALF_MARRAY -#undef __SYCL_DEF_BUILTIN_HALFN -#undef __SYCL_DEF_BUILTIN_GENFLOATH -#undef __SYCL_DEF_BUILTIN_GENGEOFLOATH -#undef __SYCL_DEF_BUILTIN_SGENFLOAT -#undef __SYCL_DEF_BUILTIN_VGENFLOAT -#undef __SYCL_DEF_BUILTIN_GENFLOAT -#undef __SYCL_DEF_BUILTIN_GENGEOFLOAT -#undef __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT -#undef __SYCL_DEF_BUILTIN_SGENTYPE -#undef __SYCL_DEF_BUILTIN_GENTYPE } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl From db3a256908221c8314e9d746ae0a3b20bad4ef7d Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 3 Jul 2023 02:14:31 -0700 Subject: [PATCH 6/6] Keep valid fixes Signed-off-by: Larsen, Steffen --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 8 ++++---- sycl/include/sycl/detail/type_traits.hpp | 9 ++++----- .../ext/intel/esimd/detail/math_intrin.hpp | 18 +++++++++--------- sycl/test-e2e/Basic/half_builtins.cpp | 11 ++++++----- .../test-e2e/Basic/sycl_2020_images/common.hpp | 2 +- .../discard_events_check_images.cpp | 3 ++- sycl/test-e2e/ESIMD/kmeans/kmeans.cpp | 8 ++++---- 7 files changed, 30 insertions(+), 29 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 0c29357f7672f..1185dd35c4a91 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -53,10 +53,10 @@ static const char *LegalSYCLFunctions[] = { "^sycl::_V1::sub_group<.+>::.+", "^sycl::_V1::range<.+>::.+", "^sycl::_V1::kernel_handler::.+", - "^sycl::_V1::cos<.+>", - "^sycl::_V1::sin<.+>", - "^sycl::_V1::log<.+>", - "^sycl::_V1::exp<.+>", + "^sycl::_V1::cos", + "^sycl::_V1::sin", + "^sycl::_V1::log", + "^sycl::_V1::exp", "^sycl::_V1::bit_cast<.+>", "^sycl::_V1::operator.+<.+>", "^sycl::_V1::ext::oneapi::sub_group::.+", diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp index be072531a7a14..a9778b97ce1d9 100644 --- a/sycl/include/sycl/detail/type_traits.hpp +++ b/sycl/include/sycl/detail/type_traits.hpp @@ -261,15 +261,14 @@ using is_gen_based_on_type_sizeof = std::bool_constant::value && (sizeof(vector_element_t) == N)>; template struct is_vec : std::false_type {}; -template -struct is_vec> : std::true_type {}; +template struct is_vec> : std::true_type {}; template struct get_vec_size { - static constexpr std::size_t size = 1; + static constexpr int size = 1; }; -template struct get_vec_size> { - static constexpr std::size_t size = N; +template struct get_vec_size> { + static constexpr int size = N; }; // is_integral diff --git a/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp b/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp index 602baa6634037..db0af8d212c1a 100644 --- a/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp +++ b/sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp @@ -572,11 +572,11 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, SZ) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(inv, 1.f / src_cpp[i]) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(log, logf(src_cpp[i]) / logf(2.f)) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(exp, powf(2.f, src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sqrt, sqrt(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sqrt, sycl::sqrt(src_cpp[i])) __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(ieee_sqrt, sqrt(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(rsqrt, 1.f / sqrt(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sin, sin(src_cpp[i])) -__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(cos, cos(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(rsqrt, 1.f / sycl::sqrt(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sin, sycl::sin(src_cpp[i])) +__ESIMD_UNARY_EXT_MATH_HOST_INTRIN(cos, sycl::cos(src_cpp[i])) #undef __ESIMD_UNARY_EXT_MATH_HOST_INTRIN @@ -627,7 +627,7 @@ __esimd_rndd(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - retv[i] = floor(src0[i]); + retv[i] = sycl::floor(src0[i]); } return retv; } @@ -640,13 +640,13 @@ __esimd_rndu(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - if (src0[i] - floor(src0[i]) > 0.0f) { + if (src0[i] - sycl::floor(src0[i]) > 0.0f) { increment = 1; } else { increment = 0; } - retv[i] = floor(src0[i]) + increment; + retv[i] = sycl::floor(src0[i]) + increment; } return retv; @@ -682,12 +682,12 @@ __esimd_rndz(__ESIMD_DNS::vector_type_t src0) { for (int i = 0; i < SZ; i++) { SIMDCF_ELEMENT_SKIP(i); - if (fabs(src0[i]) < fabs(floor(src0[i]))) { + if (sycl::fabs(src0[i]) < sycl::fabs(sycl::floor(src0[i]))) { increment = 1; } else { increment = 0; } - retv[i] = floor(src0[i]) + increment; + retv[i] = sycl::floor(src0[i]) + increment; } return retv; diff --git a/sycl/test-e2e/Basic/half_builtins.cpp b/sycl/test-e2e/Basic/half_builtins.cpp index f92983045fa01..9637742cb67bb 100644 --- a/sycl/test-e2e/Basic/half_builtins.cpp +++ b/sycl/test-e2e/Basic/half_builtins.cpp @@ -12,7 +12,8 @@ using namespace sycl; constexpr int SZ_max = 16; bool check(float a, float b) { - return fabs(2 * (a - b) / (a + b)) < std::numeric_limits::epsilon() || + return sycl::fabs(2 * (a - b) / (a + b)) < + std::numeric_limits::epsilon() || a < std::numeric_limits::min(); } @@ -180,10 +181,10 @@ int main() { auto err = err_buf.get_access(cgh); cgh.parallel_for(SZ_max, [=](item<1> index) { size_t i = index.get_id(0); - TEST_BUILTIN_1(fabs); - TEST_BUILTIN_2(fmin); - TEST_BUILTIN_2(fmax); - TEST_BUILTIN_3(fma); + TEST_BUILTIN_1(sycl::fabs); + TEST_BUILTIN_2(sycl::fmin); + TEST_BUILTIN_2(sycl::fmax); + TEST_BUILTIN_3(sycl::fma); }); }); } diff --git a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp index 576290abfb936..875a3199ca15c 100644 --- a/sycl/test-e2e/Basic/sycl_2020_images/common.hpp +++ b/sycl/test-e2e/Basic/sycl_2020_images/common.hpp @@ -393,7 +393,7 @@ float4 CalcLinearRead(typename FormatTraits::rep_elem_type *RefData, CoordT AdjCoord = Coord; if constexpr (AddrMode == addressing_mode::repeat) { assert(Normalized); - AdjCoord -= floor(AdjCoord); + AdjCoord -= sycl::floor(AdjCoord); AdjCoord *= RangeToCoord(ImageRange); } else if constexpr (AddrMode == addressing_mode::mirrored_repeat) { assert(Normalized); diff --git a/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp b/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp index bbaf8194fc501..4f79df9a88d73 100644 --- a/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp +++ b/sycl/test-e2e/DiscardEvents/discard_events_check_images.cpp @@ -35,7 +35,8 @@ void TestHelper(sycl::queue Q, const sycl::image_channel_type ChanType = sycl::image_channel_type::signed_int32; - const sycl::range<2> ImgSize(sqrt(BUFFER_SIZE), sqrt(BUFFER_SIZE)); + const sycl::range<2> ImgSize(sycl::sqrt(static_cast(BUFFER_SIZE)), + sycl::sqrt(static_cast(BUFFER_SIZE))); std::vector ImgHostData( ImgSize.size(), {InitialVal, InitialVal, InitialVal, InitialVal}); sycl::image<2> Img(ImgHostData.data(), ChanOrder, ChanType, ImgSize); diff --git a/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp b/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp index 974edb754e1d5..180f2db54a214 100644 --- a/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp +++ b/sycl/test-e2e/ESIMD/kmeans/kmeans.cpp @@ -75,10 +75,10 @@ bool verify_result(Centroid4 *centroids4, // gpu centroids result int k = 0; int j = 0; for (auto i = 0; i < NUM_CENTROIDS_ACTUAL; i++) { - float errX = fabs(centroids4[j].x[k] - centroids[i].x) / - max(fabs(centroids4[j].x[k]), fabs(centroids[i].x)); - float errY = fabs(centroids4[j].y[k] - centroids[i].y) / - max(fabs(centroids4[j].y[k]), fabs(centroids[i].y)); + float errX = std::fabs(centroids4[j].x[k] - centroids[i].x) / + max(std::fabs(centroids4[j].x[k]), std::fabs(centroids[i].x)); + float errY = std::fabs(centroids4[j].y[k] - centroids[i].y) / + max(std::fabs(centroids4[j].y[k]), std::fabs(centroids[i].y)); float errSize = abs(centroids4[j].num_points[k] - centroids[i].num_points) / max(abs(centroids4[j].num_points[k]), abs(centroids[i].num_points));