From 826d868af1c33b9630881c615737343c01746293 Mon Sep 17 00:00:00 2001
From: Nikita Kornev <nikita.kornev@intel.com>
Date: Thu, 29 Jun 2023 16:47:20 +0200
Subject: [PATCH] [SYCL] Make SYCL math functions overloads instead of
 templates (#9753)

It was decided to change all built-in functions from templates to
overloads. This patch changes non-marray part of "4.17.5. Math
functions".

Spec: https://github.com/KhronosGroup/SYCL-Docs/pull/428

---------

Co-authored-by: Steffen Larsen <steffen.larsen@intel.com>
---
 llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp  |    8 +-
 sycl/include/sycl/builtins.hpp                | 1206 ++++++++++-------
 .../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, 709 insertions(+), 549 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/builtins.hpp b/sycl/include/sycl/builtins.hpp
index 71d24e532a3c4..f4391fd44a711 100644
--- a/sycl/include/sycl/builtins.hpp
+++ b/sycl/include/sycl/builtins.hpp
@@ -486,12 +486,9 @@ namespace __sycl_std = __host_std;
 
 // TODO: Replace with overloads.
 #ifdef __FAST_MATH__
-#define __FAST_MATH_GENFLOAT(T)                                                \
-  (detail::is_svgenfloatd<T>::value || detail::is_svgenfloath<T>::value)
 #define __FAST_MATH_SGENFLOAT(T)                                               \
   (std::is_same_v<T, double> || std::is_same_v<T, half>)
 #else
-#define __FAST_MATH_GENFLOAT(T) (detail::is_svgenfloat<T>::value)
 #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat<T>::value)
 #endif
 
@@ -517,6 +514,7 @@ 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
@@ -775,207 +773,261 @@ __SYCL_MATH_FUNCTION_3_OVERLOAD(mad) __SYCL_MATH_FUNCTION_3_OVERLOAD(mix)
 
 #undef __SYCL_MATH_FUNCTION_3_OVERLOAD
 
-    // svgenfloat acos (svgenfloat x)
-    template <typename T>
-    std::enable_if_t<detail::is_svgenfloat<T>::value, T> acos(T x) __NOEXC {
-  return __sycl_std::__invoke_acos<T>(x);
-}
+// genfloat acos (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE acos(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_acos<TYPE>(x);                                 \
+  }
+        __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat acosh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> acosh(T x) __NOEXC {
-  return __sycl_std::__invoke_acosh<T>(x);
-}
+// genfloat acosh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE acosh(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_acosh<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat acospi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> acospi(T x) __NOEXC {
-  return __sycl_std::__invoke_acospi<T>(x);
-}
+// genfloat acospi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE acospi(TYPE x) __NOEXC {                                         \
+    return __sycl_std::__invoke_acospi<TYPE>(x);                               \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat asin (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> asin(T x) __NOEXC {
-  return __sycl_std::__invoke_asin<T>(x);
-}
+// genfloat asin (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE asin(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_asin<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat asinh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> asinh(T x) __NOEXC {
-  return __sycl_std::__invoke_asinh<T>(x);
-}
+// genfloat asinh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE asinh(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_asinh<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat asinpi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> asinpi(T x) __NOEXC {
-  return __sycl_std::__invoke_asinpi<T>(x);
-}
+// genfloat asinpi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE asinpi(TYPE x) __NOEXC {                                         \
+    return __sycl_std::__invoke_asinpi<TYPE>(x);                               \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat atan (svgenfloat y_over_x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> atan(T y_over_x) __NOEXC {
-  return __sycl_std::__invoke_atan<T>(y_over_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<TYPE>(y_over_x);                          \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat atan2 (svgenfloat y, svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> atan2(T y, T x) __NOEXC {
-  return __sycl_std::__invoke_atan2<T>(y, 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<TYPE>(y, x);                             \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat atanh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> atanh(T x) __NOEXC {
-  return __sycl_std::__invoke_atanh<T>(x);
-}
+// genfloat atanh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE atanh(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_atanh<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat atanpi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> atanpi(T x) __NOEXC {
-  return __sycl_std::__invoke_atanpi<T>(x);
-}
+// genfloat atanpi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE atanpi(TYPE x) __NOEXC {                                         \
+    return __sycl_std::__invoke_atanpi<TYPE>(x);                               \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat atan2pi (svgenfloat y, svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> atan2pi(T y, T x) __NOEXC {
-  return __sycl_std::__invoke_atan2pi<T>(y, 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<TYPE>(y, x);                           \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat cbrt (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> cbrt(T x) __NOEXC {
-  return __sycl_std::__invoke_cbrt<T>(x);
-}
+// genfloat cbrt (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cbrt(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_cbrt<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat ceil (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> ceil(T x) __NOEXC {
-  return __sycl_std::__invoke_ceil<T>(x);
-}
+// genfloat ceil (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE ceil(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_ceil<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat copysign (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> copysign(T x,
-                                                              T y) __NOEXC {
-  return __sycl_std::__invoke_copysign<T>(x, y);
-}
+// genfloat copysign (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE copysign(TYPE y, TYPE x) __NOEXC {                               \
+    return __sycl_std::__invoke_copysign<TYPE>(y, x);                          \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat cos (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> cos(T x) __NOEXC {
-  return __sycl_std::__invoke_cos<T>(x);
-}
+// genfloat cos (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cos(TYPE x) __NOEXC { return __sycl_std::__invoke_cos<TYPE>(x); }
+    __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat cosh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> cosh(T x) __NOEXC {
-  return __sycl_std::__invoke_cosh<T>(x);
-}
+// genfloat cosh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cosh(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_cosh<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat cospi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> cospi(T x) __NOEXC {
-  return __sycl_std::__invoke_cospi<T>(x);
-}
+// genfloat cospi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cospi(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_cospi<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat erfc (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> erfc(T x) __NOEXC {
-  return __sycl_std::__invoke_erfc<T>(x);
-}
+// genfloat erfc (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE erfc(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_erfc<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat erf (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> erf(T x) __NOEXC {
-  return __sycl_std::__invoke_erf<T>(x);
-}
+// genfloat erf (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE erf(TYPE x) __NOEXC { return __sycl_std::__invoke_erf<TYPE>(x); }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat exp (svgenfloat x )
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp(T x) __NOEXC {
-  return __sycl_std::__invoke_exp<T>(x);
-}
+// genfloat exp (genfloat x )
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp(TYPE x) __NOEXC { return __sycl_std::__invoke_exp<TYPE>(x); }
+    __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat exp2 (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp2(T x) __NOEXC {
-  return __sycl_std::__invoke_exp2<T>(x);
-}
+// genfloat exp2 (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_exp2<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat exp10 (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> exp10(T x) __NOEXC {
-  return __sycl_std::__invoke_exp10<T>(x);
-}
+// genfloat exp10 (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_exp10<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat expm1 (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> expm1(T x) __NOEXC {
-  return __sycl_std::__invoke_expm1<T>(x);
-}
+// genfloat expm1 (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE expm1(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_expm1<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fabs (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fabs(T x) __NOEXC {
-  return __sycl_std::__invoke_fabs<T>(x);
-}
+// genfloat fabs (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE fabs(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_fabs<TYPE>(x);                                 \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fdim (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fdim(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_fdim<T>(x, y);
-}
+// genfloat fdim (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE fdim(TYPE x, TYPE y) __NOEXC {                                   \
+    return __sycl_std::__invoke_fdim<TYPE>(x, y);                              \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat floor (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> floor(T x) __NOEXC {
-  return __sycl_std::__invoke_floor<T>(x);
-}
+// genfloat floor (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE floor(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_floor<TYPE>(x);                                \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fma (svgenfloat a, svgenfloat b, svgenfloat c)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fma(T a, T b,
-                                                         T c) __NOEXC {
-  return __sycl_std::__invoke_fma<T>(a, b, c);
-}
+// 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<TYPE>(a, b, c);                            \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fmax (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fmax(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_fmax<T>(x, y);
-}
+// genfloat fmax (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE fmax(TYPE x, TYPE y) __NOEXC {                                   \
+    return __sycl_std::__invoke_fmax<TYPE>(x, y);                              \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fmax (svgenfloat x, sgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_vgenfloat<T>::value, T>
-fmax(T x, typename T::element_type y) __NOEXC {
-  return __sycl_std::__invoke_fmax<T>(x, T(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<TYPE>(x, TYPE(y));                        \
+  }
+    __SYCL_DEF_BUILTIN_VGENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fmin (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fmin(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_fmin<T>(x, 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<TYPE>(x, y);                              \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fmin (svgenfloat x, sgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_vgenfloat<T>::value, T>
-fmin(T x, typename T::element_type y) __NOEXC {
-  return __sycl_std::__invoke_fmin<T>(x, T(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<TYPE>(x, TYPE(y));                        \
+  }
+    __SYCL_DEF_BUILTIN_VGENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fmod (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> fmod(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_fmod<T>(x, 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<TYPE>(x, y);                              \
+  }
+    __SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat fract (svgenfloat x, genfloatptr iptr)
-template <typename T, typename T2>
-std::enable_if_t<
-    detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
-fract(T x, T2 iptr) __NOEXC {
+    // svgenfloat fract (svgenfloat x, genfloatptr iptr)
+    template <typename T, typename T2>
+    std::enable_if_t<detail::is_svgenfloat<T>::value &&
+                         detail::is_genfloatptr<T2>::value,
+                     T> fract(T x, T2 iptr) __NOEXC {
   detail::check_vector_size<T, T2>();
   return __sycl_std::__invoke_fract<T>(x, iptr);
 }
@@ -989,32 +1041,40 @@ frexp(T x, T2 exp) __NOEXC {
   return __sycl_std::__invoke_frexp<T>(x, exp);
 }
 
-// svgenfloat hypot (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> hypot(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_hypot<T>(x, y);
-}
+// genfloat hypot (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE hypot(TYPE x, TYPE y) __NOEXC {                                  \
+    return __sycl_std::__invoke_hypot<TYPE>(x, y);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// genint ilogb (svgenfloat x)
-template <typename T,
-          typename = std::enable_if_t<detail::is_svgenfloat<T>::value, T>>
-detail::change_base_type_t<T, int> ilogb(T x) __NOEXC {
-  return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(x);
-}
+// genint ilogb (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline detail::change_base_type_t<TYPE, int> ilogb(TYPE x) __NOEXC {         \
+    return __sycl_std::__invoke_ilogb<detail::change_base_type_t<TYPE, int>>(  \
+        x);                                                                    \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // float ldexp (float x, int k)
 // double ldexp (double x, int k)
 // half ldexp (half x, int k)
-template <typename T>
-std::enable_if_t<detail::is_sgenfloat<T>::value, T> ldexp(T x, int k) __NOEXC {
-  return __sycl_std::__invoke_ldexp<T>(x, k);
-}
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE ldexp(TYPE x, int k) __NOEXC {                                   \
+    return __sycl_std::__invoke_ldexp<TYPE>(x, k);                             \
+  }
+__SYCL_DEF_BUILTIN_SGENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // vgenfloat ldexp (vgenfloat x, int k)
-template <typename T>
-std::enable_if_t<detail::is_vgenfloat<T>::value, T> ldexp(T x, int k) __NOEXC {
-  return __sycl_std::__invoke_ldexp<T>(x, vec<int, T::size()>(k));
-}
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE ldexp(TYPE x, int k) __NOEXC {                                   \
+    return __sycl_std::__invoke_ldexp<TYPE>(x, vec<int, TYPE::size()>(k));     \
+  }
+__SYCL_DEF_BUILTIN_VGENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // vgenfloat ldexp (vgenfloat x, genint k)
 template <typename T, typename T2>
@@ -1025,11 +1085,13 @@ ldexp(T x, T2 k) __NOEXC {
   return __sycl_std::__invoke_ldexp<T>(x, k);
 }
 
-// svgenfloat lgamma (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> lgamma(T x) __NOEXC {
-  return __sycl_std::__invoke_lgamma<T>(x);
-}
+// genfloat lgamma (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE lgamma(TYPE x) __NOEXC {                                         \
+    return __sycl_std::__invoke_lgamma<TYPE>(x);                               \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat lgamma_r (svgenfloat x, genintptr signp)
 template <typename T, typename T2>
@@ -1040,54 +1102,67 @@ lgamma_r(T x, T2 signp) __NOEXC {
   return __sycl_std::__invoke_lgamma_r<T>(x, signp);
 }
 
-// svgenfloat log (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log(T x) __NOEXC {
-  return __sycl_std::__invoke_log<T>(x);
-}
+// genfloat log (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log(TYPE x) __NOEXC { return __sycl_std::__invoke_log<TYPE>(x); }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat log2 (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log2(T x) __NOEXC {
-  return __sycl_std::__invoke_log2<T>(x);
-}
+// genfloat log2 (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_log2<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat log10 (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> log10(T x) __NOEXC {
-  return __sycl_std::__invoke_log10<T>(x);
-}
+// genfloat log10 (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_log10<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat log1p (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> log1p(T x) __NOEXC {
-  return __sycl_std::__invoke_log1p<T>(x);
-}
+// genfloat log1p (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log1p(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_log1p<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat logb (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> logb(T x) __NOEXC {
-  return __sycl_std::__invoke_logb<T>(x);
-}
+// genfloat logb (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE logb(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_logb<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat mad (svgenfloat a, svgenfloat b, svgenfloat c)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> mad(T a, T b,
-                                                         T c) __NOEXC {
-  return __sycl_std::__invoke_mad<T>(a, b, c);
-}
+// 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<TYPE>(a, b, c);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat maxmag (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> maxmag(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_maxmag<T>(x, y);
-}
+// genfloat maxmag (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE maxmag(TYPE x, TYPE y) __NOEXC {                                 \
+    return __sycl_std::__invoke_maxmag<TYPE>(x, y);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat minmag (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> minmag(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_minmag<T>(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<TYPE>(x, y);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat modf (svgenfloat x, genfloatptr iptr)
 template <typename T, typename T2>
@@ -1105,18 +1180,21 @@ detail::nan_return_t<T> nan(T nancode) __NOEXC {
       detail::convert_data_type<T, detail::nan_argument_base_t<T>>()(nancode));
 }
 
-// svgenfloat nextafter (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> nextafter(T x,
-                                                               T y) __NOEXC {
-  return __sycl_std::__invoke_nextafter<T>(x, y);
-}
+// genfloat nextafter (genfloat x, genfloat y)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE nextafter(TYPE x, TYPE y) __NOEXC {                              \
+    return __sycl_std::__invoke_nextafter<TYPE>(x, y);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat pow (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> pow(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_pow<T>(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<TYPE>(x, y);                               \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat pown (svgenfloat x, genint y)
 template <typename T, typename T2>
@@ -1127,18 +1205,21 @@ pown(T x, T2 y) __NOEXC {
   return __sycl_std::__invoke_pown<T>(x, y);
 }
 
-// svgenfloat powr (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> powr(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_powr<T>(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<TYPE>(x, y);                              \
+  }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat remainder (svgenfloat x, svgenfloat y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> remainder(T x,
-                                                               T y) __NOEXC {
-  return __sycl_std::__invoke_remainder<T>(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<TYPE>(x, y);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat remquo (svgenfloat x, svgenfloat y, genintptr quo)
 template <typename T, typename T2>
@@ -1149,11 +1230,13 @@ remquo(T x, T y, T2 quo) __NOEXC {
   return __sycl_std::__invoke_remquo<T>(x, y, quo);
 }
 
-// svgenfloat rint (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> rint(T x) __NOEXC {
-  return __sycl_std::__invoke_rint<T>(x);
-}
+// genfloat rint (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE rint(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_rint<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat rootn (svgenfloat x, genint y)
 template <typename T, typename T2>
@@ -1164,23 +1247,27 @@ rootn(T x, T2 y) __NOEXC {
   return __sycl_std::__invoke_rootn<T>(x, y);
 }
 
-// svgenfloat round (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> round(T x) __NOEXC {
-  return __sycl_std::__invoke_round<T>(x);
-}
+// genfloat round (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE round(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_round<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat rsqrt (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> rsqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_rsqrt<T>(x);
-}
+// genfloat rsqrt (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE rsqrt(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_rsqrt<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat sin (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sin(T x) __NOEXC {
-  return __sycl_std::__invoke_sin<T>(x);
-}
+// genfloat sin (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sin(TYPE x) __NOEXC { return __sycl_std::__invoke_sin<TYPE>(x); }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // svgenfloat sincos (svgenfloat x, genfloatptr cosval)
 template <typename T, typename T2>
@@ -1191,53 +1278,67 @@ sincos(T x, T2 cosval) __NOEXC {
   return __sycl_std::__invoke_sincos<T>(x, cosval);
 }
 
-// svgenfloat sinh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> sinh(T x) __NOEXC {
-  return __sycl_std::__invoke_sinh<T>(x);
-}
+// genfloat sinh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sinh(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_sinh<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat sinpi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> sinpi(T x) __NOEXC {
-  return __sycl_std::__invoke_sinpi<T>(x);
-}
+// genfloat sinpi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sinpi(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_sinpi<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat sqrt (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_sqrt<T>(x);
-}
+// genfloat sqrt (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sqrt(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_sqrt<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat tan (svgenfloat x)
-template <typename T>
-std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> tan(T x) __NOEXC {
-  return __sycl_std::__invoke_tan<T>(x);
-}
+// genfloat tan (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tan(TYPE x) __NOEXC { return __sycl_std::__invoke_tan<TYPE>(x); }
+__SYCL_DEF_BUILTIN_FAST_MATH_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat tanh (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> tanh(T x) __NOEXC {
-  return __sycl_std::__invoke_tanh<T>(x);
-}
+// genfloat tanh (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tanh(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_tanh<TYPE>(x);                                 \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat tanpi (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> tanpi(T x) __NOEXC {
-  return __sycl_std::__invoke_tanpi<T>(x);
-}
+// genfloat tanpi (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tanpi(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_tanpi<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat tgamma (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> tgamma(T x) __NOEXC {
-  return __sycl_std::__invoke_tgamma<T>(x);
-}
+// genfloat tgamma (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tgamma(TYPE x) __NOEXC {                                         \
+    return __sycl_std::__invoke_tgamma<TYPE>(x);                               \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloat trunc (svgenfloat x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloat<T>::value, T> trunc(T x) __NOEXC {
-  return __sycl_std::__invoke_trunc<T>(x);
-}
+// genfloat trunc (genfloat x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE trunc(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_trunc<TYPE>(x);                                \
+  }
+__SYCL_DEF_BUILTIN_GENFLOAT
+#undef __SYCL_BUILTIN_DEF
 
 // other marray math functions
 
@@ -2651,89 +2752,117 @@ __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(powr)
 
 #undef __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD
 
-// svgenfloatf cos (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
-  return __sycl_std::__invoke_native_cos<T>(x);
-}
+// genfloatf cos (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cos(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_native_cos<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf divide (svgenfloatf x, svgenfloatf y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> divide(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_native_divide<T>(x, y);
-}
+// 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<TYPE>(x, y);                     \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
-  return __sycl_std::__invoke_native_exp<T>(x);
-}
+// genfloatf exp (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_native_exp<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
-  return __sycl_std::__invoke_native_exp2<T>(x);
-}
+// genfloatf exp2 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_native_exp2<TYPE>(x);                          \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
-  return __sycl_std::__invoke_native_exp10<T>(x);
-}
+// genfloatf exp10 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_native_exp10<TYPE>(x);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
-  return __sycl_std::__invoke_native_log<T>(x);
-}
+// genfloatf log (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_native_log<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
-  return __sycl_std::__invoke_native_log2<T>(x);
-}
+// genfloatf log2 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_native_log2<TYPE>(x);                          \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
-  return __sycl_std::__invoke_native_log10<T>(x);
-}
+// genfloatf log10 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_native_log10<TYPE>(x);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf powr (svgenfloatf x, svgenfloatf y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_native_powr<T>(x, y);
-}
+// 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<TYPE>(x, y);                       \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf recip (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> recip(T x) __NOEXC {
-  return __sycl_std::__invoke_native_recip<T>(x);
-}
+// genfloatf recip (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE recip(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_native_recip<TYPE>(x);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf rsqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_native_rsqrt<T>(x);
-}
+// genfloatf rsqrt (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE rsqrt(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_native_rsqrt<TYPE>(x);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf sin (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
-  return __sycl_std::__invoke_native_sin<T>(x);
-}
+// genfloatf sin (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sin(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_native_sin<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf sqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_native_sqrt<T>(x);
-}
+// genfloatf sqrt (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sqrt(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_native_sqrt<TYPE>(x);                          \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf tan (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
-  return __sycl_std::__invoke_native_tan<T>(x);
-}
+// genfloatf tan (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tan(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_native_tan<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
 } // namespace native
 namespace half_precision {
@@ -2791,89 +2920,117 @@ __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(powr)
 
 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD
 
-// svgenfloatf cos (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
-  return __sycl_std::__invoke_half_cos<T>(x);
-}
+// genfloatf cos (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE cos(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_half_cos<TYPE>(x);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf divide (svgenfloatf x, svgenfloatf y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> divide(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_half_divide<T>(x, y);
-}
+// 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<TYPE>(x, y);                       \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
-  return __sycl_std::__invoke_half_exp<T>(x);
-}
+// genfloatf exp (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_half_exp<TYPE>(x);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
-  return __sycl_std::__invoke_half_exp2<T>(x);
-}
+// genfloatf exp2 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_half_exp2<TYPE>(x);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf exp10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
-  return __sycl_std::__invoke_half_exp10<T>(x);
-}
+// genfloatf exp10 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE exp10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_half_exp10<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
-  return __sycl_std::__invoke_half_log<T>(x);
-}
+// genfloatf log (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_half_log<TYPE>(x);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
-  return __sycl_std::__invoke_half_log2<T>(x);
-}
+// genfloatf log2 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log2(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_half_log2<TYPE>(x);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf log10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
-  return __sycl_std::__invoke_half_log10<T>(x);
-}
+// genfloatf log10 (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE log10(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_half_log10<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf powr (svgenfloatf x, svgenfloatf y)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x, T y) __NOEXC {
-  return __sycl_std::__invoke_half_powr<T>(x, y);
-}
+// 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<TYPE>(x, y);                         \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf recip (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> recip(T x) __NOEXC {
-  return __sycl_std::__invoke_half_recip<T>(x);
-}
+// genfloatf recip (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE recip(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_half_recip<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf rsqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_half_rsqrt<T>(x);
-}
+// genfloatf rsqrt (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE rsqrt(TYPE x) __NOEXC {                                          \
+    return __sycl_std::__invoke_half_rsqrt<TYPE>(x);                           \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf sin (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
-  return __sycl_std::__invoke_half_sin<T>(x);
-}
+// genfloatf sin (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sin(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_half_sin<TYPE>(x);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf sqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
-  return __sycl_std::__invoke_half_sqrt<T>(x);
-}
+// genfloatf sqrt (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE sqrt(TYPE x) __NOEXC {                                           \
+    return __sycl_std::__invoke_half_sqrt<TYPE>(x);                            \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
-// svgenfloatf tan (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
-  return __sycl_std::__invoke_half_tan<T>(x);
-}
+// genfloatf tan (genfloatf x)
+#define __SYCL_BUILTIN_DEF(TYPE)                                               \
+  inline TYPE tan(TYPE x) __NOEXC {                                            \
+    return __sycl_std::__invoke_half_tan<TYPE>(x);                             \
+  }
+__SYCL_DEF_BUILTIN_GENFLOATF
+#undef __SYCL_BUILTIN_DEF
 
 } // namespace half_precision
 
@@ -2901,84 +3058,85 @@ __SYCL_MATH_FUNCTION_OVERLOAD_FM(sqrt)
 __SYCL_MATH_FUNCTION_OVERLOAD_FM(rsqrt)
 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
 
-template <typename T, size_t N>
-inline __SYCL_ALWAYS_INLINE
-    std::enable_if_t<std::is_same_v<T, float>, marray<T, N>>
-    powr(marray<T, N> x, marray<T, N> y) __NOEXC {
-  return native::powr(x, y);
-}
-
-// svgenfloatf cos (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
-  return native::cos(x);
-}
+// 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
 
-// svgenfloatf exp (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {
-  return native::exp(x);
-}
+// 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 exp2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp2(T x) __NOEXC {
-  return native::exp2(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 exp10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp10(T x) __NOEXC {
-  return native::exp10(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 log(svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log(T x) __NOEXC {
-  return native::log(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 log2 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log2(T x) __NOEXC {
-  return native::log2(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 log10 (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> log10(T x) __NOEXC {
-  return native::log10(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 powr (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> powr(T x, T y) __NOEXC {
+// genfloatf powr (genfloatf x, genfloatf y)
+// TODO: remove when __SYCL_DEF_BUILTIN_MARRAY is defined
+template <typename T, size_t N>
+inline __SYCL_ALWAYS_INLINE
+    std::enable_if_t<std::is_same_v<T, float>, marray<T, N>>
+    powr(marray<T, N> x, marray<T, N> y) __NOEXC {
   return native::powr(x, y);
 }
 
-// svgenfloatf rsqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> rsqrt(T x) __NOEXC {
-  return native::rsqrt(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 sin (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sin(T x) __NOEXC {
-  return native::sin(x);
-}
+// 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 sqrt (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> sqrt(T x) __NOEXC {
-  return native::sqrt(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 tan (svgenfloatf x)
-template <typename T>
-std::enable_if_t<detail::is_svgenfloatf<T>::value, T> tan(T x) __NOEXC {
-  return native::tan(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
+
+// 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
 
 #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 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<float, SZ> 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<float, SZ> 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<float, SZ> 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/include/sycl/ext/oneapi/experimental/sycl_complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp
index cc2e83d754287..792e5ea61ddbd 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<_ValueType>(0, __re);
+  return sycl::atan2(static_cast<_ValueType>(0), __re);
 }
 
 // norm
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<half>::epsilon() ||
+  return sycl::fabs(2 * (a - b) / (a + b)) <
+             std::numeric_limits<half>::epsilon() ||
          a < std::numeric_limits<half>::min();
 }
 
@@ -180,10 +181,10 @@ int main() {
       auto err = err_buf.get_access<access::mode::write>(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<Format>::rep_elem_type *RefData,
   CoordT<ImageType::Sampled, Dims> AdjCoord = Coord;
   if constexpr (AddrMode == addressing_mode::repeat) {
     assert(Normalized);
-    AdjCoord -= floor(AdjCoord);
+    AdjCoord -= sycl::floor(AdjCoord);
     AdjCoord *= RangeToCoord<ImageType::Sampled, Dims>(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<float>(BUFFER_SIZE)),
+                               sycl::sqrt(static_cast<float>(BUFFER_SIZE)));
   std::vector<sycl::int4> 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));