[SYCL] Split sycl::clamp implementation in two (#12637)

SYCL2020, revision 8 defines two separate clamp functions - once in
integer functions and once in common functions. Follow the same in the
implementation so that clamp's handling is uniform with other
common/integer functions.

Author: aelovikov-intel

sycl/include/sycl/detail/builtins/common_functions.inc

@@ -72,32 +72,15 @@ min(T x, detail::get_elem_type_t<T> y) {
              detail::simplify_if_swizzle_t<T>{y});
 }
 
-#undef BUILTIN_COMMON
-
-#ifdef __SYCL_DEVICE_ONLY__
-DEVICE_IMPL_TEMPLATE(THREE_ARGS, clamp, builtin_enable_generic_t,
-                     [](auto... xs) {
-                       using ElemTy = detail::get_elem_type_t<T0>;
-                       if constexpr (std::is_integral_v<ElemTy>) {
-                         if constexpr (std::is_signed_v<ElemTy>) {
-                           return __spirv_ocl_s_clamp(xs...);
-                         } else {
-                           return __spirv_ocl_u_clamp(xs...);
-                         }
-                       } else {
-                         return __spirv_ocl_fclamp(xs...);
-                       }
-                     })
-#else
-HOST_IMPL_TEMPLATE(THREE_ARGS, clamp, builtin_enable_generic_t, common,
-                   default_ret_type)
-#endif
+BUILTIN_COMMON(THREE_ARGS, clamp, __spirv_ocl_fclamp)
 template <typename T>
-detail::builtin_enable_generic_non_scalar_t<T>
+detail::builtin_enable_common_non_scalar_t<T>
 clamp(T x, detail::get_elem_type_t<T> y, detail::get_elem_type_t<T> z) {
   return clamp(detail::simplify_if_swizzle_t<T>{x},
                detail::simplify_if_swizzle_t<T>{y},
                detail::simplify_if_swizzle_t<T>{z});
 }
+
+#undef BUILTIN_COMMON
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/detail/builtins/integer_functions.inc

@@ -122,6 +122,15 @@ min(T x, detail::get_elem_type_t<T> y) {
              detail::simplify_if_swizzle_t<T>{y});
 }
 
+BUILTIN_GENINT_SU(THREE_ARGS, clamp)
+template <typename T>
+detail::builtin_enable_integer_non_scalar_t<T>
+clamp(T x, detail::get_elem_type_t<T> y, detail::get_elem_type_t<T> z) {
+  return clamp(detail::simplify_if_swizzle_t<T>{x},
+               detail::simplify_if_swizzle_t<T>{y},
+               detail::simplify_if_swizzle_t<T>{z});
+}
+
 BUILTIN_GENINT(ONE_ARG, clz)
 BUILTIN_GENINT(ONE_ARG, ctz)
 BUILTIN_GENINT(ONE_ARG, popcount)

sycl/source/builtins/common_functions.cpp

@@ -63,16 +63,8 @@ BUILTIN_COMMON(TWO_ARGS, max,
 BUILTIN_COMMON(TWO_ARGS, min,
                [](auto x, auto y) -> decltype(x) { return (y < x ? y : x); })
 
-// clamp is implemented for INTEGER_TYPES as well, so expand/inline
-// BUILTIN_COMMON manually.
-HOST_IMPL(clamp, [](auto x, auto y, auto z) -> decltype(x) {
-  using ElemTy = detail::get_elem_type_t<decltype(x)>;
-  if constexpr (std::is_integral_v<ElemTy>) {
-    return std::min(std::max(x, y), z);
-  } else {
-    return std::fmin(std::fmax(x, y), z);
-  }
+BUILTIN_COMMON(THREE_ARGS, clamp, [](auto x, auto y, auto z) -> decltype(x) {
+  return std::fmin(std::fmax(x, y), z);
 })
-EXPORT_SCALAR_AND_VEC_1_16(THREE_ARGS, clamp, INTEGER_TYPES, FP_TYPES)
 } // namespace _V1
 } // namespace sycl

sycl/source/builtins/integer_functions.cpp

@@ -214,6 +214,10 @@ BUILTIN_GENINT_SU(TWO_ARGS, max,
 BUILTIN_GENINT_SU(TWO_ARGS, min,
                   [](auto x, auto y) -> decltype(x) { return y < x ? y : x; })
 
+BUILTIN_GENINT_SU(THREE_ARGS, clamp, [](auto x, auto y, auto z) -> decltype(x) {
+  return std::min(std::max(x, y), z);
+})
+
 template <typename T> static inline constexpr T __clz_impl(T x, T m, T n = 0) {
   return (x & m) ? n : __clz_impl(x, T(m >> 1), ++n);
 }