[NFCI][SYCL] Eliminate sycl/builtins_utils_*.hpp

sycl/include/sycl/detail/builtins/builtins.hpp

@@ -63,11 +63,73 @@
 
 #pragma once
 
-#include <sycl/builtins_utils_vec.hpp>
+#include <sycl/detail/type_traits.hpp>
+#include <sycl/detail/type_traits/vec_marray_traits.hpp>
+#include <sycl/detail/vector_convert.hpp>
+#include <sycl/marray.hpp> // for marray
+#include <sycl/vector.hpp> // for vec
 
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
+#ifdef __FAST_MATH__
+template <typename T>
+struct use_fast_math
+    : std::is_same<std::remove_cv_t<get_elem_type_t<T>>, float> {};
+#else
+template <typename> struct use_fast_math : std::false_type {};
+#endif
+template <typename T> constexpr bool use_fast_math_v = use_fast_math<T>::value;
+
+// Utility trait for getting the decoration of a multi_ptr.
+template <typename T> struct get_multi_ptr_decoration;
+template <typename ElementType, access::address_space Space,
+          access::decorated DecorateAddress>
+struct get_multi_ptr_decoration<
+    multi_ptr<ElementType, Space, DecorateAddress>> {
+  static constexpr access::decorated value = DecorateAddress;
+};
+
+template <typename T>
+constexpr access::decorated get_multi_ptr_decoration_v =
+    get_multi_ptr_decoration<T>::value;
+
+// Utility trait for checking if a multi_ptr has a "writable" address space,
+// i.e. global, local, private or generic.
+template <typename T> struct has_writeable_addr_space : std::false_type {};
+template <typename ElementType, access::address_space Space,
+          access::decorated DecorateAddress>
+struct has_writeable_addr_space<multi_ptr<ElementType, Space, DecorateAddress>>
+    : std::bool_constant<Space == access::address_space::global_space ||
+                         Space == access::address_space::local_space ||
+                         Space == access::address_space::private_space ||
+                         Space == access::address_space::generic_space> {};
+
+template <typename T>
+constexpr bool has_writeable_addr_space_v = has_writeable_addr_space<T>::value;
+
+// Utility trait for changing the element type of a type T. If T is a scalar,
+// the new type replaces T completely.
+template <typename NewElemT, typename T, typename = void>
+struct change_elements {
+  using type = NewElemT;
+};
+template <typename NewElemT, typename T>
+struct change_elements<NewElemT, T, std::enable_if_t<is_marray_v<T>>> {
+  using type =
+      marray<typename change_elements<NewElemT, typename T::value_type>::type,
+             T::size()>;
+};
+template <typename NewElemT, typename T>
+struct change_elements<NewElemT, T, std::enable_if_t<is_vec_or_swizzle_v<T>>> {
+  using type =
+      vec<typename change_elements<NewElemT, typename T::element_type>::type,
+          T::size()>;
+};
+
+template <typename NewElemT, typename T>
+using change_elements_t = typename change_elements<NewElemT, T>::type;
+
 template <typename... Ts>
 inline constexpr bool builtin_same_shape_v =
     ((... && is_scalar_arithmetic_v<Ts>) || (... && is_marray_v<Ts>) ||
@@ -80,6 +142,23 @@ inline constexpr bool builtin_same_or_swizzle_v =
     // Use builtin_same_shape_v to filter out types unrelated to builtins.
     builtin_same_shape_v<Ts...> && all_same_v<simplify_if_swizzle_t<Ts>...>;
 
+// Utility functions for converting to/from vec/marray.
+template <class T, size_t N> vec<T, 2> to_vec2(marray<T, N> X, size_t Start) {
+  return {X[Start], X[Start + 1]};
+}
+template <class T, size_t N> vec<T, N> to_vec(marray<T, N> X) {
+  vec<T, N> Vec;
+  for (size_t I = 0; I < N; I++)
+    Vec[I] = X[I];
+  return Vec;
+}
+template <class T, int N> marray<T, N> to_marray(vec<T, N> X) {
+  marray<T, N> Marray;
+  for (size_t I = 0; I < N; I++)
+    Marray[I] = X[I];
+  return Marray;
+}
+
 namespace builtins {
 #ifdef __SYCL_DEVICE_ONLY__
 template <typename T> auto convert_arg(T &&x) {

sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp

@@ -24,6 +24,22 @@ template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 class SwizzleOp;
 
+// Utility for converting a swizzle to a vector or preserve the type if it isn't
+// a swizzle.
+template <typename T> struct simplify_if_swizzle {
+  using type = T;
+};
+
+template <typename VecT, typename OperationLeftT, typename OperationRightT,
+          template <typename> class OperationCurrentT, int... Indexes>
+struct simplify_if_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
+                                     OperationCurrentT, Indexes...>> {
+  using type = vec<typename VecT::element_type, sizeof...(Indexes)>;
+};
+
+template <typename T>
+using simplify_if_swizzle_t = typename simplify_if_swizzle<T>::type;
+
 // --------- is_* traits ------------------ //
 template <typename> struct is_vec : std::false_type {};
 template <typename T, int N> struct is_vec<vec<T, N>> : std::true_type {};

sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp

@@ -9,7 +9,6 @@
 #pragma once
 
 #include <sycl/builtins.hpp>           // for ceil, cos, exp, exp10, exp2
-#include <sycl/builtins_utils_vec.hpp> // For simplify_if_swizzle, is_swizzle
 #include <sycl/detail/memcpy.hpp>      // sycl::detail::memcpy
 #include <sycl/detail/vector_convert.hpp>
 #include <sycl/ext/oneapi/bfloat16.hpp> // for bfloat16, bfloat16ToBits

sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

@@ -10,7 +10,6 @@
 
 #include <sycl/aliases.hpp>                    // for half
 #include <sycl/builtins.hpp>                   // for to_vec2
-#include <sycl/builtins_utils_vec.hpp>         // for to_vec, to_marray...
 #include <sycl/detail/defines_elementary.hpp>  // for __SYCL_ALWAYS_INLINE
 #include <sycl/detail/generic_type_traits.hpp> // for is_svgenfloath, is_sv...
 #include <sycl/detail/memcpy.hpp>              // detail::memcpy