From a7bf1ca4e414ca10bcd8f398e03d85c05b6f5346 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Thu, 6 Feb 2025 15:11:45 -0800
Subject: [PATCH] [NFCI][SYCL] Eliminate `sycl/builtins_utils_*.hpp`

Split the helper they used to provide between the builtins
implementation and `sycl/detail/type_traits/vec_marray_traits.hpp`.
---
 sycl/include/sycl/builtins_utils_scalar.hpp   | 73 -----------------
 sycl/include/sycl/builtins_utils_vec.hpp      | 79 ------------------
 .../include/sycl/detail/builtins/builtins.hpp | 81 ++++++++++++++++++-
 .../detail/type_traits/vec_marray_traits.hpp  | 16 ++++
 .../ext/oneapi/experimental/bfloat16_math.hpp |  1 -
 .../sycl/ext/oneapi/experimental/builtins.hpp |  1 -
 6 files changed, 96 insertions(+), 155 deletions(-)
 delete mode 100644 sycl/include/sycl/builtins_utils_scalar.hpp
 delete mode 100644 sycl/include/sycl/builtins_utils_vec.hpp

diff --git a/sycl/include/sycl/builtins_utils_scalar.hpp b/sycl/include/sycl/builtins_utils_scalar.hpp
deleted file mode 100644
index 7bb9e7b2ec502..0000000000000
--- a/sycl/include/sycl/builtins_utils_scalar.hpp
+++ /dev/null
@@ -1,73 +0,0 @@
-//==--- builtins_utils_scalar.hpp - SYCL built-in function utilities -------==//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#pragma once
-
-#include <sycl/access/access.hpp>
-#include <sycl/aliases.hpp>
-#include <sycl/detail/defines_elementary.hpp>
-#include <sycl/detail/generic_type_traits.hpp>
-#include <sycl/detail/type_traits.hpp>
-#include <sycl/half_type.hpp>
-#include <sycl/multi_ptr.hpp>
-
-#include <algorithm>
-#include <cstring>
-
-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 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 T>
-using simplify_if_swizzle_t = typename simplify_if_swizzle<T>::type;
-
-// 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;
-
-} // namespace detail
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/include/sycl/builtins_utils_vec.hpp b/sycl/include/sycl/builtins_utils_vec.hpp
deleted file mode 100644
index 3fd2336008654..0000000000000
--- a/sycl/include/sycl/builtins_utils_vec.hpp
+++ /dev/null
@@ -1,79 +0,0 @@
-//==--- builtins_utils_vec.hpp - SYCL built-in function utilities for vec --==//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#pragma once
-
-#include <sycl/detail/type_traits/vec_marray_traits.hpp>
-
-#include <sycl/builtins_utils_scalar.hpp>
-
-#include <sycl/detail/type_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 {
-// Utility for converting a swizzle to a vector or preserve the type if it isn't
-// a swizzle.
-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)>;
-};
-
-// 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> struct change_elements {
-  using type = NewElemT;
-};
-template <typename NewElemT, typename T, size_t N>
-struct change_elements<NewElemT, marray<T, N>> {
-  using type = marray<typename change_elements<NewElemT, T>::type, N>;
-};
-template <typename NewElemT, typename T, int N>
-struct change_elements<NewElemT, vec<T, N>> {
-  using type = vec<typename change_elements<NewElemT, T>::type, N>;
-};
-template <typename NewElemT, typename VecT, typename OperationLeftT,
-          typename OperationRightT, template <typename> class OperationCurrentT,
-          int... Indexes>
-struct change_elements<NewElemT,
-                       SwizzleOp<VecT, OperationLeftT, OperationRightT,
-                                 OperationCurrentT, Indexes...>> {
-  // Converts to vec for simplicity.
-  using type =
-      vec<typename change_elements<NewElemT, typename VecT::element_type>::type,
-          sizeof...(Indexes)>;
-};
-
-template <typename NewElemT, typename T>
-using change_elements_t = typename change_elements<NewElemT, T>::type;
-
-// 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 detail
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp
index 1f291fc4f60ac..4ce211cb1657b 100644
--- a/sycl/include/sycl/detail/builtins/builtins.hpp
+++ b/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) {
diff --git a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
index 86e8764dc87a0..6ce39bf6a072a 100644
--- a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
+++ b/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 {};
diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp
index 11dec596b1f45..49ed23008607e 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp
+++ b/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
diff --git a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
index d42df1fee26c8..e6fdba4cdad45 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
+++ b/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