diff --git a/sycl/include/sycl/builtins_scalar.hpp b/sycl/include/sycl/builtins_scalar.hpp
new file mode 100644
index 0000000000000..760e63689a48b
--- /dev/null
+++ b/sycl/include/sycl/builtins_scalar.hpp
@@ -0,0 +1,31 @@
+//==------- builtins_scalar.hpp - SYCL scalar built-in functions -----------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Lightweight alternative to <sycl/builtins.hpp> that provides only the
+// scalar (float / double / half / integer fundamental-type) overloads of all
+// SYCL built-in functions, without pulling in sycl::vec<> or sycl::marray<>
+// template machinery.
+//
+// Trade-offs vs <sycl/builtins.hpp>:
+//   - Significantly less front-end parse work (no vector.hpp, multi_ptr.hpp,
+//     marray.hpp, vector_convert.hpp)
+//   - Covers: sycl::sin(float), sycl::native::exp(float),
+//             sycl::fmax(float,float), sycl::abs(int),
+//             sycl::isequal(float,float), etc.
+//   - Does NOT cover: sycl::sin(sycl::float4),
+//             sycl::native::exp(sycl::marray<float,8>), etc.
+//     Use <sycl/builtins.hpp> for those.
+//
+// Including both this header and <sycl/builtins.hpp> in the same translation
+// unit is safe: the scalar inline functions are identical and ODR-compliant.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/builtins/scalar_builtins.hpp>
diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp
index d69c0b93ce7c1..0eed93ef6fb6f 100644
--- a/sycl/include/sycl/detail/builtins/builtins.hpp
+++ b/sycl/include/sycl/detail/builtins/builtins.hpp
@@ -63,103 +63,15 @@
 
 #pragma once
 
-#include <sycl/detail/helpers.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>
-#include <sycl/multi_ptr.hpp>
-#include <sycl/vector.hpp>
+#include <sycl/detail/builtins/gentype_utilities.hpp>
+#include <sycl/detail/builtins/scalar_infrastructure.hpp>
+#include <sycl/detail/memcpy.hpp>
+
+#include <algorithm>
 
 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>) ||
-     (... && is_vec_or_swizzle_v<Ts>)) &&
-    (... && (num_elements<Ts>::value ==
-             num_elements<typename first_type<Ts...>::type>::value));
-
-template <typename... Ts>
-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__
@@ -187,34 +99,6 @@ template <typename T> auto convert_arg(T &&x) {
 #endif
 } // namespace builtins
 
-template <typename FuncTy, typename... Ts>
-auto builtin_marray_impl(FuncTy F, const Ts &...x) {
-  using ret_elem_type = decltype(F(x[0]...));
-  using T = typename first_type<Ts...>::type;
-  marray<ret_elem_type, T::size()> Res;
-  constexpr auto N = T::size();
-  for (size_t I = 0; I < N / 2; ++I) {
-    auto PartialRes = [&]() {
-      using elem_ty = get_elem_type_t<T>;
-      if constexpr (std::is_integral_v<elem_ty>)
-        return F(
-            to_vec2(x, I * 2)
-                .template as<vec<
-                    std::conditional_t<std::is_signed_v<elem_ty>,
-                                       fixed_width_signed<sizeof(elem_ty)>,
-                                       fixed_width_unsigned<sizeof(elem_ty)>>,
-                    2>>()...);
-      else
-        return F(to_vec2(x, I * 2)...);
-    }();
-    sycl::detail::memcpy_no_adl(&Res[I * 2], &PartialRes,
-                                sizeof(decltype(PartialRes)));
-  }
-  if (N % 2)
-    Res[N - 1] = F(x[N - 1]...);
-  return Res;
-}
-
 template <typename FuncTy, typename... Ts>
 auto builtin_default_host_impl(FuncTy F, const Ts &...x) {
   // We implement support for marray/swizzle in the headers and export symbols
@@ -229,80 +113,6 @@ auto builtin_default_host_impl(FuncTy F, const Ts &...x) {
     return F(simplify_if_swizzle_t<Ts>{x}...);
   }
 }
-
-template <typename FuncTy, typename... Ts>
-auto builtin_delegate_to_scalar(FuncTy F, const Ts &...x) {
-  using T = typename first_type<Ts...>::type;
-  static_assert(is_vec_or_swizzle_v<T> || is_marray_v<T>);
-
-  constexpr auto Size = T::size();
-  using ret_elem_type = decltype(F(x[0]...));
-  std::conditional_t<is_marray_v<T>, marray<ret_elem_type, Size>,
-                     vec<ret_elem_type, Size>>
-      r{};
-
-  if constexpr (is_marray_v<T>) {
-    for (size_t i = 0; i < Size; ++i)
-      r[i] = F(x[i]...);
-  } else {
-    loop<Size>([&](auto idx) { r[idx] = F(x[idx]...); });
-  }
-
-  return r;
-}
-
-template <typename T>
-struct fp_elem_type
-    : std::bool_constant<
-          check_type_in_v<get_elem_type_t<T>, float, double, half>> {};
-template <typename T>
-struct float_elem_type
-    : std::bool_constant<check_type_in_v<get_elem_type_t<T>, float>> {};
-
-template <typename... Ts>
-struct same_basic_shape : std::bool_constant<builtin_same_shape_v<Ts...>> {};
-
-template <typename... Ts>
-struct same_elem_type : std::bool_constant<same_basic_shape<Ts...>::value &&
-                                           all_same_v<get_elem_type_t<Ts>...>> {
-};
-
-template <typename> struct any_shape : std::true_type {};
-
-template <typename T>
-struct scalar_only : std::bool_constant<is_scalar_arithmetic_v<T>> {};
-
-template <typename T>
-struct non_scalar_only : std::bool_constant<!is_scalar_arithmetic_v<T>> {};
-
-template <typename T> struct default_ret_type {
-  using type = T;
-};
-
-template <typename T> struct scalar_ret_type {
-  using type = get_elem_type_t<T>;
-};
-
-template <template <typename> typename RetTypeTrait,
-          template <typename> typename ElemTypeChecker,
-          template <typename> typename ShapeChecker,
-          template <typename...> typename ExtraConditions, typename... Ts>
-struct builtin_enable
-    : std::enable_if<
-          ElemTypeChecker<typename first_type<Ts...>::type>::value &&
-              ShapeChecker<typename first_type<Ts...>::type>::value &&
-              ExtraConditions<Ts...>::value,
-          typename RetTypeTrait<
-              simplify_if_swizzle_t<typename first_type<Ts...>::type>>::type> {
-};
-#define BUILTIN_CREATE_ENABLER(NAME, RET_TYPE_TRAIT, ELEM_TYPE_CHECKER,        \
-                               SHAPE_CHECKER, EXTRA_CONDITIONS)                \
-  namespace detail {                                                           \
-  template <typename... Ts>                                                    \
-  using NAME##_t =                                                             \
-      typename builtin_enable<RET_TYPE_TRAIT, ELEM_TYPE_CHECKER,               \
-                              SHAPE_CHECKER, EXTRA_CONDITIONS, Ts...>::type;   \
-  }
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/detail/builtins/common_functions.inc b/sycl/include/sycl/detail/builtins/common_functions.inc
index 8c8e3e46277a9..87e065ff8ca37 100644
--- a/sycl/include/sycl/detail/builtins/common_functions.inc
+++ b/sycl/include/sycl/detail/builtins/common_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/gentype_utilities.hpp b/sycl/include/sycl/detail/builtins/gentype_utilities.hpp
new file mode 100644
index 0000000000000..617165af5616c
--- /dev/null
+++ b/sycl/include/sycl/detail/builtins/gentype_utilities.hpp
@@ -0,0 +1,114 @@
+//==------------- gentype_utilities.hpp -----------------------------------==//
+//
+// 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/builtins/scalar_infrastructure.hpp>
+#include <sycl/detail/memcpy.hpp>
+
+#ifndef SYCL_DETAIL_BUILTINS_SCALAR_ONLY
+#include <sycl/detail/vector_convert.hpp>
+#include <sycl/marray.hpp>
+#include <sycl/multi_ptr.hpp>
+#include <sycl/vector.hpp>
+#endif
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+// Utility functions for converting to/from vec/marray.
+#ifndef SYCL_DETAIL_BUILTINS_SCALAR_ONLY
+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;
+}
+#else
+template <class T, size_t N> void to_vec2(marray<T, N>, size_t) = delete;
+template <class T, size_t N> void to_vec(marray<T, N>) = delete;
+template <class T, int N> void to_marray(vec<T, N>) = delete;
+#endif
+
+#ifndef SYCL_DETAIL_BUILTINS_SCALAR_ONLY
+template <typename FuncTy, typename... Ts>
+auto builtin_marray_impl(FuncTy F, const Ts &...x) {
+  using ret_elem_type = decltype(F(x[0]...));
+  using T = typename first_type<Ts...>::type;
+  marray<ret_elem_type, T::size()> Res;
+  constexpr auto N = T::size();
+  for (size_t I = 0; I < N / 2; ++I) {
+    auto PartialRes = [&]() {
+      using elem_ty = get_elem_type_t<T>;
+      if constexpr (std::is_integral_v<elem_ty>) {
+        return F(
+            to_vec2(x, I * 2)
+                .template as<vec<
+                    std::conditional_t<std::is_signed_v<elem_ty>,
+                                       fixed_width_signed<sizeof(elem_ty)>,
+                                       fixed_width_unsigned<sizeof(elem_ty)>>,
+                    2>>()...);
+      } else {
+        return F(to_vec2(x, I * 2)...);
+      }
+    }();
+    sycl::detail::memcpy_no_adl(&Res[I * 2], &PartialRes,
+                                sizeof(decltype(PartialRes)));
+  }
+  if (N % 2) {
+    Res[N - 1] = F(x[N - 1]...);
+  }
+  return Res;
+}
+#else
+template <typename FuncTy, typename... Ts>
+void builtin_marray_impl(FuncTy, const Ts &...) = delete;
+#endif
+
+#ifndef SYCL_DETAIL_BUILTINS_SCALAR_ONLY
+template <typename FuncTy, typename... Ts>
+auto builtin_delegate_to_scalar(FuncTy F, const Ts &...x) {
+  using T = typename first_type<Ts...>::type;
+  static_assert(is_vec_or_swizzle_v<T> || is_marray_v<T>);
+
+  constexpr auto Size = T::size();
+  using ret_elem_type = decltype(F(x[0]...));
+  std::conditional_t<is_marray_v<T>, marray<ret_elem_type, Size>,
+                     vec<ret_elem_type, Size>>
+      r{};
+
+  if constexpr (is_marray_v<T>) {
+    for (size_t i = 0; i < Size; ++i) {
+      r[i] = F(x[i]...);
+    }
+  } else {
+    loop<Size>([&](auto idx) { r[idx] = F(x[idx]...); });
+  }
+
+  return r;
+}
+#else
+template <typename FuncTy, typename... Ts>
+void builtin_delegate_to_scalar(FuncTy, const Ts &...) = delete;
+#endif
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
\ No newline at end of file
diff --git a/sycl/include/sycl/detail/builtins/geometric_functions.inc b/sycl/include/sycl/detail/builtins/geometric_functions.inc
index 1ee7d34fded87..93006d61c5ca9 100644
--- a/sycl/include/sycl/detail/builtins/geometric_functions.inc
+++ b/sycl/include/sycl/detail/builtins/geometric_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/half_precision_math_functions.inc b/sycl/include/sycl/detail/builtins/half_precision_math_functions.inc
index 2717f6f72071b..970fce2e88341 100644
--- a/sycl/include/sycl/detail/builtins/half_precision_math_functions.inc
+++ b/sycl/include/sycl/detail/builtins/half_precision_math_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/helper_macros.hpp b/sycl/include/sycl/detail/builtins/helper_macros.hpp
index 3b320d4e0772a..b6fe8f50aef46 100644
--- a/sycl/include/sycl/detail/builtins/helper_macros.hpp
+++ b/sycl/include/sycl/detail/builtins/helper_macros.hpp
@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <sycl/detail/builtins/scalar_infrastructure.hpp>
+
 // Usage:
 //   #define HANDLE_TYPE(INVARIANT_ARG1, INVARIANT_ARG2, TYPE) ...
 //   FOR_EACH2(HANDLE_TYPE, A1, A2, TYPE1, TYPE2, ...)
@@ -184,41 +186,6 @@
 #define FIXED_WIDTH_INTEGER_TYPES                                              \
   int8_t, int16_t, int32_t, int64_t, uint8_t, uint16_t, uint32_t, uint64_t
 
-// Use (NAME)/(NS::NAME) to deal win min/max macros in windows.h throughout this
-// file.
-
-#define DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE(                                  \
-    NUM_ARGS, NAME, ENABLER, DELEGATOR, NS, /*SCALAR_VEC_IMPL*/...)            \
-  template <NUM_ARGS##_TYPENAME_TYPE>                                          \
-  detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>(NAME)(                             \
-      NUM_ARGS##_TEMPLATE_TYPE_ARG) {                                          \
-    using ToTy = detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>;                    \
-    if constexpr (detail::is_marray_v<T0>) {                                   \
-      return detail::DELEGATOR(                                                \
-          [](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); },      \
-          NUM_ARGS##_ARG);                                                     \
-    } else if constexpr (detail::is_vec_v<ToTy>) {                             \
-      if constexpr (ToTy::size() == 3) {                                       \
-        /* For vectors of length 3, make sure to only copy 3 elements, not 4,  \
-           to work around code generation issues, see LLVM #144454. */         \
-        auto From = __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG);                     \
-        ToTy To;                                                               \
-        constexpr auto N =                                                     \
-            ToTy::size() * sizeof(detail::get_elem_type_t<ToTy>);              \
-        sycl::detail::memcpy_no_adl(&To, &From, N);                            \
-        return To;                                                             \
-      } else {                                                                 \
-        return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));          \
-      }                                                                        \
-    } else {                                                                   \
-      return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));            \
-    }                                                                          \
-  }
-
-#define DEVICE_IMPL_TEMPLATE(NUM_ARGS, NAME, ENABLER, /*SCALAR_VEC_IMPL*/...)  \
-  DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE(NUM_ARGS, NAME, ENABLER,                \
-                                       builtin_marray_impl, sycl, __VA_ARGS__)
-
 #ifdef __SYCL_BUILD_SYCL_DLL
 #define SYCL_BUILTIN_EXPORT __SYCL_EXPORT
 #else
diff --git a/sycl/include/sycl/detail/builtins/helper_macros_gentype.hpp b/sycl/include/sycl/detail/builtins/helper_macros_gentype.hpp
new file mode 100644
index 0000000000000..e9577fbfd5fb7
--- /dev/null
+++ b/sycl/include/sycl/detail/builtins/helper_macros_gentype.hpp
@@ -0,0 +1,46 @@
+//==-- helper_macros_gentype.hpp -- Gentype helper macros for sycl builtins ==//
+//
+// 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/builtins/helper_macros.hpp>
+
+// Use (NAME)/(NS::NAME) to deal win min/max macros in windows.h throughout
+// this file.
+
+#define DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE(                                  \
+    NUM_ARGS, NAME, ENABLER, DELEGATOR, NS, /*SCALAR_VEC_IMPL*/...)            \
+  template <NUM_ARGS##_TYPENAME_TYPE>                                          \
+  detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>(NAME)(                             \
+      NUM_ARGS##_TEMPLATE_TYPE_ARG) {                                          \
+    using ToTy = detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>;                    \
+    if constexpr (detail::is_marray_v<T0>) {                                   \
+      return detail::DELEGATOR(                                                \
+          [](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); },      \
+          NUM_ARGS##_ARG);                                                     \
+    } else if constexpr (detail::is_vec_v<ToTy>) {                             \
+      if constexpr (ToTy::size() == 3) {                                       \
+        /* For vectors of length 3, make sure to only copy 3 elements, not 4,  \
+           to work around code generation issues, see LLVM #144454. */         \
+        auto From = __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG);                     \
+        ToTy To;                                                               \
+        constexpr auto N =                                                     \
+            ToTy::size() * sizeof(detail::get_elem_type_t<ToTy>);              \
+        sycl::detail::memcpy_no_adl(&To, &From, N);                            \
+        return To;                                                             \
+      } else {                                                                 \
+        return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));          \
+      }                                                                        \
+    } else {                                                                   \
+      return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));            \
+    }                                                                          \
+  }
+
+#define DEVICE_IMPL_TEMPLATE(NUM_ARGS, NAME, ENABLER, /*SCALAR_VEC_IMPL*/...)  \
+  DEVICE_IMPL_TEMPLATE_CUSTOM_DELEGATE(NUM_ARGS, NAME, ENABLER,                \
+                                       builtin_marray_impl, sycl, __VA_ARGS__)
diff --git a/sycl/include/sycl/detail/builtins/integer_functions.inc b/sycl/include/sycl/detail/builtins/integer_functions.inc
index 65b91160ac260..8370b12d92a88 100644
--- a/sycl/include/sycl/detail/builtins/integer_functions.inc
+++ b/sycl/include/sycl/detail/builtins/integer_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/math_functions.inc b/sycl/include/sycl/detail/builtins/math_functions.inc
index 8387fe09e6b69..fe91f58b85ef6 100644
--- a/sycl/include/sycl/detail/builtins/math_functions.inc
+++ b/sycl/include/sycl/detail/builtins/math_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/native_math_functions.inc b/sycl/include/sycl/detail/builtins/native_math_functions.inc
index 01cb6f43e2c59..65c252779153d 100644
--- a/sycl/include/sycl/detail/builtins/native_math_functions.inc
+++ b/sycl/include/sycl/detail/builtins/native_math_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/relational_functions.inc b/sycl/include/sycl/detail/builtins/relational_functions.inc
index fea8f59e2993f..cd6ea0889e316 100644
--- a/sycl/include/sycl/detail/builtins/relational_functions.inc
+++ b/sycl/include/sycl/detail/builtins/relational_functions.inc
@@ -8,7 +8,7 @@
 
 // Intentionally insufficient set of includes and no "#pragma once".
 
-#include <sycl/detail/builtins/helper_macros.hpp>
+#include <sycl/detail/builtins/helper_macros_gentype.hpp>
 
 namespace sycl {
 inline namespace _V1 {
diff --git a/sycl/include/sycl/detail/builtins/scalar_builtins.hpp b/sycl/include/sycl/detail/builtins/scalar_builtins.hpp
new file mode 100644
index 0000000000000..4dd4e0ca0e5b4
--- /dev/null
+++ b/sycl/include/sycl/detail/builtins/scalar_builtins.hpp
@@ -0,0 +1,14 @@
+//==--------------- scalar_builtins.hpp -----------------------------------==//
+//
+// 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
+
+// Internal scalar-only entry point for SYCL builtins.
+#define SYCL_DETAIL_BUILTINS_SCALAR_ONLY
+#include <sycl/detail/builtins/builtins.hpp>
+#undef SYCL_DETAIL_BUILTINS_SCALAR_ONLY
\ No newline at end of file
diff --git a/sycl/include/sycl/detail/builtins/scalar_infrastructure.hpp b/sycl/include/sycl/detail/builtins/scalar_infrastructure.hpp
new file mode 100644
index 0000000000000..e34096e1cbc01
--- /dev/null
+++ b/sycl/include/sycl/detail/builtins/scalar_infrastructure.hpp
@@ -0,0 +1,145 @@
+//==----------- scalar_infrastructure.hpp ---------------------------------==//
+//
+// 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/generic_type_traits.hpp>
+#include <sycl/detail/helpers.hpp>
+#include <sycl/detail/type_traits.hpp>
+#include <sycl/detail/type_traits/vec_marray_traits.hpp>
+#include <sycl/half_type.hpp>
+
+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>) ||
+     (... && is_vec_or_swizzle_v<Ts>)) &&
+    (... && (num_elements<Ts>::value ==
+             num_elements<typename first_type<Ts...>::type>::value));
+
+template <typename... Ts>
+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>...>;
+
+template <typename T>
+struct fp_elem_type
+    : std::bool_constant<
+          check_type_in_v<get_elem_type_t<T>, float, double, half>> {};
+template <typename T>
+struct float_elem_type
+    : std::bool_constant<check_type_in_v<get_elem_type_t<T>, float>> {};
+
+template <typename... Ts>
+struct same_basic_shape : std::bool_constant<builtin_same_shape_v<Ts...>> {};
+
+template <typename... Ts>
+struct same_elem_type : std::bool_constant<same_basic_shape<Ts...>::value &&
+                                           all_same_v<get_elem_type_t<Ts>...>> {
+};
+
+template <typename> struct any_shape : std::true_type {};
+
+template <typename T>
+struct scalar_only : std::bool_constant<is_scalar_arithmetic_v<T>> {};
+
+template <typename T>
+struct non_scalar_only : std::bool_constant<!is_scalar_arithmetic_v<T>> {};
+
+template <typename T> struct default_ret_type {
+  using type = T;
+};
+
+template <typename T> struct scalar_ret_type {
+  using type = get_elem_type_t<T>;
+};
+
+template <template <typename> typename RetTypeTrait,
+          template <typename> typename ElemTypeChecker,
+          template <typename> typename ShapeChecker,
+          template <typename...> typename ExtraConditions, typename... Ts>
+struct builtin_enable
+    : std::enable_if<
+          ElemTypeChecker<typename first_type<Ts...>::type>::value &&
+              ShapeChecker<typename first_type<Ts...>::type>::value &&
+              ExtraConditions<Ts...>::value,
+          typename RetTypeTrait<
+              simplify_if_swizzle_t<typename first_type<Ts...>::type>>::type> {
+};
+
+#define BUILTIN_CREATE_ENABLER(NAME, RET_TYPE_TRAIT, ELEM_TYPE_CHECKER,        \
+                               SHAPE_CHECKER, EXTRA_CONDITIONS)                \
+  namespace detail {                                                           \
+  template <typename... Ts>                                                    \
+  using NAME##_t =                                                             \
+      typename builtin_enable<RET_TYPE_TRAIT, ELEM_TYPE_CHECKER,               \
+                              SHAPE_CHECKER, EXTRA_CONDITIONS, Ts...>::type;   \
+  }
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
\ No newline at end of file
diff --git a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
index 6b0b4af4e44f1..8f4c80ef6a5bd 100644
--- a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
@@ -10,27 +10,29 @@
 // CHECK-NEXT: feature_test.hpp
 // CHECK-NEXT: builtins.hpp
 // CHECK-NEXT: detail/builtins/builtins.hpp
+// CHECK-NEXT: detail/builtins/gentype_utilities.hpp
+// CHECK-NEXT: detail/builtins/scalar_infrastructure.hpp
+// CHECK-NEXT: detail/generic_type_traits.hpp
+// CHECK-NEXT: access/access.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/fwd/half.hpp
+// CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
+// CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: detail/helpers.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: access/access.hpp
 // CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: detail/type_traits.hpp
-// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: detail/fwd/multi_ptr.hpp
-// CHECK-NEXT: detail/vector_convert.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
-// CHECK-NEXT: detail/fwd/half.hpp
-// CHECK-NEXT: detail/memcpy.hpp
-// CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: half_type.hpp
 // CHECK-NEXT: aspects.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
+// CHECK-NEXT: detail/memcpy.hpp
+// CHECK-NEXT: detail/vector_convert.hpp
+// CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
@@ -40,6 +42,7 @@
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
+// CHECK-NEXT: detail/builtins/helper_macros_gentype.hpp
 // CHECK-NEXT: detail/builtins/helper_macros.hpp
 // CHECK-NEXT: detail/builtins/geometric_functions.inc
 // CHECK-NEXT: detail/builtins/half_precision_math_functions.inc
diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
index 024a9c10557d2..1d2c64c4b7c14 100644
--- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
@@ -97,12 +97,15 @@
 // CHECK-NEXT: detail/memcpy.hpp
 // CHECK-NEXT: builtins.hpp
 // CHECK-NEXT: detail/builtins/builtins.hpp
+// CHECK-NEXT: detail/builtins/gentype_utilities.hpp
+// CHECK-NEXT: detail/builtins/scalar_infrastructure.hpp
 // CHECK-NEXT: detail/vector_convert.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: marray.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
+// CHECK-NEXT: detail/builtins/helper_macros_gentype.hpp
 // CHECK-NEXT: detail/builtins/helper_macros.hpp
 // CHECK-NEXT: detail/builtins/geometric_functions.inc
 // CHECK-NEXT: detail/builtins/half_precision_math_functions.inc
diff --git a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
index a950b5a4aae3f..1c0cbd9d85585 100644
--- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
@@ -91,13 +91,16 @@
 // CHECK-NEXT: __spirv/spirv_ops.hpp
 // CHECK-NEXT: builtins.hpp
 // CHECK-NEXT: detail/builtins/builtins.hpp
-// CHECK-NEXT: detail/vector_convert.hpp
+// CHECK-NEXT: detail/builtins/gentype_utilities.hpp
+// CHECK-NEXT: detail/builtins/scalar_infrastructure.hpp
 // CHECK-NEXT: detail/memcpy.hpp
+// CHECK-NEXT: detail/vector_convert.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: marray.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
+// CHECK-NEXT: detail/builtins/helper_macros_gentype.hpp
 // CHECK-NEXT: detail/builtins/helper_macros.hpp
 // CHECK-NEXT: detail/builtins/geometric_functions.inc
 // CHECK-NEXT: detail/builtins/half_precision_math_functions.inc
diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
index e26c64af65dfb..b33a7420744c9 100644
--- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
@@ -11,27 +11,29 @@
 // CHECK-NEXT: usm.hpp
 // CHECK-NEXT: builtins.hpp
 // CHECK-NEXT: detail/builtins/builtins.hpp
+// CHECK-NEXT: detail/builtins/gentype_utilities.hpp
+// CHECK-NEXT: detail/builtins/scalar_infrastructure.hpp
+// CHECK-NEXT: detail/generic_type_traits.hpp
+// CHECK-NEXT: access/access.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/fwd/half.hpp
+// CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
+// CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: detail/helpers.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: access/access.hpp
 // CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: detail/type_traits.hpp
-// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: detail/fwd/multi_ptr.hpp
-// CHECK-NEXT: detail/vector_convert.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
-// CHECK-NEXT: detail/fwd/half.hpp
-// CHECK-NEXT: detail/memcpy.hpp
-// CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: half_type.hpp
 // CHECK-NEXT: aspects.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
+// CHECK-NEXT: detail/memcpy.hpp
+// CHECK-NEXT: detail/vector_convert.hpp
+// CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
@@ -41,6 +43,7 @@
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
+// CHECK-NEXT: detail/builtins/helper_macros_gentype.hpp
 // CHECK-NEXT: detail/builtins/helper_macros.hpp
 // CHECK-NEXT: detail/builtins/geometric_functions.inc
 // CHECK-NEXT: detail/builtins/half_precision_math_functions.inc